diff --git a/src/scene/entity.rs b/src/scene/entity.rs index 898508a12..2d33a8279 100644 --- a/src/scene/entity.rs +++ b/src/scene/entity.rs @@ -557,10 +557,11 @@ impl Scene { }) }) .collect(); - for (handle, mesh, top_level) in built { + for (handle, mut mesh, top_level) in built { if top_level { self.meshes.insert(handle, mesh); } else { + mesh.prepare_instance_source(handle); self.block_meshes.insert(handle, mesh); } } @@ -2232,10 +2233,11 @@ impl Scene { }) }) .collect(); - for (handle, m, top_level) in built { + for (handle, mut m, top_level) in built { if top_level { self.meshes.insert(handle, m); } else { + m.prepare_instance_source(handle); self.block_meshes.insert(handle, m); } } diff --git a/src/scene/mod.rs b/src/scene/mod.rs index 9cb496f2c..daa42653e 100644 --- a/src/scene/mod.rs +++ b/src/scene/mod.rs @@ -651,10 +651,11 @@ fn build_derived_caches_impl( .collect(); let mut meshes: HashMap = HashMap::default(); let mut block_meshes: HashMap = HashMap::default(); - for (handle, m, top_level) in built { + for (handle, mut m, top_level) in built { if top_level { meshes.insert(handle, m); } else { + m.prepare_instance_source(handle); block_meshes.insert(handle, m); } } @@ -1037,6 +1038,7 @@ fn transform_block_mesh_lod_set( ) -> MeshLodSet { use acadrust::types::Vector3; let mut out = set.clone(); + out.instance_transform = Some(*xform); let transform_direction = |direction: [f32; 3]| { let transformed = xform.apply_rotation(Vector3::new( direction[0] as f64, diff --git a/src/scene/model/mesh_model.rs b/src/scene/model/mesh_model.rs index 14768f837..c59197ca3 100644 --- a/src/scene/model/mesh_model.rs +++ b/src/scene/model/mesh_model.rs @@ -167,6 +167,19 @@ pub struct MeshLodSet { /// triangle). `verts` carry only the high half of the double-single /// position, so the bound is f32-precise — fine for a conservative cull. pub z_aabb: [f32; 2], + /// Immutable block-local geometry shared by every INSERT instance of this + /// block entity. Top-level meshes leave this empty. + pub instance_source: Option>, + /// Accumulated block-local → world transform for this rendered instance. + pub instance_transform: Option, +} + +#[derive(Clone, Debug)] +pub struct MeshInstanceSource { + pub handle: acadrust::Handle, + pub lods: Vec, + pub edge_verts: Vec<[f32; 3]>, + pub edge_verts_low: Vec<[f32; 3]>, } /// 3D bounds of every LOD's vertices: `([min_x, min_y, max_x, max_y], [min_z, max_z])`. @@ -279,6 +292,8 @@ impl MeshLodSet { metrics, world_aabb, z_aabb, + instance_source: None, + instance_transform: None, } } @@ -297,4 +312,14 @@ impl MeshLodSet { self.world_aabb = xy; self.z_aabb = z; } + + pub fn prepare_instance_source(&mut self, handle: acadrust::Handle) { + self.instance_source = Some(std::sync::Arc::new(MeshInstanceSource { + handle, + lods: self.lods.clone(), + edge_verts: self.edge_verts.clone(), + edge_verts_low: self.edge_verts_low.clone(), + })); + self.instance_transform = None; + } } diff --git a/src/scene/pipeline/mesh_gpu.rs b/src/scene/pipeline/mesh_gpu.rs index 313c2a8a8..ce4134aba 100644 --- a/src/scene/pipeline/mesh_gpu.rs +++ b/src/scene/pipeline/mesh_gpu.rs @@ -131,55 +131,85 @@ impl MeshVertex { } } -// ── GPU handle ──────────────────────────────────────────────────────────── - -pub struct MeshGpu { - pub vertex_buffer: wgpu::Buffer, - pub index_buffer: wgpu::Buffer, - pub index_count: u32, - /// Line-list index buffer: every triangle `(a, b, c)` from the - /// solid index buffer is expanded into three segments - /// `(a,b)(b,c)(c,a)`. Used by the wireframe-mode render path so 3D - /// solids draw as their triangle edges without needing the - /// `POLYGON_MODE_LINE` device feature. - #[allow(dead_code)] // only the highlight overlay builds MeshGpu now (fill only) - pub wire_index_buffer: wgpu::Buffer, - #[allow(dead_code)] - pub wire_index_count: u32, +#[repr(C)] +#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +pub struct MeshInstanceGpu { + pub model_row_0: [f32; 4], + pub model_row_1: [f32; 4], + pub model_row_2: [f32; 4], + pub translation_low: [f32; 4], + pub normal_row_0: [f32; 4], + pub normal_row_1: [f32; 4], + pub normal_row_2: [f32; 4], } -/// GPU-side bundle of MeshLodSet — one MeshGpu per available LOD plus -/// the world-XY AABB needed to pick a level per frame. -pub struct MeshLodGpu { - pub lods: Vec, - pub world_aabb: [f32; 4], -} +impl MeshInstanceGpu { + fn identity() -> Self { + Self { + model_row_0: [1.0, 0.0, 0.0, 0.0], + model_row_1: [0.0, 1.0, 0.0, 0.0], + model_row_2: [0.0, 0.0, 1.0, 0.0], + translation_low: [0.0; 4], + normal_row_0: [1.0, 0.0, 0.0, 0.0], + normal_row_1: [0.0, 1.0, 0.0, 0.0], + normal_row_2: [0.0, 0.0, 1.0, 0.0], + } + } -/// How a solid mesh is highlighted this frame. -#[derive(Copy, Clone, PartialEq, Eq)] -pub enum Highlight { - #[allow(dead_code)] // the highlight overlay only builds Selected / Hover - None, - /// Hovered — light orange wash. - Hover, - /// Selected — stronger blue wash. - Selected, -} - -impl Highlight { - /// Blend colour and mix factor, or `None` when the mesh keeps its colour. - fn tint(self) -> Option<([f32; 4], f32)> { - match self { - Highlight::None => None, - Highlight::Hover => Some(([0.95, 0.55, 0.10, 1.0], 0.35)), - Highlight::Selected => Some(([0.15, 0.55, 1.0, 1.0], 0.60)), + fn from_transform(transform: acadrust::types::Transform) -> Self { + let m = transform.matrix.m; + let translation = [m[0][3], m[1][3], m[2][3]]; + let translation_high = [ + translation[0] as f32, + translation[1] as f32, + translation[2] as f32, + ]; + let linear = glam::DMat3::from_cols_array(&[ + m[0][0], m[1][0], m[2][0], + m[0][1], m[1][1], m[2][1], + m[0][2], m[1][2], m[2][2], + ]); + let normal = if linear.determinant().abs() > 1e-18 { + linear.inverse().transpose() + } else { + glam::DMat3::IDENTITY + }; + let n = normal.to_cols_array(); + Self { + model_row_0: [ + m[0][0] as f32, + m[0][1] as f32, + m[0][2] as f32, + translation_high[0], + ], + model_row_1: [ + m[1][0] as f32, + m[1][1] as f32, + m[1][2] as f32, + translation_high[1], + ], + model_row_2: [ + m[2][0] as f32, + m[2][1] as f32, + m[2][2] as f32, + translation_high[2], + ], + translation_low: [ + (translation[0] - translation_high[0] as f64) as f32, + (translation[1] - translation_high[1] as f64) as f32, + (translation[2] - translation_high[2] as f64) as f32, + 0.0, + ], + normal_row_0: [n[0] as f32, n[3] as f32, n[6] as f32, 0.0], + normal_row_1: [n[1] as f32, n[4] as f32, n[7] as f32, 0.0], + normal_row_2: [n[2] as f32, n[5] as f32, n[8] as f32, 0.0], } } } // ── Batched mesh buffers ────────────────────────────────────────────────── // -// One MeshGpu per solid means one vertex/index bind + draw call per solid — +// One GPU allocation per solid means one vertex/index bind + draw call per solid — // ~10k draw calls a frame on a heavy 3D model, which strangles the GPU front // end. The batch concatenates every solid's LOD0 geometry into a handful of // large buffers (split only to stay under the 256 MB per-buffer cap), so the @@ -205,10 +235,26 @@ pub struct MeshBatchChunk { /// buffer (pairs of endpoints), drawn non-indexed. Empty for plain meshes. pub edge_vertex_buffer: wgpu::Buffer, pub edge_vertex_count: u32, + pub instance_buffer: wgpu::Buffer, + pub instance_count: u32, + pub highlight_ranges: Vec, + pub handles: rustc_hash::FxHashSet, + pub world_aabb: [f32; 6], + pub visible: bool, pub material: Option, pub material_bind_group: Option, } +#[derive(Clone, Copy)] +pub struct MeshBatchRange { + pub handle: acadrust::Handle, + pub index_start: u32, + pub index_count: u32, + pub transparent: bool, + pub instance_start: u32, + pub instance_count: u32, +} + fn make_chunk( device: &wgpu::Device, verts: &[MeshVertex], @@ -216,6 +262,10 @@ fn make_chunk( transp_indices: &[u32], wire_indices: &[u32], edge_verts: &[MeshVertex], + highlight_ranges: &[MeshBatchRange], + instances: &[MeshInstanceGpu], + handles: &rustc_hash::FxHashSet, + bounds_override: Option<[f32; 6]>, material: Option<&crate::scene::model::material_model::MeshMaterial>, ) -> MeshBatchChunk { // `create_buffer_init` with an empty slice yields a zero-sized buffer that @@ -257,6 +307,38 @@ fn make_chunk( usage: wgpu::BufferUsages::VERTEX, }) }; + let identity = [MeshInstanceGpu::identity()]; + let instance_data = if instances.is_empty() { + &identity[..] + } else { + instances + }; + let instance_usage = + if device.limits().max_storage_buffers_per_shader_stage > 0 { + wgpu::BufferUsages::STORAGE + } else { + wgpu::BufferUsages::UNIFORM + }; + let mut computed_aabb = [ + f32::INFINITY, + f32::INFINITY, + f32::INFINITY, + f32::NEG_INFINITY, + f32::NEG_INFINITY, + f32::NEG_INFINITY, + ]; + for vertex in verts.iter().chain(edge_verts) { + let point = [ + vertex.position[0] + vertex.position_low[0], + vertex.position[1] + vertex.position_low[1], + vertex.position[2] + vertex.position_low[2], + ]; + for axis in 0..3 { + computed_aabb[axis] = computed_aabb[axis].min(point[axis]); + computed_aabb[axis + 3] = computed_aabb[axis + 3].max(point[axis]); + } + } + let world_aabb = bounds_override.unwrap_or(computed_aabb); MeshBatchChunk { vertex_buffer: mk_vertex(verts, "mesh.batch.vbuf"), index_buffer: mk_index(indices, "mesh.batch.ibuf"), @@ -267,6 +349,16 @@ fn make_chunk( wire_index_count: wire_indices.len() as u32, edge_vertex_buffer: mk_vertex(edge_verts, "mesh.batch.edge_vbuf"), edge_vertex_count: edge_verts.len() as u32, + instance_buffer: device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("mesh.batch.instances"), + contents: bytemuck::cast_slice(instance_data), + usage: instance_usage, + }), + instance_count: instance_data.len() as u32, + highlight_ranges: highlight_ranges.to_vec(), + handles: handles.clone(), + world_aabb, + visible: true, material: material.cloned(), material_bind_group: None, } @@ -336,6 +428,7 @@ pub fn create_material_bind_group( queue: &wgpu::Queue, layout: &wgpu::BindGroupLayout, material: Option<&crate::scene::model::material_model::MeshMaterial>, + instance_buffer: Option<&wgpu::Buffer>, ) -> wgpu::BindGroup { let diffuse = material.and_then(|material| material.diffuse_map.image.as_deref()); let specular = material.and_then(|material| material.specular_map.image.as_deref()); @@ -440,6 +533,25 @@ pub fn create_material_bind_group( contents: bytemuck::bytes_of(¶ms), usage: wgpu::BufferUsages::UNIFORM, }); + let fallback_instances; + let instance_buffer = match instance_buffer { + Some(buffer) => buffer, + None => { + let usage = + if device.limits().max_storage_buffers_per_shader_stage > 0 { + wgpu::BufferUsages::STORAGE + } else { + wgpu::BufferUsages::UNIFORM + }; + fallback_instances = + device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("mesh.instances.identity"), + contents: bytemuck::bytes_of(&MeshInstanceGpu::identity()), + usage, + }); + &fallback_instances + } + }; device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("mesh.material.bind_group"), layout, @@ -504,6 +616,10 @@ pub fn create_material_bind_group( binding: 14, resource: params_buffer.as_entire_binding(), }, + wgpu::BindGroupEntry { + binding: 15, + resource: instance_buffer.as_entire_binding(), + }, ], }) } @@ -556,14 +672,233 @@ fn material_key( ]) } +fn morton_axis(mut value: u32) -> u64 { + value &= 0x3ff; + let mut result = 0u64; + for bit in 0..10 { + result |= ((value >> bit) as u64 & 1) << (bit * 3); + } + result +} + +fn mesh_spatial_key(set: &MeshLodSet, bounds: [f32; 6]) -> u64 { + let center = [ + (set.world_aabb[0] + set.world_aabb[2]) * 0.5, + (set.world_aabb[1] + set.world_aabb[3]) * 0.5, + (set.z_aabb[0] + set.z_aabb[1]) * 0.5, + ]; + let quantize = |value: f32, min: f32, max: f32| { + let span = max - min; + if !value.is_finite() || !span.is_finite() || span <= f32::EPSILON { + 0 + } else { + (((value - min) / span).clamp(0.0, 1.0) * 1023.0).round() as u32 + } + }; + let x = quantize(center[0], bounds[0], bounds[3]); + let y = quantize(center[1], bounds[1], bounds[4]); + let z = quantize(center[2], bounds[2], bounds[5]); + morton_axis(x) | (morton_axis(y) << 1) | (morton_axis(z) << 2) +} + struct MeshBatchPart<'a> { set: &'a MeshLodSet, mesh: &'a MeshModel, + display_mesh: &'a MeshModel, material: Option<&'a crate::scene::model::material_model::MeshMaterial>, color: [f32; 4], indices: Vec, include_faces: bool, include_edges: bool, + part_slot: u32, +} + +#[derive(Clone, Copy, PartialEq, Eq, PartialOrd, Ord)] +struct InstanceGroupKey { + material: MaterialBatchKey, + source: u64, + part_slot: u32, + index_count: usize, + index_hash: u64, + include_faces: bool, + include_edges: bool, +} + +fn index_hash(indices: &[u32]) -> u64 { + use std::hash::{Hash, Hasher}; + let mut hasher = rustc_hash::FxHasher::default(); + indices.hash(&mut hasher); + hasher.finish() +} + +fn build_instanced_chunk( + device: &wgpu::Device, + parts: &[MeshBatchPart<'_>], +) -> Option<(MeshBatchChunk, u64)> { + let first = parts.first()?; + let source = first.set.instance_source.as_ref()?; + let mesh = first.mesh; + let material = first.material; + let color = first.color; + let has_normals = mesh.normals.len() == mesh.verts.len(); + let (material_params, specular, ambient, advanced, flags) = + material_vertex_params(material); + let vertex = |index: usize, edge: bool| { + let normal = if edge { + [0.0, 1.0, 0.0] + } else if has_normals { + mesh.normals[index] + } else { + [0.0, 1.0, 0.0] + }; + let position = if edge { + source.edge_verts[index] + } else { + mesh.verts[index] + }; + let position_low = if edge { + source + .edge_verts_low + .get(index) + .copied() + .unwrap_or([0.0; 3]) + } else { + mesh.verts_low.get(index).copied().unwrap_or([0.0; 3]) + }; + let uvs = if edge { + [[0.0; 2]; 7] + } else { + material_uvs(material, position, normal) + }; + MeshVertex { + position, + normal, + color, + position_low, + material: material_params, + specular, + uv_diffuse: uvs[0], + ambient, + advanced, + flags, + uv_specular: uvs[1], + uv_reflection: uvs[2], + uv_opacity: uvs[3], + uv_bump: uvs[4], + uv_refraction: uvs[5], + uv_normal: uvs[6], + } + }; + let verts: Vec<_> = (0..mesh.verts.len()) + .map(|index| vertex(index, false)) + .collect(); + let edge_color = first + .set + .visual_style + .as_ref() + .map_or(first.display_mesh.color, |style| { + style.edge_color(first.display_mesh.color) + }); + let edge_verts: Vec<_> = if first.include_edges { + (0..source.edge_verts.len()) + .map(|index| { + let mut out = vertex(index, true); + out.color = edge_color; + out + }) + .collect() + } else { + Vec::new() + }; + let has_feature_edges = !edge_verts.is_empty(); + let mut wire_indices = Vec::new(); + if first.include_edges && !has_feature_edges { + wire_indices.reserve(first.indices.len() * 2); + for triangle in first.indices.chunks_exact(3) { + wire_indices.extend_from_slice(&[ + triangle[0], + triangle[1], + triangle[1], + triangle[2], + triangle[2], + triangle[0], + ]); + } + } + let transparent = color[3] < 0.999; + let face_indices = first.include_faces.then_some(first.indices.as_slice()); + let opaque_indices = if transparent { + &[][..] + } else { + face_indices.unwrap_or(&[]) + }; + let transparent_indices = if transparent { + face_indices.unwrap_or(&[]) + } else { + &[][..] + }; + let mut instances = Vec::with_capacity(parts.len()); + let mut highlights = Vec::with_capacity(parts.len()); + let mut handles = rustc_hash::FxHashSet::default(); + let mut bounds = [ + f32::INFINITY, + f32::INFINITY, + f32::INFINITY, + f32::NEG_INFINITY, + f32::NEG_INFINITY, + f32::NEG_INFINITY, + ]; + for part in parts { + let transform = part.set.instance_transform?; + let instance_start = instances.len() as u32; + instances.push(MeshInstanceGpu::from_transform(transform)); + bounds[0] = bounds[0].min(part.set.world_aabb[0]); + bounds[1] = bounds[1].min(part.set.world_aabb[1]); + bounds[2] = bounds[2].min(part.set.z_aabb[0]); + bounds[3] = bounds[3].max(part.set.world_aabb[2]); + bounds[4] = bounds[4].max(part.set.world_aabb[3]); + bounds[5] = bounds[5].max(part.set.z_aabb[1]); + if let Some(handle) = part + .display_mesh + .name + .parse::() + .ok() + .map(acadrust::Handle::new) + { + handles.insert(handle); + if first.include_faces { + highlights.push(MeshBatchRange { + handle, + index_start: 0, + index_count: first.indices.len() as u32, + transparent, + instance_start, + instance_count: 1, + }); + } + } + } + let triangles = if first.include_faces { + (first.indices.len() / 3) as u64 * instances.len() as u64 + } else { + 0 + }; + Some(( + make_chunk( + device, + &verts, + opaque_indices, + transparent_indices, + &wire_indices, + &edge_verts, + &highlights, + &instances, + &handles, + Some(bounds), + material, + ), + triangles, + )) } fn material_map_uv( @@ -681,10 +1016,19 @@ fn material_vertex_params( /// bounded; a single mesh too large for one chunk is split into triangle-soup /// sub-chunks so an XREF-heavy model can never overflow a single buffer (#203). pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec, u64) { + build_mesh_batch_filtered(device, sets, None) +} + +pub fn build_mesh_batch_filtered( + device: &wgpu::Device, + sets: &[MeshLodSet], + handles: Option<&rustc_hash::FxHashSet>, +) -> (Vec, u64) { // Derive the caps from the real device limit and vertex size. The previous // fixed 6 M-vertex cap assumed 40 B/vertex, but `position_low` (RTE) grew // MeshVertex to 52 B, so 6 M × 52 B = 312 MB blew past the 256 MB cap. - let budget = (device.limits().max_buffer_size as usize / 10) * 9; // 10% headroom + let hard_budget = (device.limits().max_buffer_size as usize / 10) * 9; + let budget = hard_budget.min(32 * 1024 * 1024); let vsize = std::mem::size_of::(); let max_verts = (budget / vsize).max(3); let max_tris = (budget / (6 * 4)).max(1); // wire-index buffer: 6 u32 per tri @@ -695,17 +1039,36 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec = Vec::new(); let mut wire_indices: Vec = Vec::new(); let mut edge_verts: Vec = Vec::new(); + let mut highlight_ranges: Vec = Vec::new(); + let mut chunk_handles: rustc_hash::FxHashSet = + rustc_hash::FxHashSet::default(); let mut total_tris = 0u64; let mut ordered: Vec> = Vec::new(); for set in sets { - let Some(mesh) = set.lods.iter().find(|mesh| !mesh.indices.is_empty()) else { + let Some(display_mesh) = set.lods.iter().find(|mesh| !mesh.indices.is_empty()) else { continue; }; - let triangle_count = mesh.indices.len() / 3; - let has_face_materials = mesh.triangle_material_handles.len() == triangle_count + let display_handle = display_mesh + .name + .parse::() + .ok() + .map(acadrust::Handle::new); + if handles.is_some_and(|wanted| { + display_handle.is_none_or(|handle| !wanted.contains(&handle)) + }) { + continue; + } + let mesh = set + .instance_source + .as_ref() + .and_then(|source| source.lods.iter().find(|mesh| !mesh.indices.is_empty())) + .unwrap_or(display_mesh); + let triangle_count = display_mesh.indices.len() / 3; + let has_face_materials = + display_mesh.triangle_material_handles.len() == triangle_count && !set.face_materials.is_empty(); - let has_face_colors = mesh.triangle_colors.len() == triangle_count - && mesh.triangle_colors.iter().any(Option::is_some); + let has_face_colors = display_mesh.triangle_colors.len() == triangle_count + && display_mesh.triangle_colors.iter().any(Option::is_some); let include_faces = set .visual_style .as_ref() @@ -716,7 +1079,7 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec (Vec (Vec, ), > = std::collections::BTreeMap::new(); - for (triangle, indices) in mesh.indices.chunks_exact(3).enumerate() { + for (triangle, indices) in display_mesh.indices.chunks_exact(3).enumerate() { let material = if has_face_materials { - mesh.triangle_material_handles[triangle] + display_mesh.triangle_material_handles[triangle] .and_then(|handle| set.face_materials.get(&handle)) .or(set.material.as_ref()) } else { set.material.as_ref() }; - let base_color = material.map_or(mesh.color, |material| material.diffuse); + let base_color = + material.map_or(display_mesh.color, |material| material.diffuse); let base_color = if has_face_colors { - mesh.triangle_colors[triangle].unwrap_or(base_color) + display_mesh.triangle_colors[triangle].unwrap_or(base_color) } else { base_color }; @@ -768,17 +1134,131 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec 0; + let mut instance_groups: std::collections::BTreeMap< + InstanceGroupKey, + Vec>, + > = std::collections::BTreeMap::new(); + let mut direct_parts = Vec::with_capacity(ordered.len()); + for mut part in ordered { + let eligible = storage_instancing + && part.set.instance_transform.is_some() + && part.set.instance_source.is_some() + && part.mesh.verts.len() <= max_verts + && part.indices.len() / 3 <= max_tris + && part + .set + .instance_source + .as_ref() + .is_some_and(|source| source.edge_verts.len() <= max_verts) + && part + .indices + .iter() + .all(|index| (*index as usize) < part.mesh.verts.len()); + if eligible { + let source = part + .set + .instance_source + .as_ref() + .expect("checked above") + .handle + .value(); + instance_groups + .entry(InstanceGroupKey { + material: material_key(part.material, part.color), + source, + part_slot: part.part_slot, + index_count: part.indices.len(), + index_hash: index_hash(&part.indices), + include_faces: part.include_faces, + include_edges: part.include_edges, + }) + .or_default() + .push(part); + } else { + // Compatibility and non-instanced meshes keep their already + // transformed display vertices in the ordinary static batch. + part.mesh = part.display_mesh; + direct_parts.push(part); + } + } + let sparse_groups: Vec<_> = instance_groups + .iter() + .filter_map(|(key, parts)| { + let first = parts.first()?; + let geometry_bytes = first + .mesh + .verts + .len() + .saturating_mul(std::mem::size_of::()) + .saturating_add(first.indices.len().saturating_mul(12)) + .saturating_add( + first + .set + .instance_source + .as_ref() + .map_or(0, |source| { + source + .edge_verts + .len() + .saturating_mul(std::mem::size_of::()) + }), + ); + let saved_bytes = geometry_bytes.saturating_mul(parts.len().saturating_sub(1)); + (parts.len() < 4 && saved_bytes < 512 * 1024).then_some(*key) + }) + .collect(); + for key in sparse_groups { + if let Some(parts) = instance_groups.remove(&key) { + for mut part in parts { + part.mesh = part.display_mesh; + direct_parts.push(part); + } + } + } + direct_parts.sort_by_key(|part| { + ( + material_key(part.material, part.color), + mesh_spatial_key(part.set, spatial_bounds), + ) + }); + let ordered = direct_parts; let mut active_key: Option = None; let mut active_material: Option<&crate::scene::model::material_model::MeshMaterial> = None; for part in ordered { @@ -786,6 +1266,12 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec() + .ok() + .map(acadrust::Handle::new); let key = material_key(material, part_color); if active_key.is_some_and(|active| active != key) && (!verts.is_empty() || !edge_verts.is_empty()) @@ -797,6 +1283,10 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec (Vec (Vec (Vec (Vec (Vec (Vec (Vec (Vec = if part.include_faces { + entity_handle + .map(|handle| { + vec![MeshBatchRange { + handle, + index_start: 0, + index_count: si.len() as u32, + transparent: is_transp, + instance_start: 0, + instance_count: 1, + }] + }) + .unwrap_or_default() + } else { + Vec::new() + }; if is_transp { + let sub_handles: rustc_hash::FxHashSet<_> = + entity_handle.into_iter().collect(); chunks.push(make_chunk( device, &sv, @@ -970,9 +1510,15 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec = + entity_handle.into_iter().collect(); chunks.push(make_chunk( device, &sv, @@ -980,11 +1526,16 @@ pub fn build_mesh_batch(device: &wgpu::Device, sets: &[MeshLodSet]) -> (Vec (Vec (Vec (Vec (Vec()) + .saturating_add(first.indices.len().saturating_mul(12)) + .saturating_add( + first + .set + .instance_source + .as_ref() + .map_or(0, |source| { + source + .edge_verts + .len() + .saturating_mul(std::mem::size_of::()) + }), + ) + .max(1); + // Keep repeated geometry bounded while retaining useful culling + // granularity. Small block definitions get spatial clusters of roughly + // 64–256 INSERTs; a huge source is duplicated only a few times. + let max_clusters = ((64 * 1024 * 1024) / source_bytes).clamp(1, 64); + let cluster_len = parts + .len() + .div_ceil(max_clusters) + .max(64); + for cluster in parts.chunks(cluster_len) { + if let Some((chunk, triangles)) = build_instanced_chunk(device, cluster) { + total_tris += triangles; + chunks.push(chunk); + } + } + } (chunks, total_tris) } - -impl MeshLodGpu { - #[allow(dead_code)] // built by the bypassed per-mesh upload_meshes path - pub fn new(device: &wgpu::Device, set: &MeshLodSet, highlight: Highlight) -> Self { - Self { - lods: set - .lods - .iter() - .filter(|m| !m.indices.is_empty()) - .map(|m| MeshGpu::new(device, m, set.material.as_ref(), highlight)) - .collect(), - world_aabb: set.world_aabb, - } - } -} - -impl MeshGpu { - pub fn new( - device: &wgpu::Device, - mesh: &MeshModel, - material: Option<&crate::scene::model::material_model::MeshMaterial>, - highlight: Highlight, - ) -> Self { - let has_normals = mesh.normals.len() == mesh.verts.len(); - // Blend the base colour toward the highlight so a selected / hovered - // solid reads clearly while keeping some shape shading. - let color = match highlight.tint() { - Some((hl, t)) => { - let mut c = [0.0f32; 4]; - for k in 0..4 { - c[k] = mesh.color[k] * (1.0 - t) + hl[k] * t; - } - c - } - None => mesh.color, - }; - let (material_params, specular, ambient, advanced, flags) = - material_vertex_params(material); - let vertices: Vec = mesh - .verts - .iter() - .enumerate() - .map(|(i, &pos)| { - let normal = if has_normals { - mesh.normals[i] - } else { - [0.0, 1.0, 0.0] - }; - let uv = material_uvs(material, pos, normal); - MeshVertex { - position: pos, - normal, - color, - position_low: mesh.verts_low.get(i).copied().unwrap_or([0.0; 3]), - material: material_params, - specular, - uv_diffuse: uv[0], - ambient, - advanced, - flags, - uv_specular: uv[1], - uv_reflection: uv[2], - uv_opacity: uv[3], - uv_bump: uv[4], - uv_refraction: uv[5], - uv_normal: uv[6], - } - }) - .collect(); - - let vertex_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some(&format!("mesh.vbuf.{}", mesh.name)), - contents: bytemuck::cast_slice(&vertices), - usage: wgpu::BufferUsages::VERTEX, - }); - - let index_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some(&format!("mesh.ibuf.{}", mesh.name)), - contents: bytemuck::cast_slice(&mesh.indices), - usage: wgpu::BufferUsages::INDEX, - }); - - // Wireframe-mode index buffer: expand each triangle into its - // three edge segments. Allocates ~2× the solid index count but - // is cheap compared to mesh tessellation and only happens when - // a new mesh is uploaded. - let mut wire_indices: Vec = Vec::with_capacity(mesh.indices.len() * 2); - for tri in mesh.indices.chunks_exact(3) { - let (a, b, c) = (tri[0], tri[1], tri[2]); - wire_indices.extend_from_slice(&[a, b, b, c, c, a]); - } - let wire_index_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some(&format!("mesh.wire_ibuf.{}", mesh.name)), - contents: bytemuck::cast_slice(&wire_indices), - usage: wgpu::BufferUsages::INDEX, - }); - - Self { - vertex_buffer, - index_buffer, - index_count: mesh.indices.len() as u32, - wire_index_buffer, - wire_index_count: wire_indices.len() as u32, - } - } -} diff --git a/src/scene/pipeline/mod.rs b/src/scene/pipeline/mod.rs index f18907653..329f0d54d 100644 --- a/src/scene/pipeline/mod.rs +++ b/src/scene/pipeline/mod.rs @@ -23,7 +23,6 @@ use iced::{Rectangle, Size}; pub use face3d_gpu::Face3DGpu; pub use wipeout_gpu::WipeoutGpu; pub use image_gpu::ImageGpu; -pub use mesh_gpu::MeshLodGpu; pub use uniforms::Uniforms; pub use viewcube::ViewCubePipeline; pub use wire_gpu::WireGpu; @@ -36,6 +35,46 @@ use crate::scene::model::wire_model::WireModel; /// MSAA sample count for the main drawing pipelines. const MSAA_SAMPLES: u32 = 4; +#[derive(Clone, Copy, PartialEq, Eq)] +enum MeshHighlightKind { + Selected, + Hover, +} + +#[derive(Clone, Copy)] +struct MeshResidentRange { + chunk: usize, + dynamic: bool, + index_start: u32, + index_count: u32, + transparent: bool, + instance_start: u32, + instance_count: u32, +} + +#[derive(Clone, Copy)] +struct MeshHighlightDraw { + range: MeshResidentRange, + kind: MeshHighlightKind, +} + +#[repr(C)] +#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)] +struct MeshCullItem { + min: [f32; 4], + max: [f32; 4], + counts: [u32; 4], + meta: [u32; 4], +} + +#[repr(C)] +#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)] +struct MeshCullUniform { + view_rot: [f32; 16], + eye: [f32; 4], + count: [u32; 4], +} + pub struct Pipeline { wire_pipeline: wgpu::RenderPipeline, /// Stamps clip-boundary polygons into the stencil buffer (viewports + XCLIP). @@ -66,7 +105,8 @@ pub struct Pipeline { mesh_pipeline: wgpu::RenderPipeline, /// Depth-write-disabled variant of `mesh_pipeline` for non-opaque solids. mesh_transparent_pipeline: wgpu::RenderPipeline, - mesh_highlight_pipeline: wgpu::RenderPipeline, + mesh_selected_pipeline: wgpu::RenderPipeline, + mesh_hover_pipeline: wgpu::RenderPipeline, /// Wireframe variant of the mesh pipeline (LineList topology, same /// vertex layout / shader). Used when the active render mode is /// Wireframe 2D or Wireframe 3D so 3D solids draw as their @@ -80,6 +120,16 @@ pub struct Pipeline { mesh_depth_pipeline: wgpu::RenderPipeline, mesh_material_bgl: wgpu::BindGroupLayout, mesh_default_material_bind_group: wgpu::BindGroup, + mesh_cull_pipeline: Option, + mesh_cull_bgl: Option, + mesh_cull_uniform: Option, + mesh_cull_items: Option, + mesh_cull_bind_group: Option, + mesh_opaque_indirect: Option, + mesh_transparent_indirect: Option, + mesh_wire_indirect: Option, + mesh_edge_indirect: Option, + mesh_cull_count: u32, face3d_pipeline: wgpu::RenderPipeline, /// Depth-only variant of the face3d pipeline (no color writes, /// writes depth). Paired with `mesh_depth_pipeline` for HiddenLine. @@ -206,29 +256,26 @@ pub struct Pipeline { /// viewport rect. Recomputed by `compute_wipeout_lod`. wipeout_skip_flags: Vec, gpu_images: Vec, - gpu_meshes: Vec, /// Batched mesh geometry — every solid's LOD0 concatenated into a few large /// buffers so the whole set draws in a handful of calls instead of one per - /// solid. Rebuilt only when the geometry epoch changes (see - /// `cached_mesh_batch_epoch`), so hover / selection never re-pack it. + /// solid. Hover / selection never re-pack it. gpu_mesh_batch: Vec, - /// `geometry_epoch` the batch was built for. `u64::MAX` = not built yet. - pub cached_mesh_batch_epoch: u64, - /// Tinted copies of just the selected / hovered solids, drawn over the - /// static base batch so highlight shows without re-packing the whole set. - /// Rebuilt only when the selection/hover or geometry changes — O(highlighted). - gpu_mesh_highlight: Vec, + /// Small mutable working set rebuilt after entity-only edits. Static chunks + /// touched by an edit are disabled wholesale and their current handles are + /// repacked here, keeping the rest of a multi-million-triangle scene resident. + gpu_mesh_dynamic: Vec, + mesh_disabled_chunks: rustc_hash::FxHashSet, + mesh_dynamic_handles: rustc_hash::FxHashSet, + /// Wire content generation whose geometry the static+dynamic mesh state + /// mirrors. It gates replay of the same per-entity journal handoff. + pub cached_mesh_content_id: u64, + /// Draw ranges for each entity inside the resident mesh chunks. Highlight + /// overlays reuse these buffers instead of uploading duplicate geometry. + mesh_ranges_by_handle: + rustc_hash::FxHashMap>, + mesh_highlight_draws: Vec, /// `(geometry_epoch, selection_generation)` the highlight overlay was built for. pub cached_highlight_key: (u64, u64), - /// Per-mesh LOD level (0=high, 1=mid, 2=low) picked each frame from - /// the projected pixel diagonal. Mirrors `hatch_pixel_scissors` — - /// recomputed in `compute_mesh_lod`. - mesh_lod_levels: Vec, - /// Per-mesh frustum-visible flag (Phase 2.2). `false` when the - /// mesh's projected AABB falls entirely outside the viewport rect - /// — the draw loop skips it. Mirrors `mesh_lod_levels` length / - /// index space; recomputed alongside it. - mesh_visible: Vec, /// Batched 3DFACE fill (all faces in one buffer) and edges (merged wire). gpu_face3d_fill: Option, gpu_face3d_edges: Vec, @@ -263,11 +310,6 @@ pub struct Pipeline { /// a pick bumps only `selection_generation`, refreshing the overlay without /// touching the main wire buffers. pub cached_selection: (u64, u64), - /// `(geometry_epoch, selection_generation)` the resident mesh buffers were - /// uploaded for. Bumps on a selection/hover change so highlighted solids - /// re-upload with their tint. - #[allow(dead_code)] // per-mesh LOD path, bypassed by the batched mesh draw - pub cached_mesh_key: (u64, u64), /// `(wire_content_id, face3d_fill_active)` the Face3D edge/fill buffers were /// uploaded for. A stable content id avoids retaining the resident wire Arc: /// that Arc must stay uniquely owned by Scene so a small edit can splice it @@ -834,12 +876,135 @@ impl Pipeline { }; // ── Mesh pipeline ────────────────────────────────────────────────── + let mesh_storage_instancing = + device.limits().max_storage_buffers_per_shader_stage > 0; + let mesh_source = include_str!("../../shaders/mesh.wgsl"); + let mesh_source = if mesh_storage_instancing { + std::borrow::Cow::Borrowed(mesh_source) + } else { + std::borrow::Cow::Owned(mesh_source.replace( + "var mesh_instances: array;", + "var mesh_instances: array;", + )) + }; let mesh_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("mesh.shader"), - source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!( - "../../shaders/mesh.wgsl" - ))), + source: wgpu::ShaderSource::Wgsl(mesh_source), }); + let (mesh_cull_bgl, mesh_cull_pipeline, mesh_cull_uniform) = + if mesh_storage_instancing { + let bgl = + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("mesh.cull.bgl"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: true, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: false, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 3, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: false, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 4, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: false, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 5, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: false, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + let shader = + device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("mesh.cull.shader"), + source: wgpu::ShaderSource::Wgsl( + std::borrow::Cow::Borrowed(include_str!( + "../../shaders/mesh_cull.wgsl" + )), + ), + }); + let layout = + device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("mesh.cull.layout"), + bind_group_layouts: &[&bgl], + push_constant_ranges: &[], + }); + let pipeline = + device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("mesh.cull.pipeline"), + layout: Some(&layout), + module: &shader, + entry_point: Some("main"), + compilation_options: + wgpu::PipelineCompilationOptions::default(), + cache: None, + }); + let uniform = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("mesh.cull.uniform"), + size: std::mem::size_of::() as u64, + usage: wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + (Some(bgl), Some(pipeline), Some(uniform)) + } else { + (None, None, None) + }; let mesh_material_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { @@ -967,6 +1132,20 @@ impl Pipeline { }, count: None, }, + wgpu::BindGroupLayoutEntry { + binding: 15, + visibility: wgpu::ShaderStages::VERTEX, + ty: wgpu::BindingType::Buffer { + ty: if mesh_storage_instancing { + wgpu::BufferBindingType::Storage { read_only: true } + } else { + wgpu::BufferBindingType::Uniform + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, ], }); let mesh_default_material_bind_group = mesh_gpu::create_material_bind_group( @@ -974,6 +1153,7 @@ impl Pipeline { queue, &mesh_material_bgl, None, + None, ); let mesh_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { @@ -1075,49 +1255,56 @@ impl Pipeline { cache: None, }); - // Highlight variant — same shader / vertex layout, but `depth_compare = - // Always` and no depth write, so a hovered / selected solid's tint is - // drawn on top of everything (it shows through occluding geometry) and - // doesn't disturb the depth buffer for later passes. - let mesh_highlight_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { - label: Some("mesh.highlight.pipeline"), - layout: Some(&mesh_layout), - vertex: wgpu::VertexState { - module: &mesh_shader, - entry_point: Some("vs_main"), - buffers: &[mesh_gpu::MeshVertex::layout()], - compilation_options: wgpu::PipelineCompilationOptions::default(), - }, - primitive: wgpu::PrimitiveState { - topology: wgpu::PrimitiveTopology::TriangleList, - cull_mode: None, - ..Default::default() - }, - depth_stencil: Some(wgpu::DepthStencilState { - format: wgpu::TextureFormat::Depth24PlusStencil8, - depth_write_enabled: false, - depth_compare: wgpu::CompareFunction::Always, - stencil: content_stencil.clone(), - bias: wgpu::DepthBiasState::default(), - }), - multisample: wgpu::MultisampleState { - count: MSAA_SAMPLES, - mask: !0, - alpha_to_coverage_enabled: false, - }, - fragment: Some(wgpu::FragmentState { - module: &mesh_shader, - entry_point: Some("fs_main"), - targets: &[Some(wgpu::ColorTargetState { - format, - blend: Some(wgpu::BlendState::ALPHA_BLENDING), - write_mask: wgpu::ColorWrites::ALL, - })], - compilation_options: wgpu::PipelineCompilationOptions::default(), - }), - multiview: None, - cache: None, - }); + // Highlight variants reuse the resident mesh buffers and differ only in + // their fixed fragment tint. No selected mesh geometry is re-uploaded. + let make_mesh_highlight_pipeline = + |label: &'static str, fragment_entry: &'static str| { + device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some(label), + layout: Some(&mesh_layout), + vertex: wgpu::VertexState { + module: &mesh_shader, + entry_point: Some("vs_main"), + buffers: &[mesh_gpu::MeshVertex::layout()], + compilation_options: wgpu::PipelineCompilationOptions::default(), + }, + primitive: wgpu::PrimitiveState { + topology: wgpu::PrimitiveTopology::TriangleList, + cull_mode: None, + ..Default::default() + }, + depth_stencil: Some(wgpu::DepthStencilState { + format: wgpu::TextureFormat::Depth24PlusStencil8, + depth_write_enabled: false, + depth_compare: wgpu::CompareFunction::Always, + stencil: content_stencil.clone(), + bias: wgpu::DepthBiasState::default(), + }), + multisample: wgpu::MultisampleState { + count: MSAA_SAMPLES, + mask: !0, + alpha_to_coverage_enabled: false, + }, + fragment: Some(wgpu::FragmentState { + module: &mesh_shader, + entry_point: Some(fragment_entry), + targets: &[Some(wgpu::ColorTargetState { + format, + blend: Some(wgpu::BlendState::ALPHA_BLENDING), + write_mask: wgpu::ColorWrites::ALL, + })], + compilation_options: wgpu::PipelineCompilationOptions::default(), + }), + multiview: None, + cache: None, + }) + }; + let mesh_selected_pipeline = make_mesh_highlight_pipeline( + "mesh.highlight.selected.pipeline", + "fs_highlight_selected", + ); + let mesh_hover_pipeline = + make_mesh_highlight_pipeline("mesh.highlight.hover.pipeline", "fs_highlight_hover"); // Wireframe variant — same shader / vertex layout / depth state, // only the input topology changes (LineList) and back-face @@ -1570,12 +1757,23 @@ impl Pipeline { silhouette_vcount: 0, mesh_pipeline, mesh_transparent_pipeline, - mesh_highlight_pipeline, + mesh_selected_pipeline, + mesh_hover_pipeline, mesh_wireframe_pipeline, mesh_edge_black_pipeline, mesh_depth_pipeline, mesh_material_bgl, mesh_default_material_bind_group, + mesh_cull_pipeline, + mesh_cull_bgl, + mesh_cull_uniform, + mesh_cull_items: None, + mesh_cull_bind_group: None, + mesh_opaque_indirect: None, + mesh_transparent_indirect: None, + mesh_wire_indirect: None, + mesh_edge_indirect: None, + mesh_cull_count: 0, face3d_pipeline, face3d_depth_pipeline, uniform_buffer, @@ -1619,13 +1817,14 @@ impl Pipeline { gpu_wipeouts: vec![], wipeout_skip_flags: vec![], gpu_images: vec![], - gpu_meshes: vec![], gpu_mesh_batch: vec![], - cached_mesh_batch_epoch: u64::MAX, - gpu_mesh_highlight: vec![], + gpu_mesh_dynamic: vec![], + mesh_disabled_chunks: rustc_hash::FxHashSet::default(), + mesh_dynamic_handles: rustc_hash::FxHashSet::default(), + cached_mesh_content_id: u64::MAX, + mesh_ranges_by_handle: rustc_hash::FxHashMap::default(), + mesh_highlight_draws: vec![], cached_highlight_key: (u64::MAX, u64::MAX), - mesh_lod_levels: vec![], - mesh_visible: vec![], gpu_face3d_fill: None, gpu_face3d_edges: vec![], viewcube, @@ -1640,7 +1839,6 @@ impl Pipeline { cached_epoch: (u64::MAX, u64::MAX, u64::MAX), cached_wire_id: u64::MAX, cached_selection: (u64::MAX, u64::MAX), - cached_mesh_key: (u64::MAX, u64::MAX), cached_face3d_key: (u64::MAX, false), wire_handle_index: std::sync::Arc::new(rustc_hash::FxHashMap::default()), render_sig: u64::MAX, @@ -2160,35 +2358,6 @@ impl Pipeline { } } - #[allow(dead_code)] // per-mesh upload path, bypassed by upload_mesh_batch - pub fn upload_meshes( - &mut self, - device: &wgpu::Device, - meshes: &[MeshLodSet], - selected: &rustc_hash::FxHashSet, - hover: Option, - ) { - self.gpu_meshes = meshes - .iter() - .filter(|s| s.lods.iter().any(|m| !m.indices.is_empty())) - .map(|s| { - // A mesh's name is its source entity handle (decimal). Selection - // wins over hover when both apply to the same solid. - let h = s - .lods - .first() - .and_then(|m| m.name.parse::().ok()) - .map(acadrust::Handle::new); - let mode = match h { - Some(h) if selected.contains(&h) => mesh_gpu::Highlight::Selected, - Some(h) if Some(h) == hover => mesh_gpu::Highlight::Hover, - _ => mesh_gpu::Highlight::None, - }; - MeshLodGpu::new(device, s, mode) - }) - .collect(); - } - /// Build the batched mesh buffers (a few large buffers for the whole solid /// set, drawn in a handful of calls). Selection/hover tint is intentionally /// not applied here — the batch is geometry-only and stays resident across @@ -2199,67 +2368,400 @@ impl Pipeline { queue: &wgpu::Queue, meshes: &[MeshLodSet], ) { - let (mut chunks, _tris) = mesh_gpu::build_mesh_batch(device, meshes); + let started = iced::time::Instant::now(); + let (mut chunks, triangles) = mesh_gpu::build_mesh_batch(device, meshes); for chunk in &mut chunks { chunk.material_bind_group = Some(mesh_gpu::create_material_bind_group( device, queue, &self.mesh_material_bgl, chunk.material.as_ref(), + Some(&chunk.instance_buffer), )); } + self.mesh_ranges_by_handle.clear(); + for (chunk_index, chunk) in chunks.iter().enumerate() { + for range in &chunk.highlight_ranges { + self.mesh_ranges_by_handle + .entry(range.handle) + .or_default() + .push(MeshResidentRange { + chunk: chunk_index, + dynamic: false, + index_start: range.index_start, + index_count: range.index_count, + transparent: range.transparent, + instance_start: range.instance_start, + instance_count: range.instance_count, + }); + } + } + self.mesh_highlight_draws.clear(); + self.cached_highlight_key = (u64::MAX, u64::MAX); + self.gpu_mesh_dynamic.clear(); + self.mesh_disabled_chunks.clear(); + self.mesh_dynamic_handles.clear(); self.gpu_mesh_batch = chunks; + self.rebuild_mesh_cull_resources(device); + if crate::perf::enabled() { + let instances: u64 = self + .gpu_mesh_batch + .iter() + .map(|chunk| chunk.instance_count as u64) + .sum(); + crate::perf_record!( + "[perf] mesh-batch {:>7.1}ms sets={} chunks={} instances={} triangles={}", + started.elapsed().as_secs_f64() * 1000.0, + meshes.len(), + self.gpu_mesh_batch.len(), + instances, + triangles, + ); + } } - /// Build tinted copies of just the selected / hovered solids. Drawn over the - /// static base batch with the LessEqual mesh pipeline so the tint shows on - /// top at the same depth. O(highlighted), so a pick never touches the rest. - pub fn upload_mesh_highlight( + fn rebuild_mesh_cull_resources(&mut self, device: &wgpu::Device) { + let count = self.gpu_mesh_batch.len() + self.gpu_mesh_dynamic.len(); + // Below this point CPU projection is cheaper than a compute pass plus + // indirect command reads. The large-scene path turns on automatically. + if count < 128 + || self.mesh_cull_pipeline.is_none() + || self.mesh_cull_bgl.is_none() + || self.mesh_cull_uniform.is_none() + { + self.mesh_cull_bind_group = None; + self.mesh_cull_items = None; + self.mesh_opaque_indirect = None; + self.mesh_transparent_indirect = None; + self.mesh_wire_indirect = None; + self.mesh_edge_indirect = None; + self.mesh_cull_count = 0; + return; + } + let mut items = Vec::with_capacity(count); + for (index, chunk) in self.gpu_mesh_batch.iter().enumerate() { + items.push(MeshCullItem { + min: [ + chunk.world_aabb[0], + chunk.world_aabb[1], + chunk.world_aabb[2], + 0.0, + ], + max: [ + chunk.world_aabb[3], + chunk.world_aabb[4], + chunk.world_aabb[5], + 0.0, + ], + counts: [ + chunk.index_count, + chunk.transp_index_count, + chunk.wire_index_count, + chunk.edge_vertex_count, + ], + meta: [ + chunk.instance_count, + (!self.mesh_disabled_chunks.contains(&index)) as u32, + 0, + 0, + ], + }); + } + for chunk in &self.gpu_mesh_dynamic { + items.push(MeshCullItem { + min: [ + chunk.world_aabb[0], + chunk.world_aabb[1], + chunk.world_aabb[2], + 0.0, + ], + max: [ + chunk.world_aabb[3], + chunk.world_aabb[4], + chunk.world_aabb[5], + 0.0, + ], + counts: [ + chunk.index_count, + chunk.transp_index_count, + chunk.wire_index_count, + chunk.edge_vertex_count, + ], + meta: [chunk.instance_count, 1, 0, 0], + }); + } + let items_buffer = + device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("mesh.cull.items"), + contents: bytemuck::cast_slice(&items), + usage: wgpu::BufferUsages::STORAGE, + }); + let indexed_size = (count * std::mem::size_of::()) + as u64; + let draw_size = + (count * std::mem::size_of::()) as u64; + let make_output = |label: &'static str, size: u64| { + device.create_buffer(&wgpu::BufferDescriptor { + label: Some(label), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::INDIRECT, + mapped_at_creation: false, + }) + }; + let opaque = make_output("mesh.cull.opaque", indexed_size); + let transparent = make_output("mesh.cull.transparent", indexed_size); + let wire = make_output("mesh.cull.wire", indexed_size); + let edge = make_output("mesh.cull.edge", draw_size); + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("mesh.cull.bind_group"), + layout: self.mesh_cull_bgl.as_ref().expect("checked above"), + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: self + .mesh_cull_uniform + .as_ref() + .expect("checked above") + .as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: items_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: opaque.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: transparent.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 4, + resource: wire.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 5, + resource: edge.as_entire_binding(), + }, + ], + }); + self.mesh_cull_bind_group = Some(bind_group); + self.mesh_cull_items = Some(items_buffer); + self.mesh_opaque_indirect = Some(opaque); + self.mesh_transparent_indirect = Some(transparent); + self.mesh_wire_indirect = Some(wire); + self.mesh_edge_indirect = Some(edge); + self.mesh_cull_count = count as u32; + } + + fn rebuild_mesh_range_map(&mut self) { + self.mesh_ranges_by_handle.clear(); + for (chunk_index, chunk) in self.gpu_mesh_batch.iter().enumerate() { + if self.mesh_disabled_chunks.contains(&chunk_index) { + continue; + } + for range in &chunk.highlight_ranges { + self.mesh_ranges_by_handle + .entry(range.handle) + .or_default() + .push(MeshResidentRange { + chunk: chunk_index, + dynamic: false, + index_start: range.index_start, + index_count: range.index_count, + transparent: range.transparent, + instance_start: range.instance_start, + instance_count: range.instance_count, + }); + } + } + for (chunk_index, chunk) in self.gpu_mesh_dynamic.iter().enumerate() { + for range in &chunk.highlight_ranges { + self.mesh_ranges_by_handle + .entry(range.handle) + .or_default() + .push(MeshResidentRange { + chunk: chunk_index, + dynamic: true, + index_start: range.index_start, + index_count: range.index_count, + transparent: range.transparent, + instance_start: range.instance_start, + instance_count: range.instance_count, + }); + } + } + self.mesh_highlight_draws.clear(); + self.cached_highlight_key = (u64::MAX, u64::MAX); + } + + /// Patch an entity-only geometry change without rebuilding the resident + /// mesh set. Any 32 MiB static chunk containing a changed handle is retired; + /// the current versions of all its handles are rebuilt into a small dynamic + /// working set. A bounded threshold falls back to a full rebuild before the + /// working set can grow into a second copy of the drawing. + pub fn patch_mesh_batch( &mut self, device: &wgpu::Device, + queue: &wgpu::Queue, meshes: &[MeshLodSet], + changes: &[(acadrust::Handle, crate::scene::ChangeKind)], + ) -> bool { + let started = iced::time::Instant::now(); + if self.gpu_mesh_batch.is_empty() || changes.is_empty() { + return false; + } + let current_handles: rustc_hash::FxHashSet<_> = meshes + .iter() + .filter_map(|set| { + set.lods + .first() + .and_then(|mesh| mesh.name.parse::().ok()) + .map(acadrust::Handle::new) + }) + .collect(); + let changed: rustc_hash::FxHashSet<_> = changes + .iter() + .map(|(handle, _)| *handle) + .filter(|handle| { + current_handles.contains(handle) + || self.mesh_dynamic_handles.contains(handle) + || self + .gpu_mesh_batch + .iter() + .any(|chunk| chunk.handles.contains(handle)) + }) + .collect(); + if changed.is_empty() { + return true; + } + for (chunk_index, chunk) in self.gpu_mesh_batch.iter().enumerate() { + if chunk.handles.iter().any(|handle| changed.contains(handle)) { + self.mesh_disabled_chunks.insert(chunk_index); + self.mesh_dynamic_handles + .extend(chunk.handles.iter().copied()); + } + } + self.mesh_dynamic_handles.extend(changed.iter().copied()); + let disabled_limit = self.gpu_mesh_batch.len().div_ceil(2).max(8); + if self.mesh_disabled_chunks.len() > disabled_limit + || self.mesh_dynamic_handles.len() > 4096 + { + return false; + } + let (mut chunks, _) = mesh_gpu::build_mesh_batch_filtered( + device, + meshes, + Some(&self.mesh_dynamic_handles), + ); + for chunk in &mut chunks { + chunk.material_bind_group = Some(mesh_gpu::create_material_bind_group( + device, + queue, + &self.mesh_material_bgl, + chunk.material.as_ref(), + Some(&chunk.instance_buffer), + )); + } + self.gpu_mesh_dynamic = chunks; + self.rebuild_mesh_cull_resources(device); + self.rebuild_mesh_range_map(); + if crate::perf::enabled() { + crate::perf_record!( + "[perf] mesh-patch {:>7.1}ms changed={} retired={} dynamic_handles={} dynamic_chunks={}", + started.elapsed().as_secs_f64() * 1000.0, + changed.len(), + self.mesh_disabled_chunks.len(), + self.mesh_dynamic_handles.len(), + self.gpu_mesh_dynamic.len(), + ); + } + true + } + + /// Refresh selection/hover draw ranges. The overlay references resident + /// chunk buffers; changing hover never allocates or uploads mesh geometry. + pub fn update_mesh_highlight( + &mut self, selected: &rustc_hash::FxHashSet, hover: Option, ) { - use mesh_gpu::{Highlight, MeshGpu}; let mut out = Vec::new(); - if selected.is_empty() && hover.is_none() { - self.gpu_mesh_highlight = out; - return; + for handle in selected { + if let Some(ranges) = self.mesh_ranges_by_handle.get(handle) { + out.extend(ranges.iter().copied().map(|range| MeshHighlightDraw { + range, + kind: MeshHighlightKind::Selected, + })); + } } - for set in meshes { - let Some(m) = set.lods.iter().find(|m| !m.indices.is_empty()) else { - continue; - }; - let h = m.name.parse::().ok().map(acadrust::Handle::new); - let mode = match h { - Some(h) if selected.contains(&h) => Highlight::Selected, - Some(h) if Some(h) == hover => Highlight::Hover, - _ => continue, - }; - out.push(MeshGpu::new(device, m, set.material.as_ref(), mode)); + if let Some(handle) = hover.filter(|handle| !selected.contains(handle)) { + if let Some(ranges) = self.mesh_ranges_by_handle.get(&handle) { + out.extend(ranges.iter().copied().map(|range| MeshHighlightDraw { + range, + kind: MeshHighlightKind::Hover, + })); + } } - self.gpu_mesh_highlight = out; + self.mesh_highlight_draws = out; } - /// Per-frame mesh LOD selector. Picks slot 0/1/2 based on the - /// projected pixel diagonal of each mesh's `world_aabb` (Phase 3.4 - /// ladder: >200 px → 0, 50–200 → 1, <50 → 2). Falls back to the - /// nearest available lower slot when a level wasn't generated. - pub fn compute_mesh_lod(&mut self, view_rot: glam::Mat4, eye: glam::DVec3, clip_w: u32, clip_h: u32) { - self.mesh_lod_levels = self - .gpu_meshes + fn active_mesh_chunks_indexed( + &self, + ) -> impl Iterator { + let dynamic_offset = self.gpu_mesh_batch.len(); + self.gpu_mesh_batch .iter() - .map(|m| pick_mesh_lod(m, view_rot, eye, clip_w, clip_h)) - .collect(); - // Phase 2.2 — frustum-visibility flag per mesh. Cheap: same - // 4-corner projection used for LOD selection, just answering a - // different question (any corner inside the viewport rect?). - self.mesh_visible = self - .gpu_meshes - .iter() - .map(|m| !aabb_offscreen(m.world_aabb, view_rot, eye, clip_w, clip_h)) - .collect(); + .enumerate() + .filter(|(index, _)| !self.mesh_disabled_chunks.contains(index)) + .chain( + self.gpu_mesh_dynamic + .iter() + .enumerate() + .map(move |(index, chunk)| (dynamic_offset + index, chunk)), + ) + } + + /// Per-frame coarse frustum culling for the spatially sorted resident + /// chunks. It changes only draw eligibility and never rebuilds GPU data. + pub fn compute_mesh_lod( + &mut self, + queue: &wgpu::Queue, + view_rot: glam::Mat4, + eye: glam::DVec3, + clip_w: u32, + clip_h: u32, + ) { + if self.mesh_cull_count != 0 { + if let Some(uniform) = &self.mesh_cull_uniform { + queue.write_buffer( + uniform, + 0, + bytemuck::bytes_of(&MeshCullUniform { + view_rot: view_rot.to_cols_array(), + eye: [eye.x as f32, eye.y as f32, eye.z as f32, 0.0], + count: [self.mesh_cull_count, 0, 0, 0], + }), + ); + } + for (index, chunk) in self.gpu_mesh_batch.iter_mut().enumerate() { + chunk.visible = !self.mesh_disabled_chunks.contains(&index); + } + for chunk in &mut self.gpu_mesh_dynamic { + chunk.visible = true; + } + return; + } + for (index, chunk) in self.gpu_mesh_batch.iter_mut().enumerate() { + chunk.visible = + !self.mesh_disabled_chunks.contains(&index) + && !aabb3_offscreen(chunk.world_aabb, view_rot, eye, clip_w, clip_h); + } + for chunk in &mut self.gpu_mesh_dynamic { + chunk.visible = + !aabb3_offscreen(chunk.world_aabb, view_rot, eye, clip_w, clip_h); + } } pub fn upload_hatches( @@ -2516,8 +3018,23 @@ impl Pipeline { } } + if self.mesh_cull_count != 0 { + if let (Some(pipeline), Some(bind_group)) = + (&self.mesh_cull_pipeline, &self.mesh_cull_bind_group) + { + let mut pass = + encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("mesh.cull.compute_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(pipeline); + pass.set_bind_group(0, bind_group, &[]); + pass.dispatch_workgroups(self.mesh_cull_count.div_ceil(64), 1, 1); + } + } + // ── Pass 4: solid meshes (batched) ──────────────────────────────── - if !self.gpu_mesh_batch.is_empty() { + if !self.gpu_mesh_batch.is_empty() || !self.gpu_mesh_dynamic.is_empty() { let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("mesh.render_pass"), color_attachments: &[Some(wgpu::RenderPassColorAttachment { @@ -2562,7 +3079,10 @@ impl Pipeline { // Depth-only prepass: every solid surface occludes hidden edges, // so both the opaque and the transparent tris write depth here. pass.set_pipeline(&self.mesh_depth_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2576,18 +3096,45 @@ impl Pipeline { c.index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_opaque_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed(0..c.index_count, 0, 0..c.instance_count); + } } if c.transp_index_count != 0 { pass.set_index_buffer( c.transp_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.transp_index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_transparent_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed( + 0..c.transp_index_count, + 0, + 0..c.instance_count, + ); + } } } pass.set_pipeline(&self.mesh_wireframe_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2602,12 +3149,36 @@ impl Pipeline { c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.wire_index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_wire_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed( + 0..c.wire_index_count, + 0, + 0..c.instance_count, + ); + } } // ACIS solid B-rep feature edges (LineList, non-indexed). if c.edge_vertex_count != 0 { pass.set_vertex_buffer(0, c.edge_vertex_buffer.slice(..)); - pass.draw(0..c.edge_vertex_count, 0..1); + if let Some(indirect) = &self.mesh_edge_indirect { + pass.draw_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndirectArgs, + >() as u64, + ); + } else { + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); + } } } // DISPSILH silhouettes (whole batch, one buffer). @@ -2618,7 +3189,10 @@ impl Pipeline { } else { if mesh_wireframe { pass.set_pipeline(&self.mesh_wireframe_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2632,11 +3206,35 @@ impl Pipeline { c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.wire_index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_wire_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed( + 0..c.wire_index_count, + 0, + 0..c.instance_count, + ); + } } if c.edge_vertex_count != 0 { pass.set_vertex_buffer(0, c.edge_vertex_buffer.slice(..)); - pass.draw(0..c.edge_vertex_count, 0..1); + if let Some(indirect) = &self.mesh_edge_indirect { + pass.draw_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndirectArgs, + >() as u64, + ); + } else { + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); + } } } // DISPSILH silhouettes. @@ -2647,7 +3245,10 @@ impl Pipeline { } else { // Opaque fills first (they write depth). pass.set_pipeline(&self.mesh_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2660,13 +3261,26 @@ impl Pipeline { } pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); pass.set_index_buffer(c.index_buffer.slice(..), wgpu::IndexFormat::Uint32); - pass.draw_indexed(0..c.index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_opaque_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed(0..c.index_count, 0, 0..c.instance_count); + } } // Transparent fills last, with depth writes disabled, so they // blend over the opaque geometry behind them instead of // culling it via the depth buffer. pass.set_pipeline(&self.mesh_transparent_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2682,22 +3296,80 @@ impl Pipeline { c.transp_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.transp_index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_transparent_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed( + 0..c.transp_index_count, + 0, + 0..c.instance_count, + ); + } } } - // Selection / hover highlight: tinted copies of the picked - // solids, drawn last with the Always-depth highlight pipeline so - // the tint shows on top even when the solid is behind others. - if !self.gpu_mesh_highlight.is_empty() { - pass.set_pipeline(&self.mesh_highlight_pipeline); - pass.set_bind_group(1, &self.mesh_default_material_bind_group, &[]); - for g in &self.gpu_mesh_highlight { - if g.index_count == 0 { + // Selection / hover highlight reuses index ranges already + // resident in the chunk buffers; hover never uploads geometry. + for kind in [MeshHighlightKind::Selected, MeshHighlightKind::Hover] { + if !self + .mesh_highlight_draws + .iter() + .any(|draw| draw.kind == kind) + { + continue; + } + pass.set_pipeline(match kind { + MeshHighlightKind::Selected => &self.mesh_selected_pipeline, + MeshHighlightKind::Hover => &self.mesh_hover_pipeline, + }); + for draw in self + .mesh_highlight_draws + .iter() + .filter(|draw| draw.kind == kind) + { + let chunk = if draw.range.dynamic { + self.gpu_mesh_dynamic.get(draw.range.chunk) + } else { + self.gpu_mesh_batch.get(draw.range.chunk) + }; + let Some(chunk) = chunk else { + continue; + }; + if !chunk.visible || draw.range.index_count == 0 { continue; } - pass.set_vertex_buffer(0, g.vertex_buffer.slice(..)); - pass.set_index_buffer(g.index_buffer.slice(..), wgpu::IndexFormat::Uint32); - pass.draw_indexed(0..g.index_count, 0, 0..1); + pass.set_bind_group( + 1, + chunk + .material_bind_group + .as_ref() + .unwrap_or(&self.mesh_default_material_bind_group), + &[], + ); + pass.set_vertex_buffer(0, chunk.vertex_buffer.slice(..)); + if draw.range.transparent { + pass.set_index_buffer( + chunk.transp_index_buffer.slice(..), + wgpu::IndexFormat::Uint32, + ); + } else { + pass.set_index_buffer( + chunk.index_buffer.slice(..), + wgpu::IndexFormat::Uint32, + ); + } + pass.draw_indexed( + draw.range.index_start + ..draw.range.index_start + draw.range.index_count, + 0, + draw.range.instance_start + ..draw.range.instance_start + draw.range.instance_count, + ); } } // *WithEdges variants: overlay edge segments on top of the shaded @@ -2705,7 +3377,10 @@ impl Pipeline { // test keeps the edges visible over the fragments the fill wrote. if want_solid_with_edges { pass.set_pipeline(&self.mesh_edge_black_pipeline); - for c in &self.gpu_mesh_batch { + for (mesh_command, c) in self.active_mesh_chunks_indexed() { + if !c.visible { + continue; + } pass.set_bind_group( 1, c.material_bind_group @@ -2719,11 +3394,35 @@ impl Pipeline { c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - pass.draw_indexed(0..c.wire_index_count, 0, 0..1); + if let Some(indirect) = &self.mesh_wire_indirect { + pass.draw_indexed_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndexedIndirectArgs, + >() as u64, + ); + } else { + pass.draw_indexed( + 0..c.wire_index_count, + 0, + 0..c.instance_count, + ); + } } if c.edge_vertex_count != 0 { pass.set_vertex_buffer(0, c.edge_vertex_buffer.slice(..)); - pass.draw(0..c.edge_vertex_count, 0..1); + if let Some(indirect) = &self.mesh_edge_indirect { + pass.draw_indirect( + indirect, + mesh_command as u64 + * std::mem::size_of::< + wgpu::util::DrawIndirectArgs, + >() as u64, + ); + } else { + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); + } } } // DISPSILH silhouettes over the shaded fill. @@ -3164,70 +3863,45 @@ fn round_up_tex(n: u32) -> u32 { ((n.max(1) + GRID - 1) / GRID) * GRID } -/// Project a mesh's `world_aabb` and pick a LOD slot: -/// diagonal > 200 px → 0 (HIGH) -/// diagonal 50–200 px → 1 (MID) -/// diagonal < 50 px → 2 (LOW) -/// When the picked slot is missing from `gpu_meshes[i].lods`, walks down -/// to the next available coarser/finer level so the mesh always renders. -fn pick_mesh_lod( - mesh: &MeshLodGpu, +fn aabb3_offscreen( + aabb: [f32; 6], view_rot: glam::Mat4, eye: glam::DVec3, clip_w: u32, clip_h: u32, -) -> usize { - let diag_px = aabb_diagonal_pixels(mesh.world_aabb, view_rot, eye, clip_w, clip_h); - let target = if diag_px > 200.0 { - 0 - } else if diag_px > 50.0 { - 1 - } else { - 2 - }; - // Walk down to nearest available LOD (some entities won't have all 3). - for level in (0..=target).rev() { - if mesh.lods.get(level).is_some() { - return level; - } - } - 0 -} - -fn aabb_diagonal_pixels( - aabb: [f32; 4], - view_rot: glam::Mat4, - eye: glam::DVec3, - clip_w: u32, - clip_h: u32, -) -> f32 { - let [x0, y0, x1, y1] = aabb; - if !x0.is_finite() || !y0.is_finite() || !x1.is_finite() || !y1.is_finite() { - return f32::INFINITY; +) -> bool { + if aabb.iter().any(|value| !value.is_finite()) { + return false; } let w = clip_w as f32; let h = clip_h as f32; - let corners = [ - view_rot.project_point3((glam::DVec3::new(x0 as f64, y0 as f64, 0.0) - eye).as_vec3()), - view_rot.project_point3((glam::DVec3::new(x1 as f64, y0 as f64, 0.0) - eye).as_vec3()), - view_rot.project_point3((glam::DVec3::new(x0 as f64, y1 as f64, 0.0) - eye).as_vec3()), - view_rot.project_point3((glam::DVec3::new(x1 as f64, y1 as f64, 0.0) - eye).as_vec3()), - ]; let mut min_px = f32::INFINITY; let mut max_px = f32::NEG_INFINITY; let mut min_py = f32::INFINITY; let mut max_py = f32::NEG_INFINITY; - for c in &corners { - let px = (c.x + 1.0) * 0.5 * w; - let py = (1.0 - c.y) * 0.5 * h; - if px < min_px { min_px = px; } - if px > max_px { max_px = px; } - if py < min_py { min_py = py; } - if py > max_py { max_py = py; } + for x in [aabb[0], aabb[3]] { + for y in [aabb[1], aabb[4]] { + for z in [aabb[2], aabb[5]] { + let relative = + (glam::DVec3::new(x as f64, y as f64, z as f64) - eye).as_vec3(); + let clip = view_rot * relative.extend(1.0); + if !clip.is_finite() || clip.w <= f32::EPSILON { + return false; + } + let ndc = clip.truncate() / clip.w; + let px = (ndc.x + 1.0) * 0.5 * w; + let py = (1.0 - ndc.y) * 0.5 * h; + min_px = min_px.min(px); + max_px = max_px.max(px); + min_py = min_py.min(py); + max_py = max_py.max(py); + } + } } - let dx = max_px - min_px; - let dy = max_py - min_py; - (dx * dx + dy * dy).sqrt() + const MARGIN_FRAC: f32 = 0.25; + let mx = w * MARGIN_FRAC; + let my = h * MARGIN_FRAC; + max_px < -mx || min_px > w + mx || max_py < -my || min_py > h + my } /// `true` when the world-XY AABB projects entirely outside the diff --git a/src/scene/view/render.rs b/src/scene/view/render.rs index 187b9cc39..f1266acae 100644 --- a/src/scene/view/render.rs +++ b/src/scene/view/render.rs @@ -262,7 +262,7 @@ impl shader::Primitive for Primitive { inner.cached_epoch = (u64::MAX, u64::MAX, u64::MAX); inner.cached_wire_id = u64::MAX; inner.cached_selection = (u64::MAX, u64::MAX); - inner.cached_mesh_key = (u64::MAX, u64::MAX); + inner.cached_mesh_content_id = u64::MAX; inner.cached_face3d_key = (u64::MAX, false); inner.cached_hatch_source = None; inner.cached_wipeout_source = None; @@ -768,10 +768,21 @@ impl shader::Primitive for Primitive { .as_ref() .map_or(true, |source| !Arc::ptr_eq(source, &vp.meshes)) { - inner.upload_mesh_batch(device, queue, &vp.meshes[..]); + let patched = vp.wire_patch.as_ref().is_some_and(|(previous, patch)| { + *previous == inner.cached_mesh_content_id + && inner.patch_mesh_batch( + device, + queue, + &vp.meshes[..], + &patch.changes, + ) + }); + if !patched { + inner.upload_mesh_batch(device, queue, &vp.meshes[..]); + } inner.cached_mesh_source = Some(Arc::clone(&vp.meshes)); - inner.cached_mesh_batch_epoch = vp.geometry_epoch; } + inner.cached_mesh_content_id = vp.wire_content_id; // Selection / hover highlight overlay — tinted copies of just the // picked solids, rebuilt only when the highlight set (or geometry) // changes. Drawn over the static batch so the base never re-packs. @@ -780,12 +791,7 @@ impl shader::Primitive for Primitive { vp.selection_generation, ); if hl_key != inner.cached_highlight_key { - inner.upload_mesh_highlight( - device, - &vp.meshes[..], - &vp.selected_handles, - vp.hover_handle, - ); + inner.update_mesh_highlight(&vp.selected_handles, vp.hover_handle); inner.cached_highlight_key = hl_key; } // Live overlay (command preview / interim / grip drag) — small and @@ -812,7 +818,13 @@ impl shader::Primitive for Primitive { inner.upload_clip_boundary(device, &vp.clip_boundary_ndc); inner.compute_hatch_lod(queue, view_rot, eye, clip_size.width, clip_size.height); inner.compute_wipeout_lod(view_rot, eye, clip_size.width, clip_size.height); - inner.compute_mesh_lod(view_rot, eye, clip_size.width, clip_size.height); + inner.compute_mesh_lod( + queue, + view_rot, + eye, + clip_size.width, + clip_size.height, + ); #[cfg(not(target_arch = "wasm32"))] { let cull_key = ( diff --git a/src/shaders/mesh.wgsl b/src/shaders/mesh.wgsl index 239c6b854..4d4258b2c 100644 --- a/src/shaders/mesh.wgsl +++ b/src/shaders/mesh.wgsl @@ -55,6 +55,19 @@ struct MaterialMaps { @group(1) @binding(13) var normal_sampler: sampler; @group(1) @binding(14) var maps: MaterialMaps; +struct MeshInstance { + model_row_0: vec4, + model_row_1: vec4, + model_row_2: vec4, + translation_low: vec4, + normal_row_0: vec4, + normal_row_1: vec4, + normal_row_2: vec4, +}; + +@group(1) @binding(15) +var mesh_instances: array; + struct VertexIn { @location(0) position: vec3, @location(1) normal: vec3, @@ -100,12 +113,30 @@ struct VertexOut { }; @vertex -fn vs_main(v: VertexIn) -> VertexOut { +fn vs_main( + v: VertexIn, + @builtin(instance_index) instance_index: u32, +) -> VertexOut { var out: VertexOut; - let rel = (v.position - u.eye_high) + (v.position_low - u.eye_low); + let instance = mesh_instances[instance_index]; + let world_high = vec3( + dot(instance.model_row_0.xyz, v.position) + instance.model_row_0.w, + dot(instance.model_row_1.xyz, v.position) + instance.model_row_1.w, + dot(instance.model_row_2.xyz, v.position) + instance.model_row_2.w, + ); + let world_low = vec3( + dot(instance.model_row_0.xyz, v.position_low), + dot(instance.model_row_1.xyz, v.position_low), + dot(instance.model_row_2.xyz, v.position_low), + ) + instance.translation_low.xyz; + let rel = (world_high - u.eye_high) + (world_low - u.eye_low); out.clip_pos = u.view_rot * vec4(rel, 1.0); out.color = v.color; - out.normal = v.normal; + out.normal = normalize(vec3( + dot(instance.normal_row_0.xyz, v.normal), + dot(instance.normal_row_1.xyz, v.normal), + dot(instance.normal_row_2.xyz, v.normal), + )); out.world_pos = rel; out.material = v.material; out.specular = v.specular; @@ -306,3 +337,19 @@ fn fs_edge(in: VertexOut) -> @location(0) vec4 { fn fs_edge_black(in: VertexOut) -> @location(0) vec4 { return vec4(0.0, 0.0, 0.0, 1.0); } + +@fragment +fn fs_highlight_selected(in: VertexOut) -> @location(0) vec4 { + return vec4( + mix(in.color.rgb, vec3(0.15, 0.55, 1.0), 0.60), + 0.90, + ); +} + +@fragment +fn fs_highlight_hover(in: VertexOut) -> @location(0) vec4 { + return vec4( + mix(in.color.rgb, vec3(0.95, 0.55, 0.10), 0.35), + 0.82, + ); +} diff --git a/src/shaders/mesh_cull.wgsl b/src/shaders/mesh_cull.wgsl new file mode 100644 index 000000000..2a14f01a4 --- /dev/null +++ b/src/shaders/mesh_cull.wgsl @@ -0,0 +1,78 @@ +struct CullUniform { + view_rot: mat4x4, + eye: vec4, + count: vec4, +}; + +struct CullItem { + min: vec4, + max: vec4, + counts: vec4, + info: vec4, +}; + +struct DrawIndexed { + index_count: u32, + instance_count: u32, + first_index: u32, + base_vertex: u32, + first_instance: u32, +}; + +struct Draw { + vertex_count: u32, + instance_count: u32, + first_vertex: u32, + first_instance: u32, +}; + +@group(0) @binding(0) var u: CullUniform; +@group(0) @binding(1) var items: array; +@group(0) @binding(2) var opaque: array; +@group(0) @binding(3) var transparent: array; +@group(0) @binding(4) var wire: array; +@group(0) @binding(5) var edge: array; + +fn chunk_visible(item: CullItem) -> bool { + if (item.info.y == 0u) { + return false; + } + var min_ndc = vec2(1e30, 1e30); + var max_ndc = vec2(-1e30, -1e30); + for (var corner = 0u; corner < 8u; corner++) { + let point = vec3( + select(item.min.x, item.max.x, (corner & 1u) != 0u), + select(item.min.y, item.max.y, (corner & 2u) != 0u), + select(item.min.z, item.max.z, (corner & 4u) != 0u), + ); + let clip = u.view_rot * vec4(point - u.eye.xyz, 1.0); + // A box touching or crossing the eye plane is conservatively retained. + if (clip.w <= 1e-6) { + return true; + } + let ndc = clip.xy / clip.w; + min_ndc = min(min_ndc, ndc); + max_ndc = max(max_ndc, ndc); + } + // 25% viewport margin on each side is 0.5 in NDC. + return !( + max_ndc.x < -1.5 + || min_ndc.x > 1.5 + || max_ndc.y < -1.5 + || min_ndc.y > 1.5 + ); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) id: vec3) { + let index = id.x; + if (index >= u.count.x) { + return; + } + let item = items[index]; + let instances = select(0u, item.info.x, chunk_visible(item)); + opaque[index] = DrawIndexed(item.counts.x, instances, 0u, 0u, 0u); + transparent[index] = DrawIndexed(item.counts.y, instances, 0u, 0u, 0u); + wire[index] = DrawIndexed(item.counts.z, instances, 0u, 0u, 0u); + edge[index] = Draw(item.counts.w, instances, 0u, 0u); +}