From 844605d300240bddcdd813cd40ce85dc341487dc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?S=C3=A9bastien=20Crozet?= Date: Sun, 24 Nov 2024 18:08:33 +0100 Subject: [PATCH] Partial update to bevy 0.15.rc.3 + add hot-reloading + fix macos native (#2) * feat: hot-reloading * feat: start updating to bevy 0.15 + workaround for macos support * feat: update to wgcore 0.2 --- Cargo.toml | 39 +---- crates/wgsparkl-testbed2d/Cargo.toml | 16 +- crates/wgsparkl-testbed3d/Cargo.toml | 16 +- crates/wgsparkl2d/Cargo.toml | 10 +- crates/wgsparkl3d/Cargo.toml | 10 +- src/grid/grid.rs | 11 +- src/grid/grid.wgsl | 4 + src/grid/sort.rs | 25 +++ src/grid/sort.wgsl | 6 +- src/grid/touch_particle_blocks2d.wgsl | 218 +++++++++++++++++++++++++ src/grid/touch_particle_blocks3d.wgsl | 227 ++++++++++++++++++++++++++ src/lib.rs | 10 +- src/pipeline.rs | 63 +++++-- src/solver/grid_update.wgsl | 1 - src_testbed/hot_reload.rs | 19 +++ src_testbed/instancing2d.rs | 11 +- src_testbed/instancing3d.rs | 11 +- src_testbed/lib.rs | 27 ++- src_testbed/startup.rs | 30 ++-- 19 files changed, 643 insertions(+), 111 deletions(-) create mode 100644 src/grid/touch_particle_blocks2d.wgsl create mode 100644 src/grid/touch_particle_blocks3d.wgsl create mode 100644 src_testbed/hot_reload.rs diff --git a/Cargo.toml b/Cargo.toml index 9f910dd..d3f288c 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -5,11 +5,11 @@ resolver = "2" [workspace.dependencies] nalgebra = { version = "0.33", features = ["convert-bytemuck"] } -wgpu = { version = "22.1", features = ["naga-ir"] } +wgpu = { version = "23", features = ["naga-ir"] } bytemuck = { version = "1", features = ["derive"] } anyhow = "1" async-channel = "2" -naga_oil = "0.15" +naga_oil = "0.16" thiserror = "1" encase = { version = "0.10.0", features = ["nalgebra"] } @@ -19,40 +19,7 @@ rust.unexpected_cfgs = { level = "warn", check-cfg = [ ] } [patch.crates-io] -nalgebra = { git = "https://github.com/dimforge/nalgebra", branch = "more-bytemuck" } -bevy_wasm_window_resize = { git = "https://github.com/Vrixyz/bevy_wasm_window_resize", rev = "770a679316ae24772d278360635e086278c70fa2" } -bevy_editor_cam = { git = "https://github.com/Vrixyz/bevy_editor_cam", rev = "4dce484" } -bevy_egui = { git = "https://github.com/Vrixyz/bevy_egui", rev = "9edc10c" } -bevy_app = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_color = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_derive = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_ecs = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_input = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_log = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_math = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_reflect = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_render = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_time = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_text = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_pbr = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_sprite = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_ui = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_transform = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_utils = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_window = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -# Optional -bevy_asset = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_core_pipeline = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_gizmos = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -# Dev -bevy = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -# Upstream update needed -bevy_picking_core = { git = "https://github.com/vrixyz/bevy_mod_picking.git", branch = "bevy_main" } -bevy_mod_picking = { git = "https://github.com/vrixyz/bevy_mod_picking.git", branch = "bevy_main" } -# transitive dependencies -bevy_hierarchy = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_core = { git = "https://github.com/bevyengine/bevy.git", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_eventlistener = { git = "https://github.com/vrixyz/bevy_eventlistener.git", branch = "bevy_main" } +bevy_egui = { git = "https://github.com/Vrixyz/bevy_egui", branch = "bevy_main" } [profile.release] opt-level = 'z' diff --git a/crates/wgsparkl-testbed2d/Cargo.toml b/crates/wgsparkl-testbed2d/Cargo.toml index 2863a6c..966e231 100644 --- a/crates/wgsparkl-testbed2d/Cargo.toml +++ b/crates/wgsparkl-testbed2d/Cargo.toml @@ -23,17 +23,17 @@ naga_oil = { workspace = true } bytemuck = { workspace = true } async-channel = { workspace = true } -wgcore = { path = "../../../wgmath/crates/wgcore" } -wgebra = { path = "../../../wgmath/crates/wgebra" } -wgparry2d = { path = "../../../wgmath/crates/wgparry/crates/wgparry2d" } +wgcore = { version = "0.2", path = "../../../wgmath/crates/wgcore" } +wgebra = { version = "0.2", path = "../../../wgmath/crates/wgebra" } +wgparry2d = { version = "0.2", path = "../../../wgmath/crates/wgparry/crates/wgparry2d" } futures-test = "0.3" serial_test = "3" approx = "0.5" async-std = { version = "1", features = ["attributes"] } -bevy = { version = "0.15.0-dev", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"], git = "https://github.com/bevyengine/bevy", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_wasm_window_resize = "0.4" -bevy_editor_cam = "0.3" -bevy_mod_picking = { version = "0.20", default-features = false } -bevy_egui = { version = "0.29", default-features = false, features = ["default_fonts", "render"] } +bevy = { version = "0.15.0-rc.3", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"] } +#bevy_wasm_window_resize = "0.4" +#bevy_editor_cam = "0.3" +#bevy_mod_picking = { version = "0.20", default-features = false } +bevy_egui = { version = "0.30", default-features = false, features = ["default_fonts", "render"] } wgsparkl2d = { path = "../wgsparkl2d" } diff --git a/crates/wgsparkl-testbed3d/Cargo.toml b/crates/wgsparkl-testbed3d/Cargo.toml index 7a51745..9c36e1f 100644 --- a/crates/wgsparkl-testbed3d/Cargo.toml +++ b/crates/wgsparkl-testbed3d/Cargo.toml @@ -23,17 +23,17 @@ naga_oil = { workspace = true } bytemuck = { workspace = true } async-channel = { workspace = true } -wgcore = { path = "../../../wgmath/crates/wgcore" } -wgebra = { path = "../../../wgmath/crates/wgebra" } -wgparry3d = { path = "../../../wgmath/crates/wgparry/crates/wgparry3d" } +wgcore = { version = "0.2", path = "../../../wgmath/crates/wgcore" } +wgebra = { version = "0.2", path = "../../../wgmath/crates/wgebra" } +wgparry3d = { version = "0.2", path = "../../../wgmath/crates/wgparry/crates/wgparry3d" } futures-test = "0.3" serial_test = "3" approx = "0.5" async-std = { version = "1", features = ["attributes"] } -bevy = { version = "0.15.0-dev", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"], git = "https://github.com/bevyengine/bevy", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } -bevy_editor_cam = "0.3" -bevy_mod_picking = { version = "0.20", default-features = false } -bevy_wasm_window_resize = "0.4" -bevy_egui = { version = "0.29", default-features = false, features = ["default_fonts", "render"] } +bevy = { version = "0.15.0-rc.3", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"] } +#bevy_editor_cam = "0.3" +#bevy_mod_picking = { version = "0.20", default-features = false } +#bevy_wasm_window_resize = "0.4" +bevy_egui = { version = "0.30", default-features = false, features = ["default_fonts", "render"] } wgsparkl3d = { path = "../wgsparkl3d" } diff --git a/crates/wgsparkl2d/Cargo.toml b/crates/wgsparkl2d/Cargo.toml index a53a4f9..2e6859c 100644 --- a/crates/wgsparkl2d/Cargo.toml +++ b/crates/wgsparkl2d/Cargo.toml @@ -23,10 +23,10 @@ naga_oil = { workspace = true } bytemuck = { workspace = true } encase = { workspace = true } -wgcore = { version = "0.1", path = "../../../wgmath/crates/wgcore" } -wgebra = { version = "0.1", path = "../../../wgmath/crates/wgebra" } -wgparry2d = { version = "0.1", path = "../../../wgmath/crates/wgparry/crates/wgparry2d" } -wgrapier2d = { version = "0.1", path = "../../../wgmath/crates/wgrapier/crates/wgrapier2d" } +wgcore = { version = "0.2", path = "../../../wgmath/crates/wgcore", features = ["hot_reloading"] } +wgebra = { version = "0.2", path = "../../../wgmath/crates/wgebra" } +wgparry2d = { version = "0.2", path = "../../../wgmath/crates/wgparry/crates/wgparry2d" } +wgrapier2d = { version = "0.2", path = "../../../wgmath/crates/wgrapier/crates/wgrapier2d" } [dev-dependencies] nalgebra = { version = "0.33", features = ["rand"] } @@ -34,5 +34,5 @@ futures-test = "0.3" serial_test = "3" approx = "0.5" async-std = { version = "1", features = ["attributes"] } -bevy = { version = "0.15.0-dev", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"], git = "https://github.com/bevyengine/bevy", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } +bevy = { version = "0.15.0-dev", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"] } wgsparkl_testbed2d = { path = "../wgsparkl-testbed2d" } diff --git a/crates/wgsparkl3d/Cargo.toml b/crates/wgsparkl3d/Cargo.toml index c71a4dd..221d482 100644 --- a/crates/wgsparkl3d/Cargo.toml +++ b/crates/wgsparkl3d/Cargo.toml @@ -23,10 +23,10 @@ naga_oil = { workspace = true } bytemuck = { workspace = true } encase = { workspace = true } -wgcore = { version = "0.1", path = "../../../wgmath/crates/wgcore" } -wgebra = { version = "0.1", path = "../../../wgmath/crates/wgebra" } -wgparry3d = { version = "0.1", path = "../../../wgmath/crates/wgparry/crates/wgparry3d" } -wgrapier3d = { version = "0.1", path = "../../../wgmath/crates/wgrapier/crates/wgrapier3d" } +wgcore = { version = "0.2", path = "../../../wgmath/crates/wgcore", features = ["hot_reloading"] } +wgebra = { version = "0.2", path = "../../../wgmath/crates/wgebra" } +wgparry3d = { version = "0.2", path = "../../../wgmath/crates/wgparry/crates/wgparry3d" } +wgrapier3d = { version = "0.2", path = "../../../wgmath/crates/wgrapier/crates/wgrapier3d" } [dev-dependencies] nalgebra = { version = "0.33", features = ["rand"] } @@ -34,5 +34,5 @@ futures-test = "0.3" serial_test = "3" approx = "0.5" async-std = { version = "1", features = ["attributes"] } -bevy = { version = "0.15.0-dev", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"], git = "https://github.com/bevyengine/bevy", rev = "9386bd0114c44c9f00a2e9c41db1225aaa78d159" } +bevy = { version = "0.15.0-rc.3", features = ["shader_format_glsl", "shader_format_spirv", "webgpu"] } wgsparkl_testbed3d = { path = "../wgsparkl-testbed3d" } \ No newline at end of file diff --git a/src/grid/grid.rs b/src/grid/grid.rs index 33852be..be9df99 100644 --- a/src/grid/grid.rs +++ b/src/grid/grid.rs @@ -1,6 +1,6 @@ use crate::dim_shader_defs; use crate::grid::prefix_sum::{PrefixSumWorkspace, WgPrefixSum}; -use crate::grid::sort::WgSort; +use crate::grid::sort::{TouchParticleBlocks, WgSort}; use crate::solver::{GpuParticles, WgParams}; use naga_oil::compose::NagaModuleDescriptor; use std::sync::Arc; @@ -26,6 +26,8 @@ impl WgGrid { grid: &GpuGrid, prefix_sum: &mut PrefixSumWorkspace, sort_module: &'a WgSort, + #[cfg(target_os = "macos")] + touch_particle_blocks: &'a TouchParticleBlocks, prefix_sum_module: &'a WgPrefixSum, queue: &mut KernelInvocationQueue<'a>, ) { @@ -45,7 +47,12 @@ impl WgGrid { .bind0([grid.meta.buffer(), grid.hmap_entries.buffer()]) .queue(grid.cpu_meta.hmap_capacity.div_ceil(GRID_WORKGROUP_SIZE)); - KernelInvocationBuilder::new(queue, &sort_module.touch_particle_blocks) + #[cfg(not(target_os = "macos"))] + let touch_particle_blocks = &sort_module.touch_particle_blocks; + #[cfg(target_os = "macos")] + let touch_particle_blocks = &touch_particle_blocks.touch_particle_blocks; + + KernelInvocationBuilder::new(queue, touch_particle_blocks) .bind_at( 0, [ diff --git a/src/grid/grid.wgsl b/src/grid/grid.wgsl index 7518675..1ccd3c7 100644 --- a/src/grid/grid.wgsl +++ b/src/grid/grid.wgsl @@ -112,6 +112,7 @@ struct GridHashMapEntry { value: BlockHeaderId } +#if MACOS == 0 // The hash map ipmelementation is inspired from https://nosferalatu.com/SimpleGPUHashTable.html fn insertion_index(capacity: u32, key: BlockVirtualId) -> u32 { let packed_key = pack_key(key); @@ -157,6 +158,7 @@ fn insertion_index(capacity: u32, key: BlockVirtualId) -> u32 { return NONE; } +#endif fn find_block_header_id(key: BlockVirtualId) -> BlockHeaderId { let packed_key = pack_key(key); @@ -284,6 +286,7 @@ fn blocks_associated_to_block(block: BlockVirtualId) -> array BlockPhysicalId { return BlockPhysicalId(hid.id * NUM_CELL_PER_BLOCK); diff --git a/src/grid/sort.rs b/src/grid/sort.rs index b170feb..0c42a54 100644 --- a/src/grid/sort.rs +++ b/src/grid/sort.rs @@ -12,6 +12,7 @@ use wgpu::ComputePipeline; shader_defs = "dim_shader_defs" )] pub struct WgSort { + #[cfg(not(target_os = "macos"))] pub(crate) touch_particle_blocks: ComputePipeline, pub(crate) update_block_particle_count: ComputePipeline, pub(crate) copy_particles_len_to_scan_value: ComputePipeline, @@ -19,4 +20,28 @@ pub struct WgSort { pub(crate) finalize_particles_sort: ComputePipeline, } +#[cfg(target_os = "macos")] +pub struct TouchParticleBlocks { + pub(crate) touch_particle_blocks: ComputePipeline, +} + +impl TouchParticleBlocks { + pub fn from_device(device: &wgpu::Device) -> Self { + #[cfg(feature = "dim2")] + let src = wgpu::include_wgsl!("touch_particle_blocks2d.wgsl"); + #[cfg(feature = "dim3")] + let src = wgpu::include_wgsl!("touch_particle_blocks3d.wgsl"); + let cs_module = device.create_shader_module(src); + let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: None, + module: &cs_module, + entry_point: Some("touch_particle_blocks"), + compilation_options: Default::default(), + cache: None, + }); + Self { touch_particle_blocks: compute_pipeline } + } +} + wgcore::test_shader_compilation!(WgSort); diff --git a/src/grid/sort.wgsl b/src/grid/sort.wgsl index faa2bc3..d46bb7c 100644 --- a/src/grid/sort.wgsl +++ b/src/grid/sort.wgsl @@ -13,7 +13,10 @@ var sorted_particle_ids: array; @group(1) @binding(3) var particle_node_linked_lists: array; - +// Disable this kernel on macos because of the underlying compareExchangeMap which is +// not working well with naga-oil. This is why we currently have the flattened +// toouch_particle_block2/3d.wgsl shaders as a workaround currently. +#if MACOS == 0 @compute @workgroup_size(Grid::GRID_WORKGROUP_SIZE, 1, 1) fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; @@ -25,6 +28,7 @@ fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3 } } } +#endif // TODO: can this kernel be combined with touch_particle_blocks? @compute @workgroup_size(Grid::GRID_WORKGROUP_SIZE, 1, 1) diff --git a/src/grid/touch_particle_blocks2d.wgsl b/src/grid/touch_particle_blocks2d.wgsl new file mode 100644 index 0000000..9d3c4d2 --- /dev/null +++ b/src/grid/touch_particle_blocks2d.wgsl @@ -0,0 +1,218 @@ +@group(0) @binding(0) +var grid: Grid; // TODO: should be uniform? Currently it can’t due to the mutable num_active_blocks atomic. +@group(0) @binding(1) +var hmap_entries: array; +@group(0) @binding(2) +var active_blocks: array; +@group(0) @binding(8) +var num_collisions: array>; + + +struct NodeLinkedListAtomic { + head: atomic, + len: atomic, +} + +// Non-atomic version of NodeLinkedListAtomic +struct NodeLinkedList { + head: u32, + len: u32, +} + + +const GRID_WORKGROUP_SIZE: u32 = 64; +const G2P_P2G_WORKGROUP_SIZE: u32 = 64; +const NUM_CELL_PER_BLOCK: u32 = 64; // 8 * 8 in 2D and 4 * 4 * 4 in 3D. + +// TODO: upstream this to wgcore? +struct DispatchIndirectArgs { + x: u32, + y: u32, + z: u32, +} + +/* + * Some index types. + */ +struct BlockVirtualId { + id: vec2, +} + +struct BlockHeaderId { + id: u32, +} + +struct BlockPhysicalId { + id: u32, +} + +struct NodePhysicalId { + id: u32, +} + +/* + * + * HashMap for the grid. + * + */ +const NONE: u32 = 0xffffffffu; + +fn pack_key(key: BlockVirtualId) -> u32 { + return (bitcast(key.id.x + 0x00007fff) & 0x0000ffffu) | + ((bitcast(key.id.y + 0x00007fff) & 0x0000ffffu) << 16); +} + +fn hash(packed_key: u32) -> u32 { + // Murmur3 hash function. + var key = packed_key; + key *= 0xcc9e2d51u; + key = (key << 15) | (key >> 17); + key *= 0x1b873593u; + return key; +} + +// IMPORTANT: if this struct is changed (including its layout), be sure to +// modify the GpuGridHashMapEntry struct on the Rust side to ensure +// it has the right size. Otherwise the hashmap will break. +struct GridHashMapEntry { + // Indicates if the entry is free or empty. + state: atomic, + // The key stored on this entry. + key: BlockVirtualId, + // The associated value. + value: BlockHeaderId +} + +// The hash map ipmelementation is inspired from https://nosferalatu.com/SimpleGPUHashTable.html +fn insertion_index(capacity: u32, key: BlockVirtualId) -> u32 { + let packed_key = pack_key(key); + var slot = hash(packed_key) & (capacity - 1u); + var retries = 0u; + + // NOTE: if there is no more room in the hashmap to store the data, we just do nothing. + // It is up to the user to detect the high occupancy, resize the hashmap, and re-run + // the failed insertion. + for (var k = 0u; k < capacity; k++) { + // TODO: would it be more efficient to move the `state` into its own + // vector with only atomics? + let entry = &hmap_entries[slot]; + + var my_retries = 0u; + loop { + let exch = atomicCompareExchangeWeak(&(*entry).state, NONE, packed_key); + if exch.exchanged { + // We found a slot. + (*entry).key = key; + // TODO: remove these atomicMax, it’s just for debugging. + atomicMax(&num_collisions[0], k); + atomicMax(&num_collisions[1], arrayLength(&hmap_entries)); + return slot; + } else if exch.old_value == packed_key { + // The entry already exists. + // TODO: remove these atomicMax, it’s just for debugging. + atomicMax(&num_collisions[0], k); + atomicMax(&num_collisions[1], arrayLength(&hmap_entries)); + return NONE; + } else if exch.old_value != NONE { + // The slot is already taken. + break; + } + // Otherwise we need to loop since we hit a case where the exchange could + // have happened but didn’t due to the weak nature of the operation. + my_retries += 1u; + } + + retries = max(retries, my_retries); + slot = (slot + 1u) % capacity & (capacity - 1u); + } + + return NONE; +} + +/* + * Sparse grid definition. + */ +const NUM_ASSOC_BLOCKS: u32 = 4; +const OFF_BY_ONE: i32 = 1; + +struct ActiveBlockHeader { + virtual_id: BlockVirtualId, // Needed to compute the world-space position of a block. + first_particle: u32, + num_particles: atomic, +} + + +struct Grid { + num_active_blocks: atomic, + cell_width: f32, + // NOTE: the hashmap capacity MUST be a power of 2. + hmap_capacity: u32, + capacity: u32, +} + +struct Node { + /// The first three components contains either the cell’s momentum or its velocity + /// (depending on the context). The fourth component contains the cell’s mass. + momentum_velocity_mass: vec3, +} + +fn block_associated_to_point(pt: vec2) -> BlockVirtualId { + let assoc_cell = round(pt / grid.cell_width) - 1.0; + let assoc_block = floor(assoc_cell / 8.0); + return BlockVirtualId(vec2( + i32(assoc_block.x), + i32(assoc_block.y), + )); +} + +fn blocks_associated_to_point(pt: vec2) -> array { + let main_block = block_associated_to_point(pt); + return blocks_associated_to_block(main_block); +} + +fn blocks_associated_to_block(block: BlockVirtualId) -> array { + return array( + BlockVirtualId(block.id + vec2(0, 0)), + BlockVirtualId(block.id + vec2(0, 1)), + BlockVirtualId(block.id + vec2(1, 0)), + BlockVirtualId(block.id + vec2(1, 1)), + ); +} + +fn mark_block_as_active(block: BlockVirtualId) { + let slot = insertion_index(grid.hmap_capacity, block); + + if slot != NONE { + let block_header_id = atomicAdd(&grid.num_active_blocks, 1u); + let active_block = &active_blocks[block_header_id]; + (*active_block).virtual_id = block; + (*active_block).first_particle = 0u; + (*active_block).num_particles = 0u; + hmap_entries[slot].value = BlockHeaderId(block_header_id); + } +} + +struct SimulationParameters { + dt: f32, + gravity: vec2, +} + +// ~~~~~~~~~~~ Copied from sort.wgsl and particle2/3d.wgsl ~~~~~~~~~~~~~ +struct Position { + pt: vec2, +} + +@group(1) @binding(0) +var particles_pos: array; + +@compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) +fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { + let id = invocation_id.x; + if id < arrayLength(&particles_pos) { + let particle = particles_pos[id]; + var blocks = blocks_associated_to_point(particle.pt); + for (var i = 0u; i < NUM_ASSOC_BLOCKS; i += 1u) { + mark_block_as_active(blocks[i]); + } + } +} \ No newline at end of file diff --git a/src/grid/touch_particle_blocks3d.wgsl b/src/grid/touch_particle_blocks3d.wgsl new file mode 100644 index 0000000..ca6950f --- /dev/null +++ b/src/grid/touch_particle_blocks3d.wgsl @@ -0,0 +1,227 @@ +@group(0) @binding(0) +var grid: Grid; // TODO: should be uniform? Currently it can’t due to the mutable num_active_blocks atomic. +@group(0) @binding(1) +var hmap_entries: array; +@group(0) @binding(2) +var active_blocks: array; +@group(0) @binding(8) +var num_collisions: array>; + + +struct NodeLinkedListAtomic { + head: atomic, + len: atomic, +} + +// Non-atomic version of NodeLinkedListAtomic +struct NodeLinkedList { + head: u32, + len: u32, +} + + +const GRID_WORKGROUP_SIZE: u32 = 64; +const G2P_P2G_WORKGROUP_SIZE: u32 = 64; +const NUM_CELL_PER_BLOCK: u32 = 64; // 8 * 8 in 2D and 4 * 4 * 4 in 3D. + +// TODO: upstream this to wgcore? +struct DispatchIndirectArgs { + x: u32, + y: u32, + z: u32, +} + +/* + * Some index types. + */ +struct BlockVirtualId { + id: vec3, +} + +struct BlockHeaderId { + id: u32, +} + +struct BlockPhysicalId { + id: u32, +} + +struct NodePhysicalId { + id: u32, +} + +/* + * + * HashMap for the grid. + * + */ +const NONE: u32 = 0xffffffffu; + +fn pack_key(key: BlockVirtualId) -> u32 { + // NOTE: we give the X and Z axis one more bit than Y. + // This is assuming Y-up and the fact that we want + // more room on the X-Z plane rather than along the up axis. + return (bitcast(key.id.x + 0x000003ff) & 0x000007ffu) | + ((bitcast(key.id.y + 0x000001ff) & 0x000003ffu) << 11) | + ((bitcast(key.id.z + 0x000003ff) & 0x000007ffu) << 21); +} + +fn hash(packed_key: u32) -> u32 { + // Murmur3 hash function. + var key = packed_key; + key *= 0xcc9e2d51u; + key = (key << 15) | (key >> 17); + key *= 0x1b873593u; + return key; +} + +// IMPORTANT: if this struct is changed (including its layout), be sure to +// modify the GpuGridHashMapEntry struct on the Rust side to ensure +// it has the right size. Otherwise the hashmap will break. +struct GridHashMapEntry { + // Indicates if the entry is free or empty. + state: atomic, + // The key stored on this entry. + key: BlockVirtualId, + // The associated value. + value: BlockHeaderId +} + +// The hash map ipmelementation is inspired from https://nosferalatu.com/SimpleGPUHashTable.html +fn insertion_index(capacity: u32, key: BlockVirtualId) -> u32 { + let packed_key = pack_key(key); + var slot = hash(packed_key) & (capacity - 1u); + var retries = 0u; + + // NOTE: if there is no more room in the hashmap to store the data, we just do nothing. + // It is up to the user to detect the high occupancy, resize the hashmap, and re-run + // the failed insertion. + for (var k = 0u; k < capacity; k++) { + // TODO: would it be more efficient to move the `state` into its own + // vector with only atomics? + let entry = &hmap_entries[slot]; + + var my_retries = 0u; + loop { + let exch = atomicCompareExchangeWeak(&(*entry).state, NONE, packed_key); + if exch.exchanged { + // We found a slot. + (*entry).key = key; + // TODO: remove these atomicMax, it’s just for debugging. + atomicMax(&num_collisions[0], k); + atomicMax(&num_collisions[1], arrayLength(&hmap_entries)); + return slot; + } else if exch.old_value == packed_key { + // The entry already exists. + // TODO: remove these atomicMax, it’s just for debugging. + atomicMax(&num_collisions[0], k); + atomicMax(&num_collisions[1], arrayLength(&hmap_entries)); + return NONE; + } else if exch.old_value != NONE { + // The slot is already taken. + break; + } + // Otherwise we need to loop since we hit a case where the exchange could + // have happened but didn’t due to the weak nature of the operation. + my_retries += 1u; + } + + retries = max(retries, my_retries); + slot = (slot + 1u) % capacity & (capacity - 1u); + } + + return NONE; +} + +/* + * Sparse grid definition. + */ +const NUM_ASSOC_BLOCKS: u32 = 8; +const OFF_BY_ONE: i32 = 1; + +struct ActiveBlockHeader { + virtual_id: BlockVirtualId, // Needed to compute the world-space position of a block. + first_particle: u32, + num_particles: atomic, +} + + +struct Grid { + num_active_blocks: atomic, + cell_width: f32, + // NOTE: the hashmap capacity MUST be a power of 2. + hmap_capacity: u32, + capacity: u32, +} + +struct Node { + /// The first three components contains either the cell’s momentum or its velocity + /// (depending on the context). The fourth component contains the cell’s mass. + momentum_velocity_mass: vec4, +} + +fn block_associated_to_point(pt: vec3) -> BlockVirtualId { + let assoc_cell = round(pt / grid.cell_width) - 1.0; + let assoc_block = floor(assoc_cell / 4.0); + return BlockVirtualId(vec3( + i32(assoc_block.x), + i32(assoc_block.y), + i32(assoc_block.z), + )); +} + +fn blocks_associated_to_point(pt: vec3) -> array { + let main_block = block_associated_to_point(pt); + return blocks_associated_to_block(main_block); +} + +fn blocks_associated_to_block(block: BlockVirtualId) -> array { + return array( + BlockVirtualId(block.id + vec3(0, 0, 0)), + BlockVirtualId(block.id + vec3(0, 0, 1)), + BlockVirtualId(block.id + vec3(0, 1, 0)), + BlockVirtualId(block.id + vec3(0, 1, 1)), + BlockVirtualId(block.id + vec3(1, 0, 0)), + BlockVirtualId(block.id + vec3(1, 0, 1)), + BlockVirtualId(block.id + vec3(1, 1, 0)), + BlockVirtualId(block.id + vec3(1, 1, 1)), + ); +} + +fn mark_block_as_active(block: BlockVirtualId) { + let slot = insertion_index(grid.hmap_capacity, block); + + if slot != NONE { + let block_header_id = atomicAdd(&grid.num_active_blocks, 1u); + let active_block = &active_blocks[block_header_id]; + (*active_block).virtual_id = block; + (*active_block).first_particle = 0u; + (*active_block).num_particles = 0u; + hmap_entries[slot].value = BlockHeaderId(block_header_id); + } +} + +struct SimulationParameters { + dt: f32, + gravity: vec3, +} + +// ~~~~~~~~~~~ Copied from sort.wgsl and particle2/3d.wgsl ~~~~~~~~~~~~~ +struct Position { + pt: vec3, +} + +@group(1) @binding(0) +var particles_pos: array; + +@compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) +fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { + let id = invocation_id.x; + if id < arrayLength(&particles_pos) { + let particle = particles_pos[id]; + var blocks = blocks_associated_to_point(particle.pt); + for (var i = 0u; i < NUM_ASSOC_BLOCKS; i += 1u) { + mark_block_as_active(blocks[i]); + } + } +} \ No newline at end of file diff --git a/src/lib.rs b/src/lib.rs index 11a24d9..478050f 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -13,4 +13,12 @@ pub mod models; pub mod pipeline; pub mod solver; -pub(crate) use wgparry::{dim_shader_defs, substitute_aliases}; +pub(crate) fn dim_shader_defs() -> HashMap { + let mut result = wgparry::dim_shader_defs(); + result.insert("MACOS".to_string(), ShaderDefValue::Int(if cfg!(target_os = "macos") { 1 } else { 0 })); + result +} + +use std::collections::HashMap; +use naga_oil::compose::ShaderDefValue; +pub(crate) use wgparry::{substitute_aliases}; diff --git a/src/pipeline.rs b/src/pipeline.rs index 232b3ca..2556eb2 100644 --- a/src/pipeline.rs +++ b/src/pipeline.rs @@ -1,11 +1,13 @@ use crate::grid::grid::{GpuGrid, WgGrid}; use crate::grid::prefix_sum::{PrefixSumWorkspace, WgPrefixSum}; -use crate::grid::sort::WgSort; +use crate::grid::sort::{TouchParticleBlocks, WgSort}; use crate::models::GpuModels; use crate::solver::{ GpuParticles, GpuSimulationParams, Particle, SimulationParams, WgG2P, WgGridUpdate, WgP2G, WgParticleUpdate, }; +use naga_oil::compose::ComposerError; +use wgcore::hot_reloading::HotReloadState; use wgcore::kernel::KernelInvocationQueue; use wgcore::Shader; use wgpu::Device; @@ -15,6 +17,8 @@ pub struct MpmPipeline { grid: WgGrid, prefix_sum: WgPrefixSum, sort: WgSort, + #[cfg(target_os = "macos")] + touch_particle_blocks: TouchParticleBlocks, p2g: WgP2G, grid_update: WgGridUpdate, particles_update: WgParticleUpdate, @@ -22,6 +26,37 @@ pub struct MpmPipeline { integrate_bodies: WgIntegrate, } +impl MpmPipeline { + pub fn init_hot_reloading(&self, state: &mut HotReloadState) { + WgGrid::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgPrefixSum::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgSort::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgP2G::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgGridUpdate::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgParticleUpdate::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgG2P::watch_sources(state).unwrap(); // TODO: don’t unwrap + WgIntegrate::watch_sources(state).unwrap(); // TODO: don’t unwrap + } + + pub fn reload_if_changed( + &mut self, + device: &Device, + state: &HotReloadState, + ) -> Result { + let mut changed = false; + changed = self.grid.reload_if_changed(device, state)? || changed; + changed = self.prefix_sum.reload_if_changed(device, state)? || changed; + changed = self.sort.reload_if_changed(device, state)? || changed; + changed = self.p2g.reload_if_changed(device, state)? || changed; + changed = self.grid_update.reload_if_changed(device, state)? || changed; + changed = self.particles_update.reload_if_changed(device, state)? || changed; + changed = self.g2p.reload_if_changed(device, state)? || changed; + changed = self.integrate_bodies.reload_if_changed(device, state)? || changed; + + Ok(changed) + } +} + pub struct MpmData { pub sim_params: GpuSimulationParams, pub grid: GpuGrid, @@ -59,17 +94,19 @@ impl MpmData { } impl MpmPipeline { - pub fn new(device: &Device) -> Self { - Self { - grid: WgGrid::from_device(device), - prefix_sum: WgPrefixSum::from_device(device), - sort: WgSort::from_device(device), - p2g: WgP2G::from_device(device), - grid_update: WgGridUpdate::from_device(device), - particles_update: WgParticleUpdate::from_device(device), - g2p: WgG2P::from_device(device), - integrate_bodies: WgIntegrate::from_device(device), - } + pub fn new(device: &Device) -> Result { + Ok(Self { + grid: WgGrid::from_device(device)?, + prefix_sum: WgPrefixSum::from_device(device)?, + sort: WgSort::from_device(device)?, + p2g: WgP2G::from_device(device)?, + grid_update: WgGridUpdate::from_device(device)?, + particles_update: WgParticleUpdate::from_device(device)?, + g2p: WgG2P::from_device(device)?, + integrate_bodies: WgIntegrate::from_device(device)?, + #[cfg(target_os = "macos")] + touch_particle_blocks: TouchParticleBlocks::from_device(device), + }) } pub fn queue_step<'a>( @@ -85,6 +122,8 @@ impl MpmPipeline { &data.grid, &mut data.prefix_sum, &self.sort, + #[cfg(target_os = "macos")] + &self.touch_particle_blocks, &self.prefix_sum, queue, ); diff --git a/src/solver/grid_update.wgsl b/src/solver/grid_update.wgsl index 110aefd..1d98b03 100644 --- a/src/solver/grid_update.wgsl +++ b/src/solver/grid_update.wgsl @@ -47,7 +47,6 @@ fn update_single_cell(cell_pos: vec2, momentum_velocity_mass: vec3) -> let mass = momentum_velocity_mass.z; let inv_mass = select(0.0, 1.0 / mass, mass > 0.0); var velocity = (momentum_velocity_mass.xy + mass * Grid::sim_params.gravity * Grid::sim_params.dt) * inv_mass; - // Clamp the velocity so it doesn’t exceed 1 grid cell in one step. let vel_limit = vec2(Grid::grid.cell_width / Grid::sim_params.dt); velocity = clamp(velocity, -vel_limit, vel_limit); diff --git a/src_testbed/hot_reload.rs b/src_testbed/hot_reload.rs new file mode 100644 index 0000000..806bc75 --- /dev/null +++ b/src_testbed/hot_reload.rs @@ -0,0 +1,19 @@ +use crate::AppState; +use bevy::prelude::*; +use bevy::render::renderer::RenderDevice; + +pub fn handle_hot_reloading(render_device: Res, mut state: ResMut) { + let device = render_device.wgpu_device(); + let state = &mut *state; + state.hot_reload.update_changes(); + match state.pipeline.reload_if_changed(device, &state.hot_reload) { + Err(e) => { + println!("Failed to hot reload the MPM pipeline: {e}"); + } + Ok(changed) => { + if changed { + println!("Hot reloaded MPM pipeline successfully."); + } + } + } +} diff --git a/src_testbed/instancing2d.rs b/src_testbed/instancing2d.rs index 0e17119..b07f0e3 100644 --- a/src_testbed/instancing2d.rs +++ b/src_testbed/instancing2d.rs @@ -27,6 +27,7 @@ use bevy::{ }; use bytemuck::{Pod, Zeroable}; use std::sync::Arc; +use bevy::render::sync_world::MainEntity; pub const INSTANCING_SHADER_HANDLE: Handle = Handle::weak_from_u128(3222377299100772450); @@ -79,7 +80,7 @@ fn queue_custom( pipeline_cache: Res, meshes: Res>, render_mesh_instances: Res, - material_meshes: Query>, + material_meshes: Query<(Entity, &MainEntity), With>, mut transparent_render_phases: ResMut>, mut views: Query<(Entity, &ExtractedView, &Msaa)>, ) { @@ -94,8 +95,8 @@ fn queue_custom( let view_key = msaa_key | MeshPipelineKey::from_hdr(view.hdr); let rangefinder = view.rangefinder3d(); - for entity in &material_meshes { - let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(entity) else { + for (entity, main_entity) in &material_meshes { + let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(*main_entity) else { continue; }; let Some(mesh) = meshes.get(mesh_instance.mesh_asset_id) else { @@ -107,7 +108,7 @@ fn queue_custom( .specialize(&pipeline_cache, &custom_pipeline, key, &mesh.layout) .unwrap(); transparent_phase.add(Transparent3d { - entity, + entity: (entity, *main_entity), pipeline, draw_function: draw_custom, distance: rangefinder.distance_translation(&mesh_instance.translation), @@ -223,7 +224,7 @@ impl RenderCommand

for DrawMeshInstanced { // A borrow check workaround. let mesh_allocator = mesh_allocator.into_inner(); - let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(item.entity()) + let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(item.main_entity()) else { return RenderCommandResult::Skip; }; diff --git a/src_testbed/instancing3d.rs b/src_testbed/instancing3d.rs index 10ca858..74aec57 100644 --- a/src_testbed/instancing3d.rs +++ b/src_testbed/instancing3d.rs @@ -27,6 +27,7 @@ use bevy::{ }; use bytemuck::{Pod, Zeroable}; use std::sync::Arc; +use bevy::render::sync_world::MainEntity; pub const INSTANCING_SHADER_HANDLE: Handle = Handle::weak_from_u128(3222377299100772450); @@ -79,7 +80,7 @@ fn queue_custom( pipeline_cache: Res, meshes: Res>, render_mesh_instances: Res, - material_meshes: Query>, + material_meshes: Query<(Entity, &MainEntity), With>, mut transparent_render_phases: ResMut>, views: Query<(Entity, &ExtractedView, &Msaa)>, ) { @@ -93,8 +94,8 @@ fn queue_custom( let msaa_key = MeshPipelineKey::from_msaa_samples(msaa.samples()); let view_key = msaa_key | MeshPipelineKey::from_hdr(view.hdr); let rangefinder = view.rangefinder3d(); - for entity in &material_meshes { - let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(entity) else { + for (entity, main_entity) in &material_meshes { + let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(*main_entity) else { continue; }; let Some(mesh) = meshes.get(mesh_instance.mesh_asset_id) else { @@ -106,7 +107,7 @@ fn queue_custom( .specialize(&pipeline_cache, &custom_pipeline, key, &mesh.layout) .unwrap(); transparent_phase.add(Transparent3d { - entity, + entity: (entity, *main_entity), pipeline, draw_function: draw_custom, distance: rangefinder.distance_translation(&mesh_instance.translation), @@ -222,7 +223,7 @@ impl RenderCommand

for DrawMeshInstanced { // A borrow check workaround. let mesh_allocator = mesh_allocator.into_inner(); - let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(item.entity()) + let Some(mesh_instance) = render_mesh_instances.render_mesh_queue_data(item.main_entity()) else { return RenderCommandResult::Skip; }; diff --git a/src_testbed/lib.rs b/src_testbed/lib.rs index 260395d..6e2454f 100644 --- a/src_testbed/lib.rs +++ b/src_testbed/lib.rs @@ -13,6 +13,7 @@ pub mod instancing2d; #[cfg(feature = "dim3")] pub mod instancing3d; +mod hot_reload; pub mod prep_vertex_buffer; pub mod startup; pub mod step; @@ -21,11 +22,11 @@ pub mod ui; use bevy::asset::load_internal_asset; use bevy::ecs::system::SystemId; use bevy::prelude::*; -use bevy_editor_cam::prelude::DefaultEditorCamPlugins; -use bevy_wasm_window_resize::WindowResizePlugin; - +// use bevy_editor_cam::prelude::DefaultEditorCamPlugins; +// use bevy_wasm_window_resize::WindowResizePlugin; use instancing::INSTANCING_SHADER_HANDLE; use prep_vertex_buffer::{GpuRenderConfig, RenderConfig, WgPrepVertexBuffer}; +use wgcore::hot_reloading::HotReloadState; use wgcore::timestamps::GpuTimestamps; use wgsparkl::{ pipeline::{MpmData, MpmPipeline}, @@ -34,16 +35,23 @@ use wgsparkl::{ pub fn init_testbed(app: &mut App) { app.add_plugins(DefaultPlugins) - .add_plugins(WindowResizePlugin) - .add_plugins(( - bevy_mod_picking::DefaultPickingPlugins, - DefaultEditorCamPlugins, - )) + // .add_plugins(WindowResizePlugin) + // .add_plugins(( + // bevy_mod_picking::DefaultPickingPlugins, + // DefaultEditorCamPlugins, + // )) .add_plugins(instancing::ParticlesMaterialPlugin) .add_plugins(bevy_egui::EguiPlugin) .init_resource::() .add_systems(Startup, startup::setup_app) - .add_systems(Update, (ui::update_ui, step::step_simulation)); + .add_systems( + Update, + ( + ui::update_ui, + step::step_simulation, + hot_reload::handle_hot_reloading, + ), + ); #[cfg(feature = "dim2")] load_internal_asset!( @@ -72,6 +80,7 @@ pub struct AppState { pub gravity_factor: f32, pub restarting: bool, pub selected_scene: usize, + pub hot_reload: HotReloadState, } #[derive(Resource)] diff --git a/src_testbed/startup.rs b/src_testbed/startup.rs index e76f4a4..b22364a 100644 --- a/src_testbed/startup.rs +++ b/src_testbed/startup.rs @@ -11,8 +11,9 @@ use bevy::prelude::*; use bevy::render::render_resource::BufferUsages; use bevy::render::renderer::RenderDevice; use bevy::render::view::NoFrustumCulling; -use bevy_editor_cam::prelude::{EditorCam, EnabledMotion}; +// use bevy_editor_cam::prelude::{EditorCam, EnabledMotion}; use std::sync::Arc; +use wgcore::hot_reloading::HotReloadState; use wgcore::tensor::GpuVector; use wgcore::timestamps::GpuTimestamps; use wgpu::Features; @@ -24,7 +25,10 @@ pub fn setup_app(mut commands: Commands, device: Res) { let render_config = RenderConfig::new(RenderMode::Velocity); let gpu_render_config = GpuRenderConfig::new(device.wgpu_device(), render_config); let prep_vertex_buffer = WgPrepVertexBuffer::new(device.wgpu_device()); - let pipeline = MpmPipeline::new(device.wgpu_device()); + + let mut hot_reload = HotReloadState::new().unwrap(); + let pipeline = MpmPipeline::new(device.wgpu_device()).unwrap(); + pipeline.init_hot_reloading(&mut hot_reload); commands.insert_resource(AppState { render_config, @@ -36,6 +40,7 @@ pub fn setup_app(mut commands: Commands, device: Res) { gravity_factor: 1.0, restarting: false, selected_scene: 0, + hot_reload, }); let (snd, rcv) = async_channel::unbounded(); @@ -68,14 +73,14 @@ pub fn setup_app(mut commands: Commands, device: Res) { // }), ..default() }, - EditorCam { - enabled_motion: EnabledMotion { - orbit: false, - ..Default::default() - }, - last_anchor_depth: -99.0, - ..Default::default() - }, + // EditorCam { + // enabled_motion: EnabledMotion { + // orbit: false, + // ..Default::default() + // }, + // last_anchor_depth: -99.0, + // ..Default::default() + // }, )); } @@ -86,8 +91,7 @@ pub fn setup_app(mut commands: Commands, device: Res) { transform: Transform::from_translation(Vec3::new(0.0, 1.5, 5.0)), ..default() }, - EditorCam::default(), - // PanOrbitCamera::default(), + // EditorCam::default(), )); } } @@ -144,7 +148,7 @@ pub fn setup_graphics( let num_instances = instances.len(); commands.spawn(( - cube, + Mesh3d(cube), SpatialBundle::INHERITED_IDENTITY, InstanceMaterialData { data: instances,