diff --git a/Cargo.lock b/Cargo.lock index f4f4422a4..c47e5224c 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -31,6 +31,7 @@ dependencies = [ "js-sys", "log", "lyon_tessellation", + "meshopt", "naga 27.0.3", "ocs_plugin_api", "open", @@ -152,7 +153,7 @@ version = "1.1.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301" dependencies = [ - "memchr 2.8.3", + "memchr", ] [[package]] @@ -1071,7 +1072,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ba5a308b75df32fe02788e748662718f03fde005016435c444eea572398219fd" dependencies = [ "bytes", - "memchr 2.8.3", + "memchr", ] [[package]] @@ -1694,6 +1695,15 @@ version = "0.9.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "98de4bbd547a563b716d8dfa9aad1cb19bfab00f4fa09a6a4ed21dbcf44ce9c4" +[[package]] +name = "float-cmp" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b09cf3155332e944990140d967ff5eceb70df778b34f77d8075db46e4704e6d8" +dependencies = [ + "num-traits", +] + [[package]] name = "float_next_after" version = "1.0.0" @@ -1741,7 +1751,7 @@ version = "0.12.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "54f0d287c53ffd184d04d8677f590f4ac5379785529e5e08b1c8083acdd5c198" dependencies = [ - "memchr 2.8.3", + "memchr", "thiserror 2.0.18", ] @@ -1927,7 +1937,7 @@ dependencies = [ "futures-macro", "futures-sink", "futures-task", - "memchr 2.8.3", + "memchr", "pin-project-lite", "slab", ] @@ -3055,15 +3065,6 @@ dependencies = [ "digest 0.10.7", ] -[[package]] -name = "memchr" -version = "1.0.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "148fab2e51b4f1cfc66da2a7c32981d1d3c083a803978268bb11fe4b86925e7a" -dependencies = [ - "libc", -] - [[package]] name = "memchr" version = "2.8.3" @@ -3088,6 +3089,18 @@ dependencies = [ "autocfg", ] +[[package]] +name = "meshopt" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e01e77ead21976b3a9f01ec1724f766923da74f0726364e2b0f425658935d71a" +dependencies = [ + "bitflags 2.13.1", + "cc", + "float-cmp 0.10.0", + "thiserror 2.0.18", +] + [[package]] name = "mime" version = "0.3.17" @@ -3254,7 +3267,7 @@ version = "0.9.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "418abd1b6d34fbf6cae440dc874771b0525a604428704c76e48b29a5e67b8003" dependencies = [ - "memchr 2.8.3", + "memchr", ] [[package]] @@ -3263,7 +3276,7 @@ version = "7.1.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a" dependencies = [ - "memchr 2.8.3", + "memchr", "minimal-lexical", ] @@ -3273,7 +3286,7 @@ version = "8.0.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "df9761775871bdef83bee530e60050f7e54b1105350d6884eb0fb4f46c2f9405" dependencies = [ - "memchr 2.8.3", + "memchr", ] [[package]] @@ -3283,7 +3296,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0b577e2d69827c4740cba2b52efaad1c4cc7c73042860b199710b3575c68438d" dependencies = [ "bytecount", - "memchr 1.0.2", + "memchr", "nom 8.0.0", ] @@ -4225,7 +4238,7 @@ checksum = "f86ba2052aebccc42cbbb3ed234b8b13ce76f75c3551a303cb2bcffcff12bb14" dependencies = [ "bitflags 2.13.1", "getopts", - "memchr 2.8.3", + "memchr", "pulldown-cmark-escape", "unicase", ] @@ -4286,7 +4299,7 @@ version = "0.39.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "cdcc8dd4e2f670d309a5f0e83fe36dfdc05af317008fea29144da1a2ac858e5e" dependencies = [ - "memchr 2.8.3", + "memchr", ] [[package]] @@ -4532,7 +4545,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d" dependencies = [ "aho-corasick", - "memchr 2.8.3", + "memchr", "regex-automata", "regex-syntax", ] @@ -4544,7 +4557,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8fcfdb36bda0c880c5931cdc7a2bcdc8ba4556847b9d912bca70bc94708711ad" dependencies = [ "aho-corasick", - "memchr 2.8.3", + "memchr", "regex-syntax", ] @@ -4963,7 +4976,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e8014e44b4736ed0538adeecded0fce2a272f22dc9578a7eb6b2d9993c74cfb9" dependencies = [ "itoa", - "memchr 2.8.3", + "memchr", "serde", "serde_core", "zmij", @@ -5260,7 +5273,7 @@ version = "0.1.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6637bab7722d379c8b41ba849228d680cc12d0a45ba1fa2b48f2a30577a06731" dependencies = [ - "float-cmp", + "float-cmp 0.9.0", ] [[package]] @@ -5629,7 +5642,7 @@ version = "3.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b8765b90061cba6c22b5831f675da109ae5561588290f9fa2317adab2714d5a6" dependencies = [ - "memchr 2.8.3", + "memchr", "nom 8.0.0", "petgraph", ] @@ -6699,7 +6712,7 @@ version = "1.0.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "23b97319f7b8343df12cc98938e5c3eb436064524c8d2b4e30a1d3a36eecdf81" dependencies = [ - "memchr 2.8.3", + "memchr", ] [[package]] diff --git a/Cargo.toml b/Cargo.toml index 46bf85949..88c084b3e 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -67,6 +67,7 @@ iced_widget = { git = "https://github.com/iced-rs/iced.git", rev = "23604ff22ab0 [target.'cfg(not(target_arch = "wasm32"))'.dependencies] ocs_plugin_api = { path = "crates/ocs_plugin_api", features = ["host"] } +meshopt = "0.6.2" i18n-embed = { version = "0.16", features = ["desktop-requester"] } rayon = "1" open = "5" diff --git a/src/scene/pipeline/device_capabilities.rs b/src/scene/pipeline/device_capabilities.rs index dfe483c9d..7650bfba0 100644 --- a/src/scene/pipeline/device_capabilities.rs +++ b/src/scene/pipeline/device_capabilities.rs @@ -20,10 +20,6 @@ impl DeviceCapabilities { /// instances, boundaries, families, dashes, and visibility. const HATCH_STORAGE_BINDINGS: u32 = 5; - /// Mesh compute culling reads one item buffer and writes four indirect - /// command buffers. - const MESH_CULL_BINDINGS: u32 = 5; - pub fn detect(device: &wgpu::Device) -> Self { Self::from_limits(&device.limits()) } @@ -36,11 +32,6 @@ impl DeviceCapabilities { } } - /// Mesh instancing needs one read-only storage buffer. - pub fn supports_mesh_storage_instancing(self) -> bool { - self.max_storage_buffers_per_shader_stage >= 1 - } - pub fn supports_wire_storage(self) -> bool { self.max_storage_buffers_per_shader_stage >= Self::WIRE_STORAGE_BINDINGS } @@ -49,10 +40,6 @@ impl DeviceCapabilities { self.max_storage_buffers_per_shader_stage >= Self::HATCH_STORAGE_BINDINGS } - /// WebGL2 reports zero storage bindings and stays on CPU mesh culling. - pub fn supports_mesh_compute_culling(self) -> bool { - self.max_storage_buffers_per_shader_stage >= Self::MESH_CULL_BINDINGS - } } #[cfg(test)] @@ -63,19 +50,15 @@ mod tests { #[test] fn webgl_limits_select_compatibility_paths() { let caps = DeviceCapabilities::from_limits(&wgpu::Limits::downlevel_webgl2_defaults()); - assert!(!caps.supports_mesh_storage_instancing()); assert!(!caps.supports_wire_storage()); assert!(!caps.supports_batched_hatch()); - assert!(!caps.supports_mesh_compute_culling()); } #[test] fn default_limits_select_storage_paths() { let caps = DeviceCapabilities::from_limits(&wgpu::Limits::default()); - assert!(caps.supports_mesh_storage_instancing()); assert!(caps.supports_wire_storage()); assert!(caps.supports_batched_hatch()); - assert!(caps.supports_mesh_compute_culling()); } #[test] @@ -86,8 +69,6 @@ mod tests { max_vertex_attributes: 16, }; assert!(caps.supports_wire_storage()); - assert!(caps.supports_mesh_storage_instancing()); assert!(!caps.supports_batched_hatch()); - assert!(!caps.supports_mesh_compute_culling()); } } diff --git a/src/scene/pipeline/mesh_gpu.rs b/src/scene/pipeline/mesh_gpu.rs index 4f6f4d996..91c519363 100644 --- a/src/scene/pipeline/mesh_gpu.rs +++ b/src/scene/pipeline/mesh_gpu.rs @@ -94,6 +94,41 @@ impl MeshVertex { } } +#[repr(C)] +#[derive(Copy, Clone, Default, bytemuck::Pod, bytemuck::Zeroable)] +pub struct MeshPlainVertex { + pub position: [f32; 3], + pub normal: [f32; 3], + pub position_low: [f32; 3], +} + +impl MeshPlainVertex { + pub fn layout<'a>() -> wgpu::VertexBufferLayout<'a> { + const ATTRS: &[wgpu::VertexAttribute] = &[ + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshPlainVertex, position) as u64, + shader_location: 0, + format: wgpu::VertexFormat::Float32x3, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshPlainVertex, normal) as u64, + shader_location: 1, + format: wgpu::VertexFormat::Float32x3, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshPlainVertex, position_low) as u64, + shader_location: 3, + format: wgpu::VertexFormat::Float32x3, + }, + ]; + wgpu::VertexBufferLayout { + array_stride: std::mem::size_of::() as u64, + step_mode: wgpu::VertexStepMode::Vertex, + attributes: ATTRS, + } + } +} + #[repr(C)] #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub struct MeshEdgeVertex { @@ -240,6 +275,51 @@ impl MeshInstanceGpu { normal_row_2: [n[2] as f32, n[5] as f32, n[8] as f32, 0.0], } } + + pub fn layout<'a>() -> wgpu::VertexBufferLayout<'a> { + const ATTRS: &[wgpu::VertexAttribute] = &[ + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, model_row_0) as u64, + shader_location: 4, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, model_row_1) as u64, + shader_location: 5, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, model_row_2) as u64, + shader_location: 7, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, translation_low) as u64, + shader_location: 8, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, normal_row_0) as u64, + shader_location: 9, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, normal_row_1) as u64, + shader_location: 13, + format: wgpu::VertexFormat::Float32x4, + }, + wgpu::VertexAttribute { + offset: std::mem::offset_of!(MeshInstanceGpu, normal_row_2) as u64, + shader_location: 14, + format: wgpu::VertexFormat::Float32x4, + }, + ]; + wgpu::VertexBufferLayout { + array_stride: std::mem::size_of::() as u64, + step_mode: wgpu::VertexStepMode::Instance, + attributes: ATTRS, + } + } } // ── Batched mesh buffers ────────────────────────────────────────────────── @@ -254,6 +334,7 @@ impl MeshInstanceGpu { pub struct MeshBatchChunk { pub vertex_buffer: wgpu::Buffer, + pub compact_vertices: bool, /// Opaque triangle indices (mesh colour alpha ≈ 1). Drawn with depth write. pub index_buffer: wgpu::Buffer, pub index_count: u32, @@ -274,8 +355,6 @@ pub struct MeshBatchChunk { 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 face_color: [f32; 4], pub material_bind_group: Option, @@ -326,7 +405,7 @@ fn make_chunk( highlight_ranges: &[MeshBatchRange], instances: &[MeshInstanceGpu], handles: &rustc_hash::FxHashSet, - bounds_override: Option<[f32; 6]>, + _bounds_override: Option<[f32; 6]>, material: Option<&crate::scene::model::material_model::MeshMaterial>, face_color: [f32; 4], vertex_buffer_override: Option, @@ -358,6 +437,28 @@ fn make_chunk( queue.write_buffer(&buffer, 0, bytemuck::cast_slice(data)); buffer }; + let compact_vertices = !material_has_textures(material) && wire_indices.is_empty(); + let mk_plain_vertex = |data: &[MeshVertex]| { + if data.is_empty() { + return stubs.vertex.clone(); + } + let compact: Vec = data + .iter() + .map(|vertex| MeshPlainVertex { + position: vertex.position, + normal: vertex.normal, + position_low: vertex.position_low, + }) + .collect(); + let buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("mesh.batch.plain_vbuf"), + size: std::mem::size_of_val(compact.as_slice()) as u64, + usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + queue.write_buffer(&buffer, 0, bytemuck::cast_slice(&compact)); + buffer + }; let mk_edge_vertex = |data: &[MeshEdgeVertex], label: &'static str| { if data.is_empty() { return stubs.vertex.clone(); @@ -377,46 +478,15 @@ fn make_chunk( } 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 { - 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]); - } - } - for vertex in 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: vertex_buffer_override - .unwrap_or_else(|| mk_vertex(verts, "mesh.batch.vbuf")), + vertex_buffer: vertex_buffer_override.unwrap_or_else(|| { + if compact_vertices { + mk_plain_vertex(verts) + } else { + mk_vertex(verts, "mesh.batch.vbuf") + } + }), + compact_vertices, index_buffer: mk_index(indices, "mesh.batch.ibuf"), index_count: indices.len() as u32, transp_index_buffer: mk_index(transp_indices, "mesh.batch.transp_ibuf"), @@ -428,9 +498,9 @@ fn make_chunk( edge_vertex_count: edge_verts.len() as u32, instance_buffer: { let buffer = device.create_buffer(&wgpu::BufferDescriptor { - label: Some("mesh.batch.instances"), + label: Some("mesh.batch.instances"), size: std::mem::size_of_val(instance_data) as u64, - usage: instance_usage | wgpu::BufferUsages::COPY_DST, + usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); queue.write_buffer(&buffer, 0, bytemuck::cast_slice(instance_data)); @@ -439,8 +509,6 @@ fn make_chunk( instance_count: instance_data.len() as u32, highlight_ranges: highlight_ranges.to_vec(), handles: handles.clone(), - world_aabb, - visible: true, material: material.cloned(), face_color, material_bind_group: None, @@ -727,29 +795,9 @@ pub fn create_material_bind_group_from_resources( device: &wgpu::Device, layout: &wgpu::BindGroupLayout, resources: &MeshMaterialResources, - instance_buffer: Option<&wgpu::Buffer>, material: Option<&crate::scene::model::material_model::MeshMaterial>, face_color: [f32; 4], ) -> wgpu::BindGroup { - 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 - } - }; let (material_params, specular, ambient, advanced, flags) = material_vertex_params(material); let surface = MeshSurfaceParams { @@ -829,10 +877,6 @@ pub fn create_material_bind_group_from_resources( binding: 14, resource: resources.params_buffer.as_entire_binding(), }, - wgpu::BindGroupEntry { - binding: 15, - resource: instance_buffer.as_entire_binding(), - }, wgpu::BindGroupEntry { binding: 16, resource: surface_buffer.as_entire_binding(), @@ -846,7 +890,6 @@ 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 resources = create_material_resources(device, queue, material); let color = material.map_or([0.8, 0.8, 0.8, 1.0], |material| material.diffuse); @@ -854,7 +897,6 @@ pub fn create_material_bind_group( device, layout, &resources, - instance_buffer, material, color, ) @@ -879,7 +921,6 @@ pub fn upload_chunk_material_bind_groups( device, layout, &resources, - Some(&chunk.instance_buffer), chunk.material.as_ref(), chunk.face_color, )); @@ -894,7 +935,6 @@ pub fn upload_chunk_material_bind_groups( device, layout, resources, - Some(&chunk.instance_buffer), chunk.material.as_ref(), chunk.face_color, )); @@ -1025,6 +1065,18 @@ fn index_hash(indices: &[u32]) -> u64 { hasher.finish() } +fn optimize_triangle_indices(indices: &[u32], vertex_count: usize) -> std::sync::Arc<[u32]> { + #[cfg(not(target_arch = "wasm32"))] + { + return meshopt::optimize::optimize_vertex_cache(indices, vertex_count).into(); + } + #[cfg(target_arch = "wasm32")] + { + let _ = vertex_count; + indices.into() + } +} + fn visual_style_partition_key( style: Option<&crate::scene::model::visual_style_model::MeshVisualStyle>, ) -> [u64; 12] { @@ -1061,6 +1113,7 @@ struct InstancedBuildProfile { struct InstancedVertexKey { source: u64, material: u64, + compact: bool, } #[derive(Default)] @@ -1102,6 +1155,8 @@ fn build_instanced_chunk( let material = first.material; let color = first.color; let source_handle = source.handle.value(); + let compact_vertices = !material_has_textures(material) + && (!first.include_edges || !source.edge_verts.is_empty()); let material_identity = if material_has_textures(material) { material.map_or(0, |material| { material @@ -1114,6 +1169,7 @@ fn build_instanced_chunk( let vertex_key = InstancedVertexKey { source: source_handle, material: material_identity, + compact: compact_vertices, }; let shared_vertex_buffer = buffers.vertices.get(&vertex_key).cloned(); let has_normals = mesh.normals.len() == mesh.verts.len(); @@ -1574,7 +1630,7 @@ pub fn build_mesh_batch_filtered( let (shared_indices, shared_hash) = source_indices .entry(source_identity) .or_insert_with(|| { - let indices = std::sync::Arc::<[u32]>::from(mesh.indices.as_slice()); + let indices = optimize_triangle_indices(&mesh.indices, mesh.verts.len()); let hash = index_hash(indices.as_ref()); (indices, hash) }) @@ -1655,7 +1711,7 @@ pub fn build_mesh_batch_filtered( let parts: Vec<_> = groups .into_values() .map(|(material, color, indices)| { - let indices = std::sync::Arc::<[u32]>::from(indices); + let indices = optimize_triangle_indices(&indices, mesh.verts.len()); let index_hash = index_hash(indices.as_ref()); CachedFacePart { material, @@ -1711,7 +1767,6 @@ pub fn build_mesh_batch_filtered( mesh_spatial_key(part.set, spatial_bounds), ) }); - let storage_instancing = device.limits().max_storage_buffers_per_shader_stage > 0; let mut instance_groups: std::collections::BTreeMap< InstanceGroupKey, Vec>, @@ -2173,59 +2228,24 @@ pub fn build_mesh_batch_filtered( let mut instanced_profile = InstancedBuildProfile::default(); let mut instanced_buffers = InstancedBufferCache::default(); for parts in instance_groups.values() { - let Some(first) = parts.first() else { - continue; - }; - let source_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::()) - }), - ) - .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 = if storage_instancing { - parts.len().div_ceil(max_clusters).max(64) - } else { - 1 - }; - for cluster in parts.chunks(cluster_len) { - if let Some((chunk, triangles, profile)) = - build_instanced_chunk( - device, - queue, - &stubs, - &mut instanced_buffers, - cluster, - perf_started.is_some(), - ) - { - total_tris += triangles; - chunks.push(chunk); - instanced_profile.vertices += profile.vertices; - instanced_profile.edges += profile.edges; - instanced_profile.instances += profile.instances; - instanced_profile.upload += profile.upload; - instanced_profile.vertex_count += profile.vertex_count; - instanced_profile.edge_count += profile.edge_count; - instanced_profile.index_count += profile.index_count; - instanced_profile.chunk_count += profile.chunk_count; - } + if let Some((chunk, triangles, profile)) = build_instanced_chunk( + device, + queue, + &stubs, + &mut instanced_buffers, + parts, + perf_started.is_some(), + ) { + total_tris += triangles; + chunks.push(chunk); + instanced_profile.vertices += profile.vertices; + instanced_profile.edges += profile.edges; + instanced_profile.instances += profile.instances; + instanced_profile.upload += profile.upload; + instanced_profile.vertex_count += profile.vertex_count; + instanced_profile.edge_count += profile.edge_count; + instanced_profile.index_count += profile.index_count; + instanced_profile.chunk_count += profile.chunk_count; } } if let Some(started) = perf_started { diff --git a/src/scene/pipeline/mod.rs b/src/scene/pipeline/mod.rs index a5307afc6..c8ab2030c 100644 --- a/src/scene/pipeline/mod.rs +++ b/src/scene/pipeline/mod.rs @@ -35,6 +35,12 @@ struct SilhouetteChunk { vertex_count: u32, instance_count: u32, } + +struct SilhouetteSourceGroup { + color: [f32; 4], + sources: Vec, + instance_buffers: Vec<(wgpu::Buffer, u32)>, +} use device_capabilities::DeviceCapabilities; /// MSAA sample count for the main drawing pipelines. @@ -64,26 +70,10 @@ struct MeshHighlightDraw { 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 { background_pipeline: wgpu::RenderPipeline, shadow_pipeline: wgpu::RenderPipeline, + shadow_plain_pipeline: wgpu::RenderPipeline, wire_pipeline: wgpu::RenderPipeline, block_wire_pipeline: wgpu::RenderPipeline, /// Stamps clip-boundary polygons into the stencil buffer (viewports + XCLIP). @@ -114,10 +104,14 @@ pub struct Pipeline { text_highlight_pipeline: wgpu::RenderPipeline, block_text_highlight_pipeline: wgpu::RenderPipeline, mesh_pipeline: wgpu::RenderPipeline, + mesh_plain_pipeline: wgpu::RenderPipeline, /// Depth-write-disabled variant of `mesh_pipeline` for non-opaque solids. mesh_transparent_pipeline: wgpu::RenderPipeline, + mesh_plain_transparent_pipeline: wgpu::RenderPipeline, mesh_selected_pipeline: wgpu::RenderPipeline, + mesh_plain_selected_pipeline: wgpu::RenderPipeline, mesh_hover_pipeline: wgpu::RenderPipeline, + mesh_plain_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 @@ -131,18 +125,9 @@ pub struct Pipeline { /// writes, writes depth). Used in HiddenLine mode so 3D solids /// occlude wires behind them without painting visible pixels. mesh_depth_pipeline: wgpu::RenderPipeline, + mesh_plain_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, block_face3d_pipeline: wgpu::RenderPipeline, /// Depth-only variant of the face3d pipeline (no color writes, @@ -189,6 +174,8 @@ pub struct Pipeline { text_preview_vcount: u32, /// Per-frame silhouette line lists from the kernel mesh and current view. silhouette_chunks: Vec, + silhouette_source_key: (usize, usize, u64), + silhouette_source_groups: Vec, /// Last requested render size (the full viewport rect, in pixels). The /// geometry passes render at this size; the blit UV is scaled by /// `depth_texture_size / alloc_size` so it samples only the filled region. @@ -244,14 +231,10 @@ pub struct Pipeline { /// Last content/camera/viewport tuple used to derive visible instance /// ranges from the resident arena. pub(crate) wire_cull_key: (u64, u64, u32, u32), - /// View/source keys for CPU visibility passes. A plain entity edit changes - /// the scene render signature, but it must not rescan every unchanged hatch, - /// wipeout, or mesh AABB when the camera and corresponding source stayed - /// identical. + /// View/source keys for CPU visibility passes. pub(crate) hatch_lod_key: (usize, u64, u32, u32, bool), pub(crate) wipeout_lod_key: (usize, u64, u32, u32, bool), - pub(crate) mesh_lod_key: (usize, u64, u32, u32), - pub(crate) silhouette_key: (usize, u64, u64, bool), + pub(crate) silhouette_key: (usize, u64, [u32; 3], bool), /// This content viewport's non-rectangular clip boundary as a triangle-fan /// vertex buffer in the render target's normalized device coords (`None` = /// rectangular / unclipped, where the viewport's own render rectangle does @@ -1070,154 +1053,29 @@ impl Pipeline { #[cfg(not(target_arch = "wasm32"))] if std::env::var_os("RUST_LOG").is_some() { eprintln!( - "renderer pipelines: wire={} hatch={} mesh={} compute-cull={} (storage buffers/stage: {})", + "renderer pipelines: wire={} hatch={} mesh={} (storage buffers/stage: {})", if wire_mode.uses_storage() { "storage" } else { "packed" }, hatch_gpu.backend_name(), - if device_caps.supports_mesh_storage_instancing() { "storage" } else { "uniform" }, - if device_caps.supports_mesh_compute_culling() { "gpu" } else { "cpu" }, + "vertex", device.limits().max_storage_buffers_per_shader_stage ); } #[cfg(target_arch = "wasm32")] log::info!( - "renderer pipelines: wire={} hatch={} mesh={} compute-cull={} (storage buffers/stage: {})", + "renderer pipelines: wire={} hatch={} mesh={} (storage buffers/stage: {})", if wire_mode.uses_storage() { "storage" } else { "packed" }, hatch_gpu.backend_name(), - if device_caps.supports_mesh_storage_instancing() { "storage" } else { "uniform" }, - if device_caps.supports_mesh_compute_culling() { "gpu" } else { "cpu" }, + "vertex", device.limits().max_storage_buffers_per_shader_stage ); // ── Mesh pipeline ────────────────────────────────────────────────── - let mesh_storage_instancing = device_caps.supports_mesh_storage_instancing(); - 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(mesh_source), + source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!( + "../../shaders/mesh.wgsl" + ))), }); - let (mesh_cull_bgl, mesh_cull_pipeline, mesh_cull_uniform) = - if device_caps.supports_mesh_compute_culling() { - 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].map(Some), - immediate_size: 0, - }); - 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 { label: Some("mesh.material.bgl"), @@ -1344,20 +1202,6 @@ 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, - }, wgpu::BindGroupLayoutEntry { binding: 16, visibility: wgpu::ShaderStages::FRAGMENT, @@ -1375,7 +1219,6 @@ impl Pipeline { queue, &mesh_material_bgl, None, - None, ); let mesh_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { @@ -1384,18 +1227,11 @@ impl Pipeline { immediate_size: 0, }); - let shadow_source = include_str!("../../shaders/shadow.wgsl"); - let shadow_source = if mesh_storage_instancing { - std::borrow::Cow::Borrowed(shadow_source) - } else { - std::borrow::Cow::Owned(shadow_source.replace( - "var mesh_instances: array;", - "var mesh_instances: array;", - )) - }; let shadow_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("shadow.shader"), - source: wgpu::ShaderSource::Wgsl(shadow_source), + source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!( + "../../shaders/shadow.wgsl" + ))), }); let shadow_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("shadow.pipeline_layout"), @@ -1408,7 +1244,10 @@ impl Pipeline { vertex: wgpu::VertexState { module: &shadow_shader, entry_point: Some("vs_main"), - buffers: &[mesh_gpu::MeshVertex::layout()], + buffers: &[ + mesh_gpu::MeshVertex::layout(), + mesh_gpu::MeshInstanceGpu::layout(), + ], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1432,6 +1271,40 @@ impl Pipeline { multiview_mask: None, cache: None, }); + let shadow_plain_pipeline = + device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some("shadow.plain.pipeline"), + layout: Some(&shadow_layout), + vertex: wgpu::VertexState { + module: &shadow_shader, + entry_point: Some("vs_main"), + buffers: &[ + mesh_gpu::MeshPlainVertex::layout(), + mesh_gpu::MeshInstanceGpu::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::Depth32Float, + depth_write_enabled: Some(true), + depth_compare: Some(wgpu::CompareFunction::LessEqual), + stencil: wgpu::StencilState::default(), + bias: wgpu::DepthBiasState { + constant: 2, + slope_scale: 2.0, + clamp: 0.0, + }, + }), + multisample: wgpu::MultisampleState::default(), + fragment: None, + multiview_mask: None, + cache: None, + }); let mesh_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { label: Some("mesh.pipeline"), @@ -1439,7 +1312,10 @@ impl Pipeline { vertex: wgpu::VertexState { module: &mesh_shader, entry_point: Some("vs_main"), - buffers: &[mesh_gpu::MeshVertex::layout()], + buffers: &[ + mesh_gpu::MeshVertex::layout(), + mesh_gpu::MeshInstanceGpu::layout(), + ], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1476,6 +1352,53 @@ impl Pipeline { multiview_mask: None, cache: None, }); + let mesh_plain_pipeline = + device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some("mesh.plain.pipeline"), + layout: Some(&mesh_layout), + vertex: wgpu::VertexState { + module: &mesh_shader, + entry_point: Some("vs_main_plain"), + buffers: &[ + mesh_gpu::MeshPlainVertex::layout(), + mesh_gpu::MeshInstanceGpu::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: Some(true), + depth_compare: Some(wgpu::CompareFunction::LessEqual), + stencil: content_stencil.clone(), + bias: wgpu::DepthBiasState { + constant: 1, + slope_scale: 1.0, + clamp: 0.0, + }, + }), + 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_mask: None, + cache: None, + }); // Transparent variant — identical to `mesh_pipeline` but with depth // writes disabled. Non-opaque solids are drawn after the opaque fills @@ -1488,7 +1411,57 @@ impl Pipeline { vertex: wgpu::VertexState { module: &mesh_shader, entry_point: Some("vs_main"), - buffers: &[mesh_gpu::MeshVertex::layout()], + buffers: &[ + mesh_gpu::MeshVertex::layout(), + mesh_gpu::MeshInstanceGpu::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: Some(false), + depth_compare: Some(wgpu::CompareFunction::LessEqual), + stencil: content_stencil.clone(), + bias: wgpu::DepthBiasState { + constant: 1, + slope_scale: 1.0, + clamp: 0.0, + }, + }), + 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_mask: None, + cache: None, + }); + let mesh_plain_transparent_pipeline = + device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some("mesh.plain.transparent.pipeline"), + layout: Some(&mesh_layout), + vertex: wgpu::VertexState { + module: &mesh_shader, + entry_point: Some("vs_main_plain"), + buffers: &[ + mesh_gpu::MeshPlainVertex::layout(), + mesh_gpu::MeshInstanceGpu::layout(), + ], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1529,14 +1502,20 @@ impl Pipeline { // 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| { + |label: &'static str, + vertex_entry: &'static str, + vertex_layout: wgpu::VertexBufferLayout<'static>, + 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()], + entry_point: Some(vertex_entry), + buffers: &[ + vertex_layout, + mesh_gpu::MeshInstanceGpu::layout(), + ], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1572,10 +1551,28 @@ impl Pipeline { }; let mesh_selected_pipeline = make_mesh_highlight_pipeline( "mesh.highlight.selected.pipeline", + "vs_main", + mesh_gpu::MeshVertex::layout(), "fs_highlight_selected", ); - let mesh_hover_pipeline = - make_mesh_highlight_pipeline("mesh.highlight.hover.pipeline", "fs_highlight_hover"); + let mesh_plain_selected_pipeline = make_mesh_highlight_pipeline( + "mesh.plain.highlight.selected.pipeline", + "vs_main_plain", + mesh_gpu::MeshPlainVertex::layout(), + "fs_highlight_selected", + ); + let mesh_hover_pipeline = make_mesh_highlight_pipeline( + "mesh.highlight.hover.pipeline", + "vs_main", + mesh_gpu::MeshVertex::layout(), + "fs_highlight_hover", + ); + let mesh_plain_hover_pipeline = make_mesh_highlight_pipeline( + "mesh.plain.highlight.hover.pipeline", + "vs_main_plain", + mesh_gpu::MeshPlainVertex::layout(), + "fs_highlight_hover", + ); // Wireframe variant — same shader / vertex layout / depth state, // only the input topology changes (LineList) and back-face @@ -1595,7 +1592,7 @@ impl Pipeline { vertex: wgpu::VertexState { module: &mesh_shader, entry_point: Some("vs_edge"), - buffers: &[vertex_layout], + buffers: &[vertex_layout, mesh_gpu::MeshInstanceGpu::layout()], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1710,7 +1707,10 @@ impl Pipeline { vertex: wgpu::VertexState { module: &mesh_shader, entry_point: Some("vs_main"), - buffers: &[mesh_gpu::MeshVertex::layout()], + buffers: &[ + mesh_gpu::MeshVertex::layout(), + mesh_gpu::MeshInstanceGpu::layout(), + ], compilation_options: wgpu::PipelineCompilationOptions::default(), }, primitive: wgpu::PrimitiveState { @@ -1747,6 +1747,53 @@ impl Pipeline { multiview_mask: None, cache: None, }); + let mesh_plain_depth_pipeline = + device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some("mesh.plain.depth.pipeline"), + layout: Some(&mesh_layout), + vertex: wgpu::VertexState { + module: &mesh_shader, + entry_point: Some("vs_main_plain"), + buffers: &[ + mesh_gpu::MeshPlainVertex::layout(), + mesh_gpu::MeshInstanceGpu::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: Some(true), + depth_compare: Some(wgpu::CompareFunction::LessEqual), + stencil: content_stencil.clone(), + bias: wgpu::DepthBiasState { + constant: 1, + slope_scale: 1.0, + clamp: 0.0, + }, + }), + 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: None, + write_mask: wgpu::ColorWrites::empty(), + })], + compilation_options: wgpu::PipelineCompilationOptions::default(), + }), + multiview_mask: None, + cache: None, + }); // ── Face3D pipeline ──────────────────────────────────────────────── let face3d_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { @@ -2159,6 +2206,7 @@ impl Pipeline { Self { background_pipeline, shadow_pipeline, + shadow_plain_pipeline, wire_pipeline, block_wire_pipeline, clip_mask_pipeline, @@ -2186,27 +2234,24 @@ impl Pipeline { text_preview_vbuf: None, text_preview_vcount: 0, silhouette_chunks: Vec::new(), + silhouette_source_key: (usize::MAX, usize::MAX, u64::MAX), + silhouette_source_groups: Vec::new(), mesh_pipeline, + mesh_plain_pipeline, mesh_transparent_pipeline, + mesh_plain_transparent_pipeline, mesh_selected_pipeline, + mesh_plain_selected_pipeline, mesh_hover_pipeline, + mesh_plain_hover_pipeline, mesh_wireframe_pipeline, mesh_edge_black_pipeline, silhouette_pipeline, silhouette_black_pipeline, mesh_depth_pipeline, + mesh_plain_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, block_face3d_pipeline, face3d_depth_pipeline, @@ -2250,8 +2295,7 @@ impl Pipeline { wire_cull_key: (u64::MAX, u64::MAX, 0, 0), hatch_lod_key: (usize::MAX, u64::MAX, 0, 0, false), wipeout_lod_key: (usize::MAX, u64::MAX, 0, 0, false), - mesh_lod_key: (usize::MAX, u64::MAX, 0, 0), - silhouette_key: (usize::MAX, u64::MAX, u64::MAX, false), + silhouette_key: (usize::MAX, u64::MAX, [u32::MAX; 3], false), clip_boundary: None, gpu_selected_wires: vec![], gpu_selected_block_wires: vec![], @@ -2674,6 +2718,7 @@ impl Pipeline { &mut self, device: &wgpu::Device, sets: &[crate::scene::model::mesh_model::MeshLodSet], + content_id: u64, view_dir: glam::Vec3, ) { let perf_started = crate::perf::enabled().then(iced::time::Instant::now); @@ -2691,83 +2736,161 @@ impl Pipeline { self.silhouette_chunks.clear(); return; } - let mut slots = rustc_hash::FxHashMap::default(); - let mut groups: Vec> = Vec::new(); - for (index, set) in sets.iter().enumerate() { - let color = set.display_color().unwrap_or([0.0, 0.0, 0.0, 1.0]); - let key = match (&set.instance_source, set.instance_transform) { - (Some(source), Some(transform)) => { - let matrix = &transform.matrix.m; - ( - true, - source.handle.value(), - [ - matrix[0][0].to_bits(), matrix[0][1].to_bits(), matrix[0][2].to_bits(), - matrix[1][0].to_bits(), matrix[1][1].to_bits(), matrix[1][2].to_bits(), - matrix[2][0].to_bits(), matrix[2][1].to_bits(), matrix[2][2].to_bits(), - ], - color.map(f32::to_bits), - ) - } - _ => (false, index as u64, [0; 9], color.map(f32::to_bits)), - }; - let slot = *slots.entry(key).or_insert_with(|| { - let slot = groups.len(); - groups.push(Vec::new()); - slot - }); - groups[slot].push(set); + let source_key = (sets.as_ptr() as usize, sets.len(), content_id); + if self.silhouette_source_key != source_key { + let mut slots = rustc_hash::FxHashMap::default(); + let mut groups: Vec> = Vec::new(); + for (index, set) in sets.iter().enumerate() { + let color = set.display_color().unwrap_or([0.0, 0.0, 0.0, 1.0]); + let key = match (&set.instance_source, set.instance_transform) { + (Some(source), Some(transform)) => { + let matrix = &transform.matrix.m; + ( + true, + source.handle.value(), + [ + matrix[0][0].to_bits(), matrix[0][1].to_bits(), + matrix[0][2].to_bits(), matrix[1][0].to_bits(), + matrix[1][1].to_bits(), matrix[1][2].to_bits(), + matrix[2][0].to_bits(), matrix[2][1].to_bits(), + matrix[2][2].to_bits(), + ], + color.map(f32::to_bits), + ) + } + _ => (false, index as u64, [0; 9], color.map(f32::to_bits)), + }; + let slot = *slots.entry(key).or_insert_with(|| { + let slot = groups.len(); + groups.push(Vec::new()); + slot + }); + groups[slot].push(set); + } + + let prepare_group = + |group: Vec<&crate::scene::model::mesh_model::MeshLodSet>| { + let source = *group.first()?; + let translation = + |set: &crate::scene::model::mesh_model::MeshLodSet| { + set.instance_transform.map_or([0.0; 3], |transform| { + let matrix = &transform.matrix.m; + [matrix[0][3], matrix[1][3], matrix[2][3]] + }) + }; + let base = translation(source); + let instances: Vec = group + .iter() + .map(|set| { + let placement = translation(set); + let delta = [ + placement[0] - base[0], + placement[1] - base[1], + placement[2] - base[2], + ]; + let high = delta.map(|value| value as f32); + SilhouetteInstance { + translation: high, + translation_low: [ + (delta[0] - high[0] as f64) as f32, + (delta[1] - high[1] as f64) as f32, + (delta[2] - high[2] as f64) as f32, + ], + } + }) + .collect(); + let max_instances = + chunk_bytes / std::mem::size_of::(); + let instance_buffers = instances + .chunks(max_instances.max(1)) + .map(|instances| { + ( + device.create_buffer_init( + &wgpu::util::BufferInitDescriptor { + label: Some("mesh.silhouette.instances"), + contents: bytemuck::cast_slice(instances), + usage: wgpu::BufferUsages::VERTEX, + }, + ), + instances.len() as u32, + ) + }) + .collect(); + let generators = source.instance_source.as_ref().map_or( + source.curved_gens.as_slice(), + |instance| instance.curved_gens.as_slice(), + ); + let mut sources = Vec::with_capacity(generators.len()); + for generator in generators { + if let Some(transform) = source.instance_transform { + let origin = + transform.apply(acadrust::types::Vector3::ZERO); + let vectors = [ + transform.apply_rotation( + acadrust::types::Vector3::UNIT_X, + ), + transform.apply_rotation( + acadrust::types::Vector3::UNIT_Y, + ), + transform.apply_rotation( + acadrust::types::Vector3::UNIT_Z, + ), + ]; + if let Some(transformed) = + cadkernel::brep::mesh::transform_silhouette_affine( + &generator.source, + vectors.map(|vector| [vector.x, vector.y, vector.z]), + [origin.x, origin.y, origin.z], + ) + { + sources.push(transformed); + } + } else { + sources.push(generator.source.clone()); + } + } + Some(SilhouetteSourceGroup { + color: source + .display_color() + .unwrap_or([0.0, 0.0, 0.0, 1.0]), + sources, + instance_buffers, + }) + }; + self.silhouette_source_groups = + groups.into_iter().filter_map(prepare_group).collect(); + self.silhouette_source_key = source_key; } - let source_count = groups.len(); + let source_count = self.silhouette_source_groups.len(); + let compute_group = |group: &SilhouetteSourceGroup| { + let mut points = Vec::new(); + for source in &group.sources { + points.extend(cadkernel::brep::mesh::silhouette( + source, + [view.x, view.y, view.z], + )); + } + points + }; + #[cfg(not(target_arch = "wasm32"))] + let points: Vec<_> = { + use rayon::prelude::*; + self.silhouette_source_groups + .par_iter() + .map(compute_group) + .collect() + }; + #[cfg(target_arch = "wasm32")] + let points: Vec<_> = self + .silhouette_source_groups + .iter() + .map(compute_group) + .collect(); let mut source_vertex_count = 0usize; let mut chunks = Vec::new(); - for group in groups { - let Some(&source) = group.first() else { - continue; - }; - let translation = |set: &crate::scene::model::mesh_model::MeshLodSet| { - set.instance_transform.map_or([0.0; 3], |transform| { - let matrix = &transform.matrix.m; - [matrix[0][3], matrix[1][3], matrix[2][3]] - }) - }; - let base = translation(source); - let instances: Vec = group - .iter() - .map(|set| { - let placement = translation(set); - let delta = [ - placement[0] - base[0], - placement[1] - base[1], - placement[2] - base[2], - ]; - let high = delta.map(|value| value as f32); - SilhouetteInstance { - translation: high, - translation_low: [ - (delta[0] - high[0] as f64) as f32, - (delta[1] - high[1] as f64) as f32, - (delta[2] - high[2] as f64) as f32, - ], - } - }) - .collect(); - let max_instances = chunk_bytes / std::mem::size_of::(); - let instance_buffers: Vec<(wgpu::Buffer, u32)> = instances - .chunks(max_instances.max(1)) - .map(|instances| { - ( - device.create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("mesh.silhouette.instances"), - contents: bytemuck::cast_slice(instances), - usage: wgpu::BufferUsages::VERTEX, - }), - instances.len() as u32, - ) - }) - .collect(); - let color = source.display_color().unwrap_or([0.0, 0.0, 0.0, 1.0]); + for (group, points) in self.silhouette_source_groups.iter().zip(points) { + let color = group.color; let mk = |w: glam::DVec3| -> SilhouetteVertex { let (hx, hy, hz) = (w.x as f32, w.y as f32, w.z as f32); SilhouetteVertex { @@ -2796,7 +2919,7 @@ impl Pipeline { usage: wgpu::BufferUsages::VERTEX, }, ); - for (instance_buffer, instance_count) in &instance_buffers { + for (instance_buffer, instance_count) in &group.instance_buffers { chunks.push(SilhouetteChunk { vertex_buffer: vertex_buffer.clone(), instance_buffer: instance_buffer.clone(), @@ -2805,44 +2928,12 @@ impl Pipeline { }); } }; - let generators = source - .instance_source - .as_ref() - .map_or(source.curved_gens.as_slice(), |instance| { - instance.curved_gens.as_slice() - }); - for generator in generators { - let transformed; - let silhouette_source = if let Some(transform) = source.instance_transform { - let origin = transform.apply(acadrust::types::Vector3::ZERO); - let vectors = [ - transform.apply_rotation(acadrust::types::Vector3::UNIT_X), - transform.apply_rotation(acadrust::types::Vector3::UNIT_Y), - transform.apply_rotation(acadrust::types::Vector3::UNIT_Z), - ]; - transformed = cadkernel::brep::mesh::transform_silhouette_affine( - &generator.source, - vectors.map(|vector| [vector.x, vector.y, vector.z]), - [origin.x, origin.y, origin.z], - ); - let Some(transformed) = transformed.as_ref() else { - continue; - }; - transformed - } else { - &generator.source - }; - let points = cadkernel::brep::mesh::silhouette( - silhouette_source, - [view.x, view.y, view.z], - ); - source_vertex_count += points.len(); - for point in points { - verts.push(mk(glam::DVec3::from_array(point))); - if verts.len() == max_vertices { - push_chunk(&verts, &mut chunks); - verts.clear(); - } + source_vertex_count += points.len(); + for point in points { + verts.push(mk(glam::DVec3::from_array(point))); + if verts.len() == max_vertices { + push_chunk(&verts, &mut chunks); + verts.clear(); } } push_chunk(&verts, &mut chunks); @@ -3040,164 +3131,32 @@ impl Pipeline { self.mesh_dynamic_handles.clear(); self.gpu_mesh_batch = chunks; let indexed_at = iced::time::Instant::now(); - self.rebuild_mesh_cull_resources(device); - let culled_at = iced::time::Instant::now(); if crate::perf::enabled() { let instances: u64 = self .gpu_mesh_batch .iter() .map(|chunk| chunk.instance_count as u64) .sum(); + let compact_chunks = self + .gpu_mesh_batch + .iter() + .filter(|chunk| chunk.compact_vertices) + .count(); crate::perf_record!( - "[perf] mesh-batch {:>7.1}ms build={:.1} material={:.1} index={:.1} cull={:.1} sets={} chunks={} instances={} triangles={}", + "[perf] mesh-batch {:>7.1}ms build={:.1} material={:.1} index={:.1} sets={} chunks={} compact={} instances={} triangles={}", started.elapsed().as_secs_f64() * 1000.0, built_at.duration_since(started).as_secs_f64() * 1000.0, materials_at.duration_since(built_at).as_secs_f64() * 1000.0, indexed_at.duration_since(materials_at).as_secs_f64() * 1000.0, - culled_at.duration_since(indexed_at).as_secs_f64() * 1000.0, meshes.len(), self.gpu_mesh_batch.len(), + compact_chunks, instances, triangles, ); } } - 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() { @@ -3301,7 +3260,6 @@ impl Pipeline { &mut chunks, ); self.gpu_mesh_dynamic = chunks; - self.rebuild_mesh_cull_resources(device); self.rebuild_mesh_range_map(); if crate::perf::enabled() { crate::perf_record!( @@ -3359,47 +3317,6 @@ impl Pipeline { ) } - /// 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( &mut self, device: &wgpu::Device, @@ -3681,12 +3598,13 @@ impl Pipeline { occlusion_query_set: None, multiview_mask: None, }); - pass.set_pipeline(&self.shadow_pipeline); pass.set_bind_group(0, &self.shadow_uniform_bind_group, &[]); for (_, chunk) in self.active_mesh_chunks_indexed() { - if !chunk.visible { - continue; - } + pass.set_pipeline(if chunk.compact_vertices { + &self.shadow_plain_pipeline + } else { + &self.shadow_pipeline + }); pass.set_bind_group( 1, chunk @@ -3696,6 +3614,7 @@ impl Pipeline { &[], ); pass.set_vertex_buffer(0, chunk.vertex_buffer.slice(..)); + pass.set_vertex_buffer(1, chunk.instance_buffer.slice(..)); if chunk.index_count != 0 { pass.set_index_buffer(chunk.index_buffer.slice(..), wgpu::IndexFormat::Uint32); pass.draw_indexed(0..chunk.index_count, 0, 0..chunk.instance_count); @@ -3814,21 +3733,6 @@ 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() || !self.gpu_mesh_dynamic.is_empty() { let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { @@ -3869,17 +3773,16 @@ impl Pipeline { // top of the shaded surface. let want_solid_with_edges = !hidden_line && !mesh_wireframe && show_3d_edges; // Each path now binds a chunk's buffers and draws the whole chunk in - // one call — a handful of draws total instead of one per solid. No - // per-solid LOD / frustum cull in the batched path (the batch is - // resident in full); that is reintroduced separately if needed. + // one call — a handful of draws total instead of one per solid. if hidden_line { // 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 (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { + pass.set_pipeline(if c.compact_vertices { + &self.mesh_plain_depth_pipeline + } else { + &self.mesh_depth_pipeline + }); pass.set_bind_group( 1, c.material_bind_group @@ -3888,51 +3791,29 @@ impl Pipeline { &[], ); pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); if c.index_count != 0 { pass.set_index_buffer( c.index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - 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); - } + 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, ); - 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.draw_indexed( + 0..c.transp_index_count, + 0, + 0..c.instance_count, + ); } } pass.set_pipeline(&self.mesh_wireframe_pipeline); pass.set_bind_group(1, &self.mesh_default_material_bind_group, &[]); - for (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { pass.set_bind_group( 1, c.material_bind_group @@ -3940,6 +3821,7 @@ impl Pipeline { .unwrap_or(&self.mesh_default_material_bind_group), &[], ); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); // Plain-mesh triangulation edges. if c.wire_index_count != 0 { pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); @@ -3947,36 +3829,16 @@ impl Pipeline { c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - 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, - ); - } + 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(..)); - 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); - } + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); } } pass.set_pipeline(&self.silhouette_black_pipeline); @@ -3989,10 +3851,7 @@ impl Pipeline { if mesh_wireframe { pass.set_pipeline(&self.mesh_wireframe_pipeline); pass.set_bind_group(1, &self.mesh_default_material_bind_group, &[]); - for (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { pass.set_bind_group( 1, c.material_bind_group @@ -4000,41 +3859,22 @@ impl Pipeline { .unwrap_or(&self.mesh_default_material_bind_group), &[], ); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); if c.wire_index_count != 0 { pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); pass.set_index_buffer( c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - 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, - ); - } + 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(..)); - 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); - } + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); } } pass.set_pipeline(&self.silhouette_pipeline); @@ -4045,11 +3885,12 @@ impl Pipeline { } } else { // Opaque fills first (they write depth). - pass.set_pipeline(&self.mesh_pipeline); - for (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { + pass.set_pipeline(if c.compact_vertices { + &self.mesh_plain_pipeline + } else { + &self.mesh_pipeline + }); pass.set_bind_group( 1, c.material_bind_group @@ -4061,27 +3902,19 @@ impl Pipeline { continue; } pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); pass.set_index_buffer(c.index_buffer.slice(..), wgpu::IndexFormat::Uint32); - 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); - } + 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 (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { + pass.set_pipeline(if c.compact_vertices { + &self.mesh_plain_transparent_pipeline + } else { + &self.mesh_transparent_pipeline + }); pass.set_bind_group( 1, c.material_bind_group @@ -4093,25 +3926,16 @@ impl Pipeline { continue; } pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); pass.set_index_buffer( c.transp_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - 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.draw_indexed( + 0..c.transp_index_count, + 0, + 0..c.instance_count, + ); } } // Selection / hover highlight reuses index ranges already @@ -4124,10 +3948,6 @@ impl Pipeline { { 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() @@ -4141,9 +3961,17 @@ impl Pipeline { let Some(chunk) = chunk else { continue; }; - if !chunk.visible || draw.range.index_count == 0 { + if draw.range.index_count == 0 { continue; } + pass.set_pipeline(match (kind, chunk.compact_vertices) { + (MeshHighlightKind::Selected, false) => &self.mesh_selected_pipeline, + (MeshHighlightKind::Selected, true) => { + &self.mesh_plain_selected_pipeline + } + (MeshHighlightKind::Hover, false) => &self.mesh_hover_pipeline, + (MeshHighlightKind::Hover, true) => &self.mesh_plain_hover_pipeline, + }); pass.set_bind_group( 1, chunk @@ -4153,6 +3981,7 @@ impl Pipeline { &[], ); pass.set_vertex_buffer(0, chunk.vertex_buffer.slice(..)); + pass.set_vertex_buffer(1, chunk.instance_buffer.slice(..)); if draw.range.transparent { pass.set_index_buffer( chunk.transp_index_buffer.slice(..), @@ -4178,10 +4007,7 @@ 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 (mesh_command, c) in self.active_mesh_chunks_indexed() { - if !c.visible { - continue; - } + for (_, c) in self.active_mesh_chunks_indexed() { pass.set_bind_group( 1, c.material_bind_group @@ -4189,41 +4015,22 @@ impl Pipeline { .unwrap_or(&self.mesh_default_material_bind_group), &[], ); + pass.set_vertex_buffer(1, c.instance_buffer.slice(..)); if c.wire_index_count != 0 { pass.set_vertex_buffer(0, c.vertex_buffer.slice(..)); pass.set_index_buffer( c.wire_index_buffer.slice(..), wgpu::IndexFormat::Uint32, ); - 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, - ); - } + 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(..)); - 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); - } + pass.draw(0..c.edge_vertex_count, 0..c.instance_count); } } pass.set_pipeline(&self.silhouette_black_pipeline); @@ -4758,47 +4565,6 @@ fn round_up_tex(n: u32) -> u32 { ((n.max(1) + GRID - 1) / GRID) * GRID } -fn aabb3_offscreen( - aabb: [f32; 6], - view_rot: glam::Mat4, - eye: glam::DVec3, - clip_w: u32, - clip_h: u32, -) -> bool { - if aabb.iter().any(|value| !value.is_finite()) { - return false; - } - let w = clip_w as f32; - let h = clip_h as f32; - 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 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); - } - } - } - 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 /// viewport rect (extended by `MARGIN_FRAC` to absorb pan inertia and /// avoid edge pop-in). Phase 2.2 mesh-frustum / Phase 2.3 hatch + diff --git a/src/scene/view/render.rs b/src/scene/view/render.rs index b07b9bc12..20d20f725 100644 --- a/src/scene/view/render.rs +++ b/src/scene/view/render.rs @@ -390,8 +390,7 @@ impl shader::Primitive for Primitive { inner.wire_cull_key = (u64::MAX, u64::MAX, 0, 0); inner.hatch_lod_key = (usize::MAX, u64::MAX, 0, 0, false); inner.wipeout_lod_key = (usize::MAX, u64::MAX, 0, 0, false); - inner.mesh_lod_key = (usize::MAX, u64::MAX, 0, 0); - inner.silhouette_key = (usize::MAX, u64::MAX, u64::MAX, false); + inner.silhouette_key = (usize::MAX, u64::MAX, [u32::MAX; 3], false); inner.render_sig = u64::MAX; } // The MSAA / depth / resolve textures are always sized to the @@ -1039,13 +1038,14 @@ impl shader::Primitive for Primitive { let silhouette_key = ( Arc::as_ptr(&vp.meshes) as usize, vp.wire_content_id, - vp.camera_generation, + vp.view_dir.to_array().map(f32::to_bits), silhouette_enabled, ); if inner.silhouette_key != silhouette_key { inner.upload_silhouettes( device, if silhouette_enabled { &vp.meshes[..] } else { &[] }, + vp.wire_content_id, vp.view_dir, ); inner.silhouette_key = silhouette_key; @@ -1073,22 +1073,6 @@ impl shader::Primitive for Primitive { inner.compute_wipeout_lod(view_rot, eye, clip_size.width, clip_size.height); inner.wipeout_lod_key = wipeout_lod_key; } - let mesh_lod_key = ( - Arc::as_ptr(&vp.meshes) as usize, - vp.camera_generation, - clip_size.width, - clip_size.height, - ); - if inner.mesh_lod_key != mesh_lod_key { - inner.compute_mesh_lod( - queue, - view_rot, - eye, - clip_size.width, - clip_size.height, - ); - inner.mesh_lod_key = mesh_lod_key; - } let cull_key = ( vp.wire_content_id, vp.camera_generation, @@ -1298,6 +1282,7 @@ fn render_signature(vp: &ViewportData, clip_w: u32, clip_h: u32) -> u64 { vp.show_2d_solid_fills.hash(&mut h); vp.mesh_fill.hash(&mut h); vp.show_3d_edges.hash(&mut h); + vp.display_silhouette.hash(&mut h); vp.hidden_line.hash(&mut h); // ViewCube visibility is excluded from the *scene* signature elsewhere only // for the live-hover pass; here it MUST invalidate the cache so toggling the @@ -3701,6 +3686,7 @@ impl Scene { } else { Arc::new(vec![]) }; + let navigating = !inst.paper_sheet && self.navigating_lod(); Some(ViewportData { instance_id, force_rasterize, @@ -3744,7 +3730,7 @@ impl Scene { // actively moving; the scene-render cache holds the full-quality // (hatched) frame once it settles. Only applied to the on-screen // Model / paper content — the paper *sheet* keeps its fills. - skip_hatch: self.hatch_lod_enabled() && !inst.paper_sheet && self.navigating_lod(), + skip_hatch: self.hatch_lod_enabled() && navigating, skip_background: !inst.paper_sheet && self.current_layout != "Model", geometry_epoch: self.geometry_epoch, camera_generation: self.camera_generation, diff --git a/src/shaders/mesh.wgsl b/src/shaders/mesh.wgsl index 849b7d22f..806abbe62 100644 --- a/src/shaders/mesh.wgsl +++ b/src/shaders/mesh.wgsl @@ -98,9 +98,6 @@ struct MeshInstance { normal_row_2: vec4, }; -@group(1) @binding(15) -var mesh_instances: array; - struct VertexIn { @location(0) position: vec3, @location(1) normal: vec3, @@ -111,6 +108,22 @@ struct VertexIn { @location(12) uv_refraction_normal: vec4, }; +struct PlainVertexIn { + @location(0) position: vec3, + @location(1) normal: vec3, + @location(3) position_low: vec3, +}; + +struct InstanceIn { + @location(4) model_row_0: vec4, + @location(5) model_row_1: vec4, + @location(7) model_row_2: vec4, + @location(8) translation_low: vec4, + @location(9) normal_row_0: vec4, + @location(13) normal_row_1: vec4, + @location(14) normal_row_2: vec4, +}; + struct VertexOut { @builtin(position) clip_pos: vec4, @location(0) normal: vec3, @@ -153,13 +166,25 @@ fn relative_position( return (world_high - u.eye_high) + (world_low - u.eye_low); } +fn mesh_instance(input: InstanceIn) -> MeshInstance { + return MeshInstance( + input.model_row_0, + input.model_row_1, + input.model_row_2, + input.translation_low, + input.normal_row_0, + input.normal_row_1, + input.normal_row_2, + ); +} + @vertex fn vs_main( v: VertexIn, - @builtin(instance_index) instance_index: u32, + i: InstanceIn, ) -> VertexOut { var out: VertexOut; - let instance = mesh_instances[instance_index]; + let instance = mesh_instance(i); let rel = relative_position(v.position, v.position_low, instance); out.clip_pos = u.view_rot * vec4(rel, 1.0); out.normal = normalize(vec3( @@ -178,13 +203,38 @@ fn vs_main( return out; } +@vertex +fn vs_main_plain( + v: PlainVertexIn, + i: InstanceIn, +) -> VertexOut { + var out: VertexOut; + let instance = mesh_instance(i); + let rel = relative_position(v.position, v.position_low, instance); + out.clip_pos = u.view_rot * vec4(rel, 1.0); + 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.uv_diffuse = vec2(0.0); + out.uv_specular = vec2(0.0); + out.uv_reflection = vec2(0.0); + out.uv_opacity = vec2(0.0); + out.uv_bump = vec2(0.0); + out.uv_refraction = vec2(0.0); + out.uv_normal = vec2(0.0); + return out; +} + @vertex fn vs_edge( v: EdgeVertexIn, - @builtin(instance_index) instance_index: u32, + i: InstanceIn, ) -> EdgeVertexOut { var out: EdgeVertexOut; - let instance = mesh_instances[instance_index]; + let instance = mesh_instance(i); let rel = relative_position(v.position, v.position_low, instance); out.clip_pos = u.view_rot * vec4(rel, 1.0); out.color = v.color; diff --git a/src/shaders/mesh_cull.wgsl b/src/shaders/mesh_cull.wgsl deleted file mode 100644 index 2a14f01a4..000000000 --- a/src/shaders/mesh_cull.wgsl +++ /dev/null @@ -1,78 +0,0 @@ -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); -} diff --git a/src/shaders/shadow.wgsl b/src/shaders/shadow.wgsl index 957e4ffab..7003221c4 100644 --- a/src/shaders/shadow.wgsl +++ b/src/shaders/shadow.wgsl @@ -51,19 +51,24 @@ struct MeshInstance { }; @group(0) @binding(0) var u: Uniforms; -@group(1) @binding(15) var mesh_instances: array; struct VertexIn { @location(0) position: vec3, @location(3) position_low: vec3, }; +struct InstanceIn { + @location(4) model_row_0: vec4, + @location(5) model_row_1: vec4, + @location(7) model_row_2: vec4, + @location(8) translation_low: vec4, +}; + @vertex fn vs_main( vertex: VertexIn, - @builtin(instance_index) instance_index: u32, + instance: InstanceIn, ) -> @builtin(position) vec4 { - let instance = mesh_instances[instance_index]; let world_high = vec3( dot(instance.model_row_0.xyz, vertex.position) + instance.model_row_0.w, dot(instance.model_row_1.xyz, vertex.position) + instance.model_row_1.w,