From ff7a34701313a46e47a35d0c3545faeccb7b698c Mon Sep 17 00:00:00 2001 From: Irakli Kourtis Date: Tue, 13 Oct 2026 16:27:16 -0400 Subject: [PATCH] Display tessellation plans (3/4): display meshes from plans, GPU evaluation in the viewport (#805) Display meshes are built from tessellation plans, with GPU evaluation in the viewport. --- Cargo.lock | 14 + Cargo.toml | 4 + crates/doc/src/eval.rs | 31 +- crates/doc/src/project.rs | 44 +- crates/geom/src/mesh.rs | 62 ++- crates/geom/src/tests.rs | 34 ++ crates/gpu/Cargo.toml | 23 + crates/gpu/src/lib.rs | 442 ++++++++++++++++++ crates/gpu/src/tess.wgsl | 289 ++++++++++++ crates/gpu/src/tests.rs | 168 +++++++ crates/kernel/src/body.rs | 40 +- crates/kernel/src/planar_sources.rs | 7 +- crates/kernel/src/smooth_edges.rs | 12 +- crates/ui-egui/Cargo.toml | 5 + crates/ui-egui/src/gpu.rs | 284 +++++++++-- crates/ui-egui/src/hole_placement_ui_tests.rs | 4 +- crates/ui-egui/src/viewport.rs | 59 ++- docs/perf.md | 25 + xtask/src/layers.rs | 4 +- 19 files changed, 1474 insertions(+), 77 deletions(-) create mode 100644 crates/gpu/Cargo.toml create mode 100644 crates/gpu/src/lib.rs create mode 100644 crates/gpu/src/tess.wgsl create mode 100644 crates/gpu/src/tests.rs diff --git a/Cargo.lock b/Cargo.lock index b5906e2a..b9a31dcc 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -3345,6 +3345,18 @@ dependencies = [ "thiserror 2.0.21", ] +[[package]] +name = "solidcraft-gpu" +version = "0.8.0" +dependencies = [ + "log", + "pollster 0.4.0", + "solidcraft-geom", + "solidcraft-kernel", + "thiserror 2.0.21", + "wgpu", +] + [[package]] name = "solidcraft-io" version = "0.9.0" @@ -3415,9 +3427,11 @@ dependencies = [ "egui-wgpu", "image", "log", + "pollster 0.4.0", "serde", "serde_json", "solidcraft-engine", + "solidcraft-gpu", "ttf-parser", "web-time", ] diff --git a/Cargo.toml b/Cargo.toml index 10fefeb0..292c303b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -18,6 +18,7 @@ solidcraft-geom = { path = "crates/geom" } solidcraft-sketch = { path = "crates/sketch" } solidcraft-kernel = { path = "crates/kernel" } solidcraft-render = { path = "crates/render" } +solidcraft-gpu = { path = "crates/gpu" } solidcraft-doc = { path = "crates/doc" } solidcraft-io = { path = "crates/io" } solidcraft-engine = { path = "crates/engine" } @@ -58,6 +59,9 @@ eframe = { version = "0.36", default-features = false, features = ["persistence" # Instant/SystemTime that also work in the browser (std panics on wasm32-unknown-unknown). web-time = "1" egui-wgpu = "0.36" +# the wgpu egui-wgpu uses, for compute (crates/gpu); backends come from the apps' features +wgpu = { version = "30", default-features = false, features = ["std", "wgsl"] } +pollster = "0.4" rfd = "0.17" [workspace.lints.rust] diff --git a/crates/doc/src/eval.rs b/crates/doc/src/eval.rs index b4ec9290..f5799db5 100644 --- a/crates/doc/src/eval.rs +++ b/crates/doc/src/eval.rs @@ -5,6 +5,7 @@ use std::hash::{Hash, Hasher}; use std::sync::{Arc, OnceLock}; use serde::Serialize; +use solidcraft_geom::tess::TessPlan; use solidcraft_geom::{Mesh, Plane, Region2, Vec2, Vec3}; use solidcraft_kernel::{self as kernel, Body, BoolOp}; use solidcraft_sketch::{CurveKind, Profile, SolveReport, find_profiles}; @@ -27,6 +28,9 @@ pub struct ModelBody { // Timeline snapshots are cloned before meshes are requested. Share the cell itself so // meshing any snapshot also warms every unchanged copy of this body. mesh: Arc>>, + /// The display plan the mesh is evaluated from (none when the body has no plan: a mesh + /// body, or one the planner refuses), shared like the mesh. A GPU viewport draws it. + plan: Arc>>>, /// Identity of this body's content for the evaluation cache: the key of the evaluation /// that produced it (kept while later features pass it through unchanged). key: u64, @@ -51,6 +55,7 @@ impl ModelBody { body, feature, mesh: Arc::new(OnceLock::new()), + plan: Arc::new(OnceLock::new()), key: 0, names: Default::default(), made: Default::default(), @@ -68,12 +73,34 @@ impl ModelBody { pub fn mesh(&self) -> Arc { self.mesh .get_or_init(|| { - let tol = (self.body.size() * 1e-3).max(1e-3); - Arc::new(self.body.display_mesh_angular(tol, kernel::DISPLAY_ANGLE).unwrap_or_default()) + let mesh = match self.display_plan() { + Some(plan) => self.body.display_mesh_of(&plan), + None => self.body.reference_display_mesh_angular(self.display_tol(), kernel::DISPLAY_ANGLE), + }; + Arc::new(mesh.unwrap_or_default()) }) .clone() } + /// The display plan behind [`ModelBody::mesh`] (same tolerance and angle), made once; `None` + /// when the body has no plan and its mesh comes from the reference mesher. + pub fn display_plan(&self) -> Option> { + self.plan + .get_or_init(|| match self.body.display_plan_angular(self.display_tol(), kernel::DISPLAY_ANGLE) { + Ok(plan) => Some(Arc::new(plan)), + Err(e) => { + log::debug!("{}: no display plan ({e})", self.name); + None + } + }) + .clone() + } + + /// The display mesh's chord tolerance: 1/1000 of the body's size. + fn display_tol(&self) -> f64 { + (self.body.size() * 1e-3).max(1e-3) + } + /// Already computed display geometry; saving must not trigger cold tessellation. pub fn cached_mesh(&self) -> Option> { self.mesh.get().cloned() diff --git a/crates/doc/src/project.rs b/crates/doc/src/project.rs index 28110829..8ee0de43 100644 --- a/crates/doc/src/project.rs +++ b/crates/doc/src/project.rs @@ -970,20 +970,32 @@ pub(crate) fn section(plane: &Plane, m: &Mesh, face: Option) -> Vec> = BTreeMap::new(); - for (t, _) in m.triangles.iter().zip(&m.tri_face) { + // Surface normal at each mesh vertex where the surface is smooth (None at creases), over + // all faces and per face. + type Key = (i64, i64, i64); + let mut normals: BTreeMap> = BTreeMap::new(); + let mut face_normals: BTreeMap<(u32, Key), Option> = BTreeMap::new(); + // Faces on either side of each mesh edge. + let mut edge_faces: BTreeMap<(Key, Key), Vec> = BTreeMap::new(); + let merge = |e: &mut Option, n: Vec3| { + if e.is_some_and(|o| o.dot(n) < 1.0 - 1e-6) { + *e = None; + } + }; + for (t, f) in m.triangles.iter().zip(&m.tri_face) { for k in t { let (Some(p), Some(n)) = (m.positions.get(*k as usize), m.normals.get(*k as usize)) else { continue }; let Some(n) = n.normalized() else { continue }; - normals - .entry(key(*p)) - .and_modify(|e| { - if e.is_some_and(|o| o.dot(n) < 1.0 - 1e-6) { - *e = None; - } - }) - .or_insert(Some(n)); + normals.entry(key(*p)).and_modify(|e| merge(e, n)).or_insert(Some(n)); + face_normals.entry((*f, key(*p))).and_modify(|e| merge(e, n)).or_insert(Some(n)); + } + for (i, j) in [(0usize, 1usize), (1, 2), (2, 0)] { + let (Some(a), Some(b)) = (t.get(i).and_then(|k| m.positions.get(*k as usize)), t.get(j).and_then(|k| m.positions.get(*k as usize))) + else { + continue; + }; + let (ka, kb) = (key(*a), key(*b)); + edge_faces.entry(if ka <= kb { (ka, kb) } else { (kb, ka) }).or_default().push(*f); } } // Where the plane crosses a mesh edge: on a cubic through the end points whose end @@ -997,8 +1009,16 @@ pub(crate) fn section(plane: &Plane, m: &Mesh, face: Option) -> Vec face_normals.get(&(f, key(p))).copied().flatten(), + None => normals.get(&key(p)).copied().flatten(), + }; // Tangent length: the arc length of a circle through both ends with these normals. - let (na, nb) = (normals.get(&key(a)).copied().flatten(), normals.get(&key(b)).copied().flatten()); + let (na, nb) = (normal(a), normal(b)); let l = match (na, nb) { (Some(x), Some(y)) => { let half = 0.5 * x.dot(y).clamp(-1.0, 1.0).acos(); diff --git a/crates/geom/src/mesh.rs b/crates/geom/src/mesh.rs index 0bf84c20..d0b07add 100644 --- a/crates/geom/src/mesh.rs +++ b/crates/geom/src/mesh.rs @@ -396,29 +396,59 @@ impl Mesh { } /// Closest intersection of a ray with the mesh: (distance along `dir`, triangle index). + /// + /// Watertight (Woop, Benthin & Wald, "Watertight Ray/Triangle Intersection", JCGT 2(1), + /// 2013): the vertices are moved into a frame where the ray runs along +z through the + /// origin, and each edge's side test is computed from those per-vertex values alone, so the + /// two triangles that share an edge get exactly opposite values there. A ray through a + /// shared edge or vertex hits at least one of them; it never slips between them. pub fn raycast(&self, o: Vec3, dir: Vec3) -> Option<(f64, usize)> { + let comp = |v: Vec3, k: usize| match k { + 0 => v.x, + 1 => v.y, + _ => v.z, + }; + // The ray's dominant axis becomes z; x and y follow it cyclically, swapped to keep + // the frame right-handed when the ray runs towards −z. + let a = [dir.x.abs(), dir.y.abs(), dir.z.abs()]; + let kz = if a[0] >= a[1] && a[0] >= a[2] { + 0 + } else if a[1] >= a[2] { + 1 + } else { + 2 + }; + let dz = comp(dir, kz); + if !(dz.is_finite() && dz != 0.0) { + return None; + } + let (mut kx, mut ky) = ((kz + 1) % 3, (kz + 2) % 3); + if dz < 0.0 { + std::mem::swap(&mut kx, &mut ky); + } + let (sx, sy, sz) = (comp(dir, kx) / dz, comp(dir, ky) / dz, 1.0 / dz); + // A vertex in the ray's frame: (x, y) across the ray, z scaled along it. + let frame = |p: Vec3| { + let q = p - o; + let z = comp(q, kz); + (comp(q, kx) - sx * z, comp(q, ky) - sy * z, sz * z) + }; let mut best: Option<(f64, usize)> = None; for (i, t) in self.triangles.iter().enumerate() { - let Some([a, b, c]) = self.tri(t) else { continue }; - let (e1, e2) = (b - a, c - a); - let p = dir.cross(e2); - let det = e1.dot(p); - if det.abs() < 1e-14 { + let Some([pa, pb, pc]) = self.tri(t) else { continue }; + let ((ax, ay, az), (bx, by, bz), (cx, cy, cz)) = (frame(pa), frame(pb), frame(pc)); + let u = cx * by - cy * bx; + let v = ax * cy - ay * cx; + let w = bx * ay - by * ax; + if (u < 0.0 || v < 0.0 || w < 0.0) && (u > 0.0 || v > 0.0 || w > 0.0) { continue; } - let inv = 1.0 / det; - let s = o - a; - let u = s.dot(p) * inv; - if !(0.0..=1.0).contains(&u) { + let det = u + v + w; + if det == 0.0 { continue; } - let q = s.cross(e1); - let v = dir.dot(q) * inv; - if v < 0.0 || u + v > 1.0 { - continue; - } - let t = e2.dot(q) * inv; - if t > 1e-9 && best.is_none_or(|(bt, _)| t < bt) { + let t = (u * az + v * bz + w * cz) / det; + if t.is_finite() && t > 1e-9 && best.is_none_or(|(bt, _)| t < bt) { best = Some((t, i)); } } diff --git a/crates/geom/src/tests.rs b/crates/geom/src/tests.rs index c71c8f3c..317a2697 100644 --- a/crates/geom/src/tests.rs +++ b/crates/geom/src/tests.rs @@ -59,6 +59,40 @@ fn raycast_hits_front_face() { assert!(m.raycast(Vec3::new(5.0, 5.0, 10.0), Vec3::new(0.0, 0.0, -1.0)).is_none()); } +/// A ray through a vertex or edge that triangles share never slips between them: a grid of +/// triangles on a square, hit at every inner vertex and edge midpoint from oblique directions +/// (on the square's own border nothing is shared, and a ray grazing it may miss). +#[test] +fn raycast_never_slips_between_triangles() { + let n = 6; + let positions: Vec = + (0..=n).flat_map(|j| (0..=n).map(move |i| Vec3::new(f64::from(i) * 40.0 / 6.0 + 100.0, f64::from(j) * 5.0, 20.0))).collect(); + let id = |i: u32, j: u32| j * (n + 1) + i; + let mut triangles = Vec::new(); + for j in 0..n { + for i in 0..n { + triangles.push([id(i, j), id(i + 1, j), id(i + 1, j + 1)]); + triangles.push([id(i, j), id(i + 1, j + 1), id(i, j + 1)]); + } + } + let tri_face = vec![0; triangles.len()]; + let m = Mesh { positions, triangles, tri_face, ..Default::default() }; + let dirs = [Vec3::new(-0.57735, 0.57735, -0.57735), Vec3::new(0.3, -0.7, -0.648), Vec3::new(-0.123456789, -0.2, -1.0), Vec3::new(0.0, 0.0, -1.0)]; + let mut targets = Vec::new(); + for j in 1..2 * n { + for i in 1..2 * n { + targets.push(Vec3::new(f64::from(i) * 20.0 / 6.0 + 100.0, f64::from(j) * 2.5, 20.0)); + } + } + for d in dirs { + for p in &targets { + let o = *p - d * 50.0; + let hit = m.raycast(o, d); + assert!(hit.is_some_and(|(t, _)| (t - 50.0).abs() < 1e-6), "{p:?} along {d:?}: {hit:?}"); + } + } +} + #[test] fn planes() { assert_eq!(Plane::XY.normal(), Vec3::Z); diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml new file mode 100644 index 00000000..1571d746 --- /dev/null +++ b/crates/gpu/Cargo.toml @@ -0,0 +1,23 @@ +[package] +name = "solidcraft-gpu" +description = "SolidCraft GPU compute: evaluates display tessellation plans (surface points and normals) in a compute shader" +version.workspace = true +edition.workspace = true +license.workspace = true + +[lints] +workspace = true + +[dependencies] +solidcraft-geom = { workspace = true } +wgpu = { workspace = true } +log = { workspace = true } +thiserror = { workspace = true } + +[target.'cfg(not(target_arch = "wasm32"))'.dependencies] +pollster = { workspace = true } +# Native backends, so tests and the CLI can open a device without a window. +wgpu = { workspace = true, features = ["metal", "vulkan", "dx12", "gles"] } + +[dev-dependencies] +solidcraft-kernel = { workspace = true } diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs new file mode 100644 index 00000000..c723f088 --- /dev/null +++ b/crates/gpu/src/lib.rs @@ -0,0 +1,442 @@ +//! GPU evaluation of display tessellation plans: the surface point and normal at every +//! vertex's parameters, in a compute shader (`tess.wgsl`), the GPU twin of +//! [`TessPlan::evaluate`]. Shared edge nodes keep their exact positions, so faces still meet +//! exactly; evaluation runs in f32. +//! +//! Use [`TessGpu::new`] with the application's device, or [`TessGpu::headless`] (native only) +//! to open one without a window, as the CLI and tests do. A viewport draws a plan straight from +//! the GPU: [`TessGpu::prepare`] uploads it, [`TessGpu::record`] adds the evaluation to the +//! frame's command encoder, and the render pass reads [`GpuMesh::vertices`] as its vertex buffer +//! with [`GpuMesh::indices`]; nothing comes back to the CPU. +#![deny(clippy::unwrap_used, clippy::expect_used, clippy::panic, clippy::unimplemented, clippy::todo, clippy::unreachable)] +#![forbid(unsafe_code)] + +#[cfg(not(target_arch = "wasm32"))] +use solidcraft_geom::Mesh; +use solidcraft_geom::Vec3; +use solidcraft_geom::tess::{EvalCurve, EvalSurface, TessPlan}; +use wgpu::util::DeviceExt; + +/// Most vertices evaluated in one dispatch (bounds the buffers; larger plans go in batches). +pub const MAX_BATCH: usize = 1 << 20; +const WORKGROUP: u32 = 64; + +#[derive(Debug, thiserror::Error)] +pub enum GpuError { + #[error("no GPU adapter is available")] + NoAdapter, + #[error("the GPU device could not be opened: {0}")] + Device(String), + #[error("the GPU evaluation failed: {0}")] + Failed(String), + #[error("not supported yet: {0}")] + Unsupported(String), +} + +pub type Result = std::result::Result; + +/// Bytes per evaluated vertex in [`GpuMesh::vertices`]: the position (x, y, z, 1) then the +/// normal (x, y, z, 0), each four f32. +pub const VERTEX_STRIDE: u64 = 32; +/// Byte offset of the normal within a vertex. +pub const NORMAL_OFFSET: u64 = 16; + +/// A plan on the GPU: its inputs, and the buffers the evaluation writes and a render pass +/// reads. Evaluated once [`TessGpu::record`] has run in a submitted encoder. +pub struct GpuMesh { + /// Per vertex: position and normal (see [`VERTEX_STRIDE`]); a vertex and storage buffer. + pub vertices: wgpu::Buffer, + /// The plan's triangles, three u32 vertex indices each. + pub indices: wgpu::Buffer, + pub index_count: u32, + pub vertex_count: u32, + /// Each vertex's face (the plan's face numbering), for colouring by face. + pub vertex_face: Vec, + bind: wgpu::BindGroup, + // The inputs, kept until the evaluation has run. + _inputs: [wgpu::Buffer; 4], +} + +/// The compute pipeline, ready to evaluate plans on one device. +pub struct TessGpu { + adapter: String, + device: wgpu::Device, + /// For the blocking [`TessGpu::evaluate`] (native only). + #[cfg(not(target_arch = "wasm32"))] + queue: wgpu::Queue, + pipeline: wgpu::ComputePipeline, + layout: wgpu::BindGroupLayout, +} + +impl TessGpu { + /// The pipeline on an existing device (`adapter`: its name, for reports). + pub fn new(adapter: impl Into, device: wgpu::Device, queue: wgpu::Queue) -> Self { + let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("solidcraft tessellation"), + source: wgpu::ShaderSource::Wgsl(include_str!("tess.wgsl").into()), + }); + let entry = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { ty, has_dynamic_offset: false, min_binding_size: None }, + count: None, + }; + let storage = |read_only| wgpu::BufferBindingType::Storage { read_only }; + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("solidcraft tessellation"), + entries: &[ + entry(0, wgpu::BufferBindingType::Uniform), + entry(1, storage(true)), + entry(2, storage(true)), + entry(3, storage(true)), + entry(4, storage(false)), + ], + }); + let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("solidcraft tessellation"), + bind_group_layouts: &[Some(&layout)], + immediate_size: 0, + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("solidcraft tessellation"), + layout: Some(&pipeline_layout), + module: &module, + entry_point: Some("main"), + compilation_options: Default::default(), + cache: None, + }); + // A viewport records the evaluation into its own frame (see `TessGpu::record`). + #[cfg(target_arch = "wasm32")] + let _ = queue; + TessGpu { + adapter: adapter.into(), + device, + #[cfg(not(target_arch = "wasm32"))] + queue, + pipeline, + layout, + } + } + + /// A device of its own, without a window (native only). + #[cfg(not(target_arch = "wasm32"))] + pub fn headless() -> Result { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); + let adapter = pollster::block_on( + instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }), + ) + .map_err(|_| GpuError::NoAdapter)?; + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + label: Some("solidcraft compute"), + required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), + ..Default::default() + })) + .map_err(|e| GpuError::Device(e.to_string()))?; + Ok(Self::new(adapter.get_info().name, device, queue)) + } + + /// The adapter's name, for reports. + pub fn adapter_name(&self) -> &str { + &self.adapter + } + + /// Evaluate `plan` on the GPU into the same mesh [`TessPlan::evaluate`] gives (positions + /// and normals in f32, fixed edge nodes exact). Blocks until the GPU is done (native). + #[cfg(not(target_arch = "wasm32"))] + pub fn evaluate(&self, plan: &TessPlan) -> Result { + let packed = pack(plan)?; + let mut mesh = Mesh::default(); + let mut base = 0usize; + for chunk in packed.vertices.chunks(MAX_BATCH * VERTEX_WORDS) { + let n = chunk.len() / VERTEX_WORDS; + let out = self.run(&packed, chunk, n)?; + for k in 0..n { + let at = |i: usize| out.get(8 * k + i).copied().map_or(0.0, f64::from); + mesh.positions.push(Vec3::new(at(0), at(1), at(2))); + mesh.normals.push(Vec3::new(at(4), at(5), at(6))); + } + base += n; + } + if base != packed.count { + return Err(GpuError::Failed("vertex count mismatch".into())); + } + // Fixed nodes exactly (the shader has them in f32), and the plan's topology. + for (k, p) in packed.fixed { + if let Some(slot) = mesh.positions.get_mut(k) { + *slot = p; + } + } + let mut offset = 0u32; + for f in &plan.faces { + let n = u32::try_from(f.uv.len()).map_err(|_| GpuError::Failed("too many vertices".into()))?; + for t in &f.triangles { + if t.iter().all(|&i| i < n) { + mesh.triangles.push([offset + t[0], offset + t[1], offset + t[2]]); + mesh.tri_face.push(f.face); + } + } + offset += n; + } + mesh.edges = plan.edges.clone(); + mesh.edge_faces = plan.edge_faces.clone(); + Ok(mesh) + } + + /// Upload `plan` and create the buffers its evaluation writes; nothing runs until + /// [`TessGpu::record`]. Refused (so the caller draws its CPU mesh) for a plan with more than + /// [`MAX_BATCH`] vertices or a face without an evaluable surface. + pub fn prepare(&self, plan: &TessPlan) -> Result { + let packed = pack(plan)?; + if packed.count > MAX_BATCH { + return Err(GpuError::Unsupported(format!("a plan of {} vertices on the GPU", packed.count))); + } + let d = &self.device; + let count = word(packed.count)?; + let init = |label: &str, contents: &[u8], usage| { + // wgpu rejects empty bindings: pad to one word. + let contents = if contents.is_empty() { &[0u8; 16][..] } else { contents }; + d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some(label), contents, usage }) + }; + let words = |v: &[u32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::>(); + let floats = |v: &[f32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::>(); + let params = init("tess params", &words(&[count, 0, 0, 0]), wgpu::BufferUsages::UNIFORM); + let surfaces = init("tess surfaces", &words(&packed.surfaces), wgpu::BufferUsages::STORAGE); + let data = init("tess data", &floats(&packed.data), wgpu::BufferUsages::STORAGE); + let verts = init("tess vertices", &words(&packed.vertices), wgpu::BufferUsages::STORAGE); + let vertices = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("tess mesh vertices"), + size: (packed.count.max(1) as u64) * VERTEX_STRIDE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let mut tri: Vec = Vec::new(); + let mut vertex_face = Vec::with_capacity(packed.count); + let mut offset = 0u32; + for f in &plan.faces { + let n = word(f.uv.len())?; + for t in f.triangles.iter().filter(|t| t.iter().all(|&i| i < n)) { + tri.extend([offset + t[0], offset + t[1], offset + t[2]]); + } + vertex_face.extend(std::iter::repeat_n(f.face, f.uv.len())); + offset += n; + } + let indices = init("tess mesh indices", &words(&tri), wgpu::BufferUsages::INDEX); + let bind = d.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("tess mesh"), + layout: &self.layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: params.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: surfaces.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: data.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: verts.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 4, resource: vertices.as_entire_binding() }, + ], + }); + Ok(GpuMesh { + vertices, + indices, + index_count: word(tri.len())?, + vertex_count: count, + vertex_face, + bind, + _inputs: [params, surfaces, data, verts], + }) + } + + /// Add `mesh`'s evaluation to `encoder` (one dispatch, no wait, nothing read back): once the + /// encoder is submitted, render passes after it read the evaluated vertices. + pub fn record(&self, encoder: &mut wgpu::CommandEncoder, mesh: &GpuMesh) { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("tess mesh"), timestamp_writes: None }); + pass.set_pipeline(&self.pipeline); + pass.set_bind_group(0, &mesh.bind, &[]); + pass.dispatch_workgroups(mesh.vertex_count.div_ceil(WORKGROUP).max(1), 1, 1); + } + + #[cfg(not(target_arch = "wasm32"))] + fn run(&self, packed: &Packed, vertices: &[u32], n: usize) -> Result> { + let d = &self.device; + let bytes_u32 = |v: &[u32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::>(); + let bytes_f32 = |v: &[f32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::>(); + let count = u32::try_from(n).map_err(|_| GpuError::Failed("batch too large".into()))?; + let init = |label: &str, contents: &[u8], usage| { + // wgpu rejects empty bindings: pad to one word. + let contents = if contents.is_empty() { &[0u8; 16][..] } else { contents }; + d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some(label), contents, usage }) + }; + let params = init("tess params", &bytes_u32(&[count, 0, 0, 0]), wgpu::BufferUsages::UNIFORM); + let surfaces = init("tess surfaces", &bytes_u32(&packed.surfaces), wgpu::BufferUsages::STORAGE); + let data = init("tess data", &bytes_f32(&packed.data), wgpu::BufferUsages::STORAGE); + let verts = init("tess vertices", &bytes_u32(vertices), wgpu::BufferUsages::STORAGE); + let size = (n.max(1) * 32) as u64; + let out = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("tess out"), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let read = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("tess read"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let bind = d.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("tess"), + layout: &self.layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: params.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: surfaces.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: data.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: verts.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 4, resource: out.as_entire_binding() }, + ], + }); + let mut enc = d.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("tess") }); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("tess"), timestamp_writes: None }); + pass.set_pipeline(&self.pipeline); + pass.set_bind_group(0, &bind, &[]); + pass.dispatch_workgroups(count.div_ceil(WORKGROUP).max(1), 1, 1); + } + enc.copy_buffer_to_buffer(&out, 0, &read, 0, size); + self.queue.submit([enc.finish()]); + let slice = read.slice(..); + let (tx, rx) = std::sync::mpsc::channel(); + slice.map_async(wgpu::MapMode::Read, move |r| { + let _ = tx.send(r); + }); + d.poll(wgpu::PollType::wait_indefinitely()).map_err(|e| GpuError::Failed(e.to_string()))?; + rx.recv().map_err(|e| GpuError::Failed(e.to_string()))?.map_err(|e| GpuError::Failed(e.to_string()))?; + let bytes = slice.get_mapped_range().map_err(|e| GpuError::Failed(e.to_string()))?; + let values = bytes.as_chunks::<4>().0.iter().map(|b| f32::from_le_bytes(*b)).collect(); + drop(bytes); + read.unmap(); + Ok(values) + } +} + +/// Words per vertex: u, v (f32 bits), surface, flags, then a vec4 position. +const VERTEX_WORDS: usize = 8; +/// Words per surface header (matches `Surface` in the shader). +const SURFACE_WORDS: usize = 12; + +/// A plan as the shader reads it. +struct Packed { + surfaces: Vec, + data: Vec, + vertices: Vec, + count: usize, + /// Exact positions of fixed vertices (index in the evaluated order), for the blocking + /// [`TessGpu::evaluate`] to restore after the f32 evaluation (native only). + #[cfg(not(target_arch = "wasm32"))] + fixed: Vec<(usize, Vec3)>, +} + +fn word(x: usize) -> Result { + u32::try_from(x).map_err(|_| GpuError::Failed("plan too large".into())) +} + +fn push3(data: &mut Vec, v: Vec3) { + data.extend([v.x as f32, v.y as f32, v.z as f32]); +} + +fn push_curve(data: &mut Vec, c: &EvalCurve) -> Result<(u32, u32)> { + let knots = word(data.len())?; + data.extend(c.knots.iter().map(|k| *k as f32)); + let ctrl = word(data.len())?; + data.extend(c.ctrl.iter().flat_map(|h| h.map(|x| x as f32))); + Ok((knots, ctrl)) +} + +fn pack(plan: &TessPlan) -> Result { + let mut surfaces = Vec::with_capacity(plan.surfaces.len() * SURFACE_WORDS); + let mut data: Vec = Vec::new(); + for s in &plan.surfaces { + let mut head = [0u32; SURFACE_WORDS]; + match s { + EvalSurface::Plane { origin, du, dv } => { + head[0] = 0; + head[8] = word(data.len())?; + push3(&mut data, *origin); + push3(&mut data, *du); + push3(&mut data, *dv); + } + EvalSurface::Nurbs { degree, knots, count, ctrl } => { + head[0] = 1; + head[1] = word(degree[0])?; + head[2] = word(degree[1])?; + head[3] = word(count[0])?; + head[4] = word(count[1])?; + head[5] = word(data.len())?; + data.extend(knots[0].iter().map(|k| *k as f32)); + head[6] = word(data.len())?; + data.extend(knots[1].iter().map(|k| *k as f32)); + head[7] = word(data.len())?; + data.extend(ctrl.iter().flat_map(|h| h.map(|x| x as f32))); + } + EvalSurface::Revolved { profile, origin, axis, transform, swap } => { + head[0] = 2; + head[1] = word(profile.degree)?; + head[3] = word(profile.ctrl.len())?; + let (k, c) = push_curve(&mut data, profile)?; + head[5] = k; + head[7] = c; + head[8] = word(data.len())?; + push3(&mut data, *origin); + push3(&mut data, *axis); + data.extend(transform.iter().flatten().map(|x| *x as f32)); + head[9] = u32::from(*swap); + let m = transform; + let det = m[0][0] * (m[1][1] * m[2][2] - m[1][2] * m[2][1]) - m[0][1] * (m[1][0] * m[2][2] - m[1][2] * m[2][0]) + + m[0][2] * (m[1][0] * m[2][1] - m[1][1] * m[2][0]); + head[10] = u32::from(det < 0.0); + } + } + surfaces.extend(head); + } + let mut vertices = Vec::with_capacity(plan.vertex_count() * VERTEX_WORDS); + #[cfg(not(target_arch = "wasm32"))] + let mut fixed = Vec::new(); + let mut count = 0usize; + for f in &plan.faces { + let Some(surface) = f.surface else { + return Err(GpuError::Unsupported("a face without an evaluable surface".into())); + }; + if surface as usize >= plan.surfaces.len() { + return Err(GpuError::Failed("surface index out of range".into())); + } + for (k, uv) in f.uv.iter().enumerate() { + let pinned = f.fixed.get(k).copied(); + let flags = u32::from(f.flip) | if pinned.is_some() { 2 } else { 0 }; + let p = pinned.unwrap_or(Vec3::ZERO); + vertices.extend([ + (uv[0] as f32).to_bits(), + (uv[1] as f32).to_bits(), + surface, + flags, + (p.x as f32).to_bits(), + (p.y as f32).to_bits(), + (p.z as f32).to_bits(), + 0, + ]); + if let Some(p) = pinned { + #[cfg(not(target_arch = "wasm32"))] + fixed.push((count, p)); + #[cfg(target_arch = "wasm32")] + let _ = p; + } + count += 1; + } + } + Ok(Packed { + surfaces, + data, + vertices, + count, + #[cfg(not(target_arch = "wasm32"))] + fixed, + }) +} + +#[cfg(test)] +mod tests; diff --git a/crates/gpu/src/tess.wgsl b/crates/gpu/src/tess.wgsl new file mode 100644 index 00000000..66266300 --- /dev/null +++ b/crates/gpu/src/tess.wgsl @@ -0,0 +1,289 @@ +// Display tessellation: surface points and normals at each vertex's (u, v). +// +// The GPU twin of solidcraft_geom::tess::EvalSurface::{derivs, eval}: the same surface forms, +// the same algorithms (The NURBS Book A2.1 span search, A2.3 basis functions and first +// derivatives, the rational quotient rule), in f32. + +const PLANE: u32 = 0u; +const NURBS: u32 = 1u; +const REVOLVED: u32 = 2u; +// Basis slots for the highest degree (MAX_SPLINE_DEGREE = 7). +const W: u32 = 8u; + +struct Surface { + kind: u32, + // Degrees and control point counts (u, v); a revolved surface's profile uses pu, nu. + pu: u32, + pv: u32, + nu: u32, + nv: u32, + // Offsets into `data` (f32 units): knots along u, knots along v, homogeneous control + // points (4 each, u fastest), and the placement (plane: origin, du, dv; revolved: origin, + // axis, row-major 3×4 transform). + knots_u: u32, + knots_v: u32, + ctrl: u32, + extra: u32, + // Revolved: parameters exchanged; mirror transform (normal reversed). + swap: u32, + mirror: u32, + _pad: u32, +}; + +struct Vertex { + u: f32, + v: f32, + surface: u32, + // bit 0: normal reversed (face against its surface); bit 1: position fixed (p). + flags: u32, + p: vec4, +}; + +struct Params { + count: u32, + _p0: u32, + _p1: u32, + _p2: u32, +}; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var surfaces: array; +@group(0) @binding(2) var data: array; +@group(0) @binding(3) var vertices: array; +// Out: position (xyz, w = 1 when evaluated) and normal (xyz) per vertex. +@group(0) @binding(4) var out: array>; + +struct Derivs { + p: vec3, + du: vec3, + dv: vec3, + ok: bool, +}; + +struct Basis { + n0: array, + n1: array, +}; + +fn vec3_at(o: u32) -> vec3 { + return vec3(data[o], data[o + 1u], data[o + 2u]); +} + +fn ctrl_at(o: u32) -> vec4 { + return vec4(data[o], data[o + 1u], data[o + 2u], data[o + 3u]); +} + +// The last non-empty span starting at or before t, clamped to p ≤ k < n (A2.1, by bisection +// over the knots, which do not decrease). +fn span_of(k0: u32, p: u32, n: u32, t: f32) -> u32 { + if (t >= data[k0 + n]) { + // At or past the end: the last non-empty span. + var k = n - 1u; + while (k > p && data[k0 + k] >= data[k0 + k + 1u]) { + k = k - 1u; + } + return k; + } + if (t < data[k0 + p]) { + return p; + } + var lo = p; + var hi = n; + // Invariant: knot[lo] ≤ t < knot[hi]. + while (hi - lo > 1u) { + let mid = (lo + hi) / 2u; + if (data[k0 + mid] <= t) { + lo = mid; + } else { + hi = mid; + } + } + return lo; +} + +// Non-zero basis functions of degree p on span `span` and their first derivatives (A2.3, n = 1). +fn basis(k0: u32, span: u32, p: u32, t: f32) -> Basis { + var ndu: array, 8>; + var left: array; + var right: array; + ndu[0][0] = 1.0; + for (var j = 1u; j <= p; j = j + 1u) { + left[j] = t - data[k0 + span + 1u - j]; + right[j] = data[k0 + span + j] - t; + var saved = 0.0; + for (var r = 0u; r < j; r = r + 1u) { + ndu[j][r] = right[r + 1u] + left[j - r]; + var temp = 0.0; + if (ndu[j][r] != 0.0) { + temp = ndu[r][j - 1u] / ndu[j][r]; + } + ndu[r][j] = saved + right[r + 1u] * temp; + saved = left[j - r] * temp; + } + ndu[j][j] = saved; + } + var b: Basis; + for (var j = 0u; j <= p; j = j + 1u) { + b.n0[j] = ndu[j][p]; + } + for (var r = 0u; r <= p; r = r + 1u) { + var d = 0.0; + if (r >= 1u && ndu[p][r - 1u] != 0.0) { + d = d + ndu[r - 1u][p - 1u] / ndu[p][r - 1u]; + } + if (r < p && ndu[p][r] != 0.0) { + d = d - ndu[r][p - 1u] / ndu[p][r]; + } + b.n1[r] = d * f32(p); + } + return b; +} + +// Point and derivative of a rational function from its homogeneous value and derivative. +fn rational_point(a: vec4) -> vec3 { + return a.xyz / a.w; +} + +fn rational_der(a: vec4, da: vec4) -> vec3 { + return (da.xyz - rational_point(a) * da.w) / a.w; +} + +fn nurbs(s: Surface, u: f32, v: f32) -> Derivs { + var r: Derivs; + let su = span_of(s.knots_u, s.pu, s.nu, u); + let sv = span_of(s.knots_v, s.pv, s.nv, v); + let bu = basis(s.knots_u, su, s.pu, u); + let bv = basis(s.knots_v, sv, s.pv, v); + var a = vec4(0.0); + var au = vec4(0.0); + var av = vec4(0.0); + for (var j = 0u; j <= s.pv; j = j + 1u) { + for (var i = 0u; i <= s.pu; i = i + 1u) { + let c = ctrl_at(s.ctrl + 4u * ((sv + j - s.pv) * s.nu + su + i - s.pu)); + a = a + bu.n0[i] * bv.n0[j] * c; + au = au + bu.n1[i] * bv.n0[j] * c; + av = av + bu.n0[i] * bv.n1[j] * c; + } + } + r.ok = abs(a.w) > 0.0; + if (r.ok) { + r.p = rational_point(a); + r.du = rational_der(a, au); + r.dv = rational_der(a, av); + } + return r; +} + +// Rodrigues rotation of x by angle (sin, cos) about the unit axis k. +fn rotate(x: vec3, k: vec3, sn: f32, cs: f32) -> vec3 { + return x * cs + cross(k, x) * sn + k * (dot(k, x) * (1.0 - cs)); +} + +fn revolved(s: Surface, u: f32, v: f32) -> Derivs { + var r: Derivs; + var t = u; + var angle = v; + if (s.swap != 0u) { + t = v; + angle = u; + } + // The profile curve: a one-row B-spline. + let span = span_of(s.knots_u, s.pu, s.nu, t); + let b = basis(s.knots_u, span, s.pu, t); + var c = vec4(0.0); + var dc = vec4(0.0); + for (var i = 0u; i <= s.pu; i = i + 1u) { + let h = ctrl_at(s.ctrl + 4u * (span + i - s.pu)); + c = c + b.n0[i] * h; + dc = dc + b.n1[i] * h; + } + r.ok = abs(c.w) > 0.0; + if (!r.ok) { + return r; + } + let o = vec3_at(s.extra); + let k = vec3_at(s.extra + 3u); + let m = s.extra + 6u; + let row0 = vec4(data[m], data[m + 1u], data[m + 2u], data[m + 3u]); + let row1 = vec4(data[m + 4u], data[m + 5u], data[m + 6u], data[m + 7u]); + let row2 = vec4(data[m + 8u], data[m + 9u], data[m + 10u], data[m + 11u]); + let sn = sin(angle); + let cs = cos(angle); + let radius = rotate(rational_point(c) - o, k, sn, cs); + let p = o + radius; + let dt = rotate(rational_der(c, dc), k, sn, cs); + let da = cross(k, radius); + let pp = vec3(dot(row0.xyz, p) + row0.w, dot(row1.xyz, p) + row1.w, dot(row2.xyz, p) + row2.w); + let ldt = vec3(dot(row0.xyz, dt), dot(row1.xyz, dt), dot(row2.xyz, dt)); + let lda = vec3(dot(row0.xyz, da), dot(row1.xyz, da), dot(row2.xyz, da)); + r.p = pp; + if (s.swap != 0u) { + r.du = lda; + r.dv = ldt; + } else { + r.du = ldt; + r.dv = lda; + } + return r; +} + +fn derivs(s: Surface, u: f32, v: f32) -> Derivs { + if (s.kind == PLANE) { + var r: Derivs; + r.du = vec3_at(s.extra + 3u); + r.dv = vec3_at(s.extra + 6u); + r.p = vec3_at(s.extra) + r.du * u + r.dv * v; + r.ok = true; + return r; + } + if (s.kind == NURBS) { + return nurbs(s, u, v); + } + return revolved(s, u, v); +} + +fn oriented(s: Surface, n: vec3) -> vec3 { + if (s.mirror != 0u) { + return -n; + } + return n; +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + if (i >= params.count) { + return; + } + let vert = vertices[i]; + let s = surfaces[vert.surface]; + let d = derivs(s, vert.u, vert.v); + var n = oriented(s, cross(d.du, d.dv)); + let scale = length(d.du) * length(d.dv); + if (!(length(n) > 1e-6 * scale) || !d.ok) { + // A degenerate point (a pole): the normal from a step inside. + let h = 1e-4 * (1.0 + max(abs(vert.u), abs(vert.v))); + let steps = array, 4>(vec2(h, h), vec2(-h, h), vec2(h, -h), vec2(-h, -h)); + for (var k = 0u; k < 4u; k = k + 1u) { + let e = derivs(s, vert.u + steps[k].x, vert.v + steps[k].y); + let m = oriented(s, cross(e.du, e.dv)); + if (e.ok && length(m) > 0.0) { + n = m; + break; + } + } + } + var nn = vec3(0.0, 0.0, 1.0); + if (length(n) > 0.0) { + nn = normalize(n); + } + if ((vert.flags & 1u) != 0u) { + nn = -nn; + } + var p = d.p; + if ((vert.flags & 2u) != 0u) { + p = vert.p.xyz; + } + out[2u * i] = vec4(p, 1.0); + out[2u * i + 1u] = vec4(nn, 0.0); +} diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs new file mode 100644 index 00000000..993c03d0 --- /dev/null +++ b/crates/gpu/src/tests.rs @@ -0,0 +1,168 @@ +use solidcraft_geom::Vec3; +use solidcraft_kernel::{Body, BoolOp}; + +use super::*; + +/// The GPU, or `None` with a note when the machine has no adapter (CI runners without a GPU). +fn gpu() -> Option { + match TessGpu::headless() { + Ok(g) => { + eprintln!("GPU: {}", g.adapter_name()); + Some(g) + } + Err(e) => { + eprintln!("skipping the GPU comparison: {e}"); + None + } + } +} + +fn ok(r: std::result::Result) -> T { + match r { + Ok(v) => v, + Err(e) => panic!("{e:?}"), + } +} + +fn bodies() -> Vec<(&'static str, Body)> { + let bx = ok(solidcraft_kernel::box_solid(Vec3::ZERO, Vec3::new(10.0, 10.0, 4.0))); + let pin = ok(solidcraft_kernel::cylinder(Vec3::new(5.0, 5.0, -1.0), Vec3::Z, 2.0, 6.0)); + let cross = ok(solidcraft_kernel::cylinder(Vec3::new(5.0, -1.0, 2.0), Vec3::Y, 1.0, 12.0)); + let cut = |a: &Body, b: &Body| ok(solidcraft_kernel::boolean(a, b, BoolOp::Cut)).unwrap_or_else(|| a.clone()); + let edges = [Vec3::new(5.0, 0.0, 4.0), Vec3::new(10.0, 5.0, 4.0), Vec3::new(5.0, 10.0, 4.0), Vec3::new(0.0, 5.0, 4.0)]; + vec![ + ("box with a hole", cut(&bx, &pin)), + ("cross hole", cut(&pin, &cross)), + ("torus", ok(solidcraft_kernel::torus(Vec3::ZERO, 6.0, 1.5))), + ("fillets", ok(solidcraft_kernel::fillet(&bx, &edges, 1.5))), + ( + "mirrored", + ok(solidcraft_kernel::transform_matrix(&pin, [[-1.0, 0.0, 0.0, 0.0], [0.0, 1.0, 0.0, 0.0], [0.0, 0.0, 1.0, 0.0], [3.0, 0.0, 1.0, 1.0]])), + ), + ] +} + +#[test] +fn gpu_evaluation_matches_the_cpu() { + let Some(gpu) = gpu() else { return }; + for (name, body) in bodies() { + let plan = ok(body.display_plan(0.01)); + let cpu = plan.evaluate(); + let got = ok(gpu.evaluate(&plan)); + assert_eq!(got.triangles, cpu.triangles, "{name}"); + assert_eq!(got.tri_face, cpu.tri_face, "{name}"); + assert_eq!(got.positions.len(), cpu.positions.len(), "{name}"); + let size = 20.0; + for (k, (p, q)) in got.positions.iter().zip(&cpu.positions).enumerate() { + assert!(p.dist(*q) < size * 1e-5, "{name}: vertex {k} at {p:?}, CPU {q:?}"); + } + for (k, (n, m)) in got.normals.iter().zip(&cpu.normals).enumerate() { + assert!(n.dist(*m) < 2e-3, "{name}: normal {k} {n:?}, CPU {m:?}"); + } + // Fixed edge nodes are exact: shared boundaries are bit-identical to the CPU mesh's. + let mut k = 0; + for f in &plan.faces { + for (j, p) in f.fixed.iter().enumerate() { + assert_eq!(got.positions.get(k + j), Some(p), "{name}"); + } + k += f.uv.len(); + } + } +} + +/// A plan drawn straight from the GPU: `prepare` and `record` into an encoder, as the viewport +/// does, give the vertices and triangles the CPU evaluator gives (read back here only to check). +#[test] +fn a_resident_mesh_matches_the_cpu() { + let Some(gpu) = gpu() else { return }; + for (name, body) in bodies() { + let plan = ok(body.display_plan_angular(0.01, std::f64::consts::PI / 9.0)); + let cpu = plan.evaluate(); + let mesh = ok(gpu.prepare(&plan)); + assert_eq!(mesh.vertex_count as usize, cpu.positions.len(), "{name}"); + assert_eq!(mesh.index_count as usize, cpu.triangles.len() * 3, "{name}"); + assert_eq!(mesh.vertex_face.len(), cpu.positions.len(), "{name}"); + for (t, f) in cpu.triangles.iter().zip(&cpu.tri_face) { + assert!(t.iter().all(|&i| mesh.vertex_face[i as usize] == *f), "{name}: a triangle's vertices belong to its face"); + } + let d = &gpu.device; + let read = |src: &wgpu::Buffer, size: u64| { + let dst = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("test read"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let mut enc = d.create_command_encoder(&Default::default()); + gpu.record(&mut enc, &mesh); + enc.copy_buffer_to_buffer(src, 0, &dst, 0, size); + gpu.queue.submit([enc.finish()]); + let slice = dst.slice(..); + slice.map_async(wgpu::MapMode::Read, |_| {}); + ok(d.poll(wgpu::PollType::wait_indefinitely())); + let bytes = ok(slice.get_mapped_range()); + bytes.as_chunks::<4>().0.iter().map(|b| f32::from_le_bytes(*b)).collect::>() + }; + let out = read(&mesh.vertices, u64::from(mesh.vertex_count) * VERTEX_STRIDE); + for (k, (p, n)) in cpu.positions.iter().zip(&cpu.normals).enumerate() { + let at = |i: usize| f64::from(out[8 * k + i]); + assert!(Vec3::new(at(0), at(1), at(2)).dist(*p) < 20.0 * 1e-5, "{name}: vertex {k}"); + assert!(Vec3::new(at(4), at(5), at(6)).dist(*n) < 2e-3, "{name}: normal {k}"); + } + } +} + +/// Display meshing of a 400-hole plate (#303's benchmark): reference mesher, plan, CPU and GPU +/// evaluation. `cargo test --release -p solidcraft-gpu bench_plate -- --ignored --nocapture`. +#[test] +#[ignore] +fn bench_plate() { + use solidcraft_geom::{Loop2, Plane, Region2, Seg2, Vec2}; + let rect = |w: f64| Loop2 { + segs: vec![ + Seg2::Line { a: Vec2::new(0.0, 0.0), b: Vec2::new(w, 0.0) }, + Seg2::Line { a: Vec2::new(w, 0.0), b: Vec2::new(w, w) }, + Seg2::Line { a: Vec2::new(w, w), b: Vec2::new(0.0, w) }, + Seg2::Line { a: Vec2::new(0.0, w), b: Vec2::new(0.0, 0.0) }, + ], + }; + let holes: Vec = (0..400) + .map(|k| { + let c = Vec2::new(5.0 + 10.0 * (k % 20) as f64, 5.0 + 10.0 * (k / 20) as f64); + Loop2 { segs: vec![Seg2::Arc { center: c, radius: 3.0, start: 0.0, sweep: -std::f64::consts::TAU }] } + }) + .collect(); + let plate = ok(solidcraft_kernel::extrude(&Plane::XY, &[Region2 { outer: rect(200.0), holes }], 0.0, 5.0)).remove(0); + let gpu = gpu(); + let ms = |t: std::time::Instant| t.elapsed().as_secs_f64() * 1e3; + let median = |mut v: Vec| { + v.sort_by(f64::total_cmp); + v[v.len() / 2] + }; + println!("faces {}", plate.face_count()); + println!("| tol | reference | plan | CPU eval | GPU eval | triangles (ref / plan) |"); + for tol in [0.3, 0.03, 0.006] { + let (mut r, mut p, mut c, mut g) = (vec![], vec![], vec![], vec![]); + let (mut nr, mut np) = (0, 0); + for _ in 0..5 { + let t = std::time::Instant::now(); + let m = ok(plate.tessellate(tol)); + r.push(ms(t)); + nr = m.triangles.len(); + let t = std::time::Instant::now(); + let plan = ok(plate.display_plan(tol)); + p.push(ms(t)); + let t = std::time::Instant::now(); + let m = plan.evaluate(); + c.push(ms(t)); + np = m.triangles.len(); + if let Some(gpu) = &gpu { + let t = std::time::Instant::now(); + let _ = ok(gpu.evaluate(&plan)); + g.push(ms(t)); + } + } + let g = if g.is_empty() { "-".to_string() } else { format!("{:.1} ms", median(g)) }; + println!("| {tol} | {:.1} ms | {:.1} ms | {:.1} ms | {g} | {nr} / {np} |", median(r), median(p), median(c)); + } +} diff --git a/crates/kernel/src/body.rs b/crates/kernel/src/body.rs index 0929eed8..211cfcd8 100644 --- a/crates/kernel/src/body.rs +++ b/crates/kernel/src/body.rs @@ -1093,9 +1093,21 @@ impl Body { has.iter().filter(|h| !**h).count() } - /// Display mesh: like [`Body::tessellate`], with seam edges flagged. + /// Display mesh: like [`Body::tessellate`], with seam edges flagged. Meshed from the + /// body's [`Body::display_plan`] (shared edge nodes, structured grids), or the way + /// [`Body::tessellate`] meshes when no plan can be made. pub fn display_mesh(&self, tol: f64) -> Result { - self.flagged(self.tessellate(tol)?) + if self.is_mesh() { + return self.flagged(self.tessellate(tol)?); + } + let m = match self.display_plan(tol) { + Ok(plan) => plan.evaluate(), + Err(e) => { + log::debug!("display mesh: {e}; meshing the reference way"); + self.tessellate(tol)? + } + }; + self.flagged(m) } /// Display mesh whose curved faces and edges look curved whatever the body's size: like @@ -1104,6 +1116,30 @@ impl Body { /// sagitta of a small radius on a large body (a sheet-metal bend, a small hole), which the /// tolerance alone leaves as one flat facet. pub fn display_mesh_angular(&self, tol: f64, angle: f64) -> Result { + if self.is_mesh() { + return self.flagged(self.tessellate(tol)?); + } + match self.display_plan_angular(tol, angle) { + Ok(plan) => self.display_mesh_of(&plan), + Err(e) => { + log::debug!("display mesh: {e}; meshing the reference way"); + self.reference_display_mesh_angular(tol, angle) + } + } + } + + /// The display mesh of a plan made from this body ([`Body::display_plan`], + /// [`Body::display_plan_angular`]): evaluated on the CPU, with its edges flagged. + pub fn display_mesh_of(&self, plan: &solidcraft_geom::tess::TessPlan) -> Result { + self.flagged(plan.evaluate()) + } + + /// The reference mesher's angular display mesh: what [`Body::display_mesh_angular`] falls + /// back to when no plan can be made. + pub fn reference_display_mesh_angular(&self, tol: f64, angle: f64) -> Result { + if self.is_mesh() { + return self.flagged(self.tessellate(tol)?); + } self.flagged(self.tessellate_within(tol, angle, false)?) } diff --git a/crates/kernel/src/planar_sources.rs b/crates/kernel/src/planar_sources.rs index e99816e1..e7739561 100644 --- a/crates/kernel/src/planar_sources.rs +++ b/crates/kernel/src/planar_sources.rs @@ -406,7 +406,9 @@ mod tests { assert!(display.contains(Vec3::new(width / 2. - thickness / 2., 0., thickness / 2.))); assert!(display.contains(Vec3::new(width / 2. - thickness / 2., 0., 15.))); assert!(!display.contains(Vec3::new(0., 0., 15.))); - assert_eq!(raw.measure().volume, display.measure().volume); + // Different meshers (the display plan and the reference), both within 0.01 mm. + let (rv, dv) = (raw.measure().volume, display.measure().volume); + assert!((rv - dv).abs() <= raw.measure().area * 0.01, "display volume {dv} vs {rv}"); let expected = (width * 60. - (4. - std::f64::consts::PI) * 36.) * 30. - ((width - 2. * thickness) * (60. - 2. * thickness) - (4. - std::f64::consts::PI) * (6. - thickness).powi(2)) * (30. - thickness); assert!((cup.tessellate(1e-4).unwrap().measure().volume - expected).abs() < 0.2); @@ -549,7 +551,8 @@ mod tests { assert!(!display.contains(Vec3::new(0., 0., 5.))); assert!(display.contains(Vec3::new(28., 0., 5.))); assert!((display.measure().volume - (60. * 54. * 10. - 4. * x * y * 9.)).abs() < 1e-7); - assert_eq!(raw.measure().volume, display.measure().volume); + // Different triangulations of the same polyhedron: equal up to rounding. + assert!((raw.measure().volume - display.measure().volume).abs() < 1e-9 * raw.measure().volume); let imported = crate::step_import(&crate::step_export(&[&pocket], "regularized pocket").unwrap()).unwrap(); assert!(imported.warnings.is_empty()); let copy = &imported.bodies[0].body; diff --git a/crates/kernel/src/smooth_edges.rs b/crates/kernel/src/smooth_edges.rs index 96ca9061..bb24f507 100644 --- a/crates/kernel/src/smooth_edges.rs +++ b/crates/kernel/src/smooth_edges.rs @@ -256,13 +256,15 @@ mod tests { assert!(mesh.shaded_edge_visible(index)); } } + // Smoothing flags are display-only: the edges and faces keep the body's identities. let raw = body.tessellate(0.02).unwrap(); - assert_eq!(mesh.edges, raw.edges); + assert_eq!(mesh.edges.len(), raw.edges.len()); assert_eq!(mesh.edge_faces, raw.edge_faces); - assert_eq!(mesh.tri_face, raw.tri_face); - assert_eq!(mesh.triangles, raw.triangles); - assert_eq!(mesh.positions, raw.positions); - assert_eq!(mesh.normals, raw.normals); + for (a, b) in mesh.edges.iter().zip(&raw.edges) { + assert_eq!((a.first(), a.last()), (b.first(), b.last())); + } + let faces = |m: &Mesh| m.tri_face.iter().copied().collect::>(); + assert_eq!(faces(&mesh), faces(&raw)); let moved = crate::transform_matrix(&body, [[1., 0., 0., 0.], [0., 0., 1., 0.], [0., -1., 0., 0.], [30., 20., 10., 1.]]).unwrap(); assert_eq!(moved.display_mesh(0.02).unwrap().smooth_edges.iter().filter(|f| **f).count(), 512); } diff --git a/crates/ui-egui/Cargo.toml b/crates/ui-egui/Cargo.toml index 5529628f..8a2f84e3 100644 --- a/crates/ui-egui/Cargo.toml +++ b/crates/ui-egui/Cargo.toml @@ -12,6 +12,7 @@ workspace = true solidcraft-engine = { workspace = true } egui = { workspace = true } egui-wgpu = { workspace = true } +solidcraft-gpu = { workspace = true } serde = { workspace = true } serde_json = { workspace = true } log = { workspace = true } @@ -19,5 +20,9 @@ web-time = { workspace = true } image = { workspace = true, features = ["jpeg"] } ttf-parser = { workspace = true } +[target.'cfg(not(target_arch = "wasm32"))'.dev-dependencies] +# The offscreen viewport render test opens a device without a window. +pollster = { workspace = true } + [build-dependencies] serde_json = { workspace = true } diff --git a/crates/ui-egui/src/gpu.rs b/crates/ui-egui/src/gpu.rs index 2a171e50..b605e6ef 100644 --- a/crates/ui-egui/src/gpu.rs +++ b/crates/ui-egui/src/gpu.rs @@ -1,6 +1,8 @@ //! GPU viewport: shaded triangles and screen-space wide lines with a depth buffer, drawn in an //! `egui_wgpu` paint callback. Geometry is uploaded when the scene changes; each frame only the -//! camera uniform changes. +//! camera uniform changes. Where the device runs compute shaders, opaque bodies are drawn +//! straight from their display plans: the plan is evaluated on the GPU into the vertex buffer +//! the pass draws (see [`PlannedBody`]), with nothing built or read back on the CPU. use std::sync::{Arc, Mutex}; @@ -13,6 +15,14 @@ pub const TRI_SIZE: usize = 28; pub const LINE_SIZE: usize = 32; const UNIFORM_SIZE: u64 = 208; +/// A body drawn from its display plan, evaluated on the GPU: the plan, and a colour per face +/// (the plan's face numbering; faces past the end take the last colour). +#[derive(Clone)] +pub struct PlannedBody { + pub plan: Arc, + pub face_colors: Vec<[u8; 4]>, +} + /// CPU-side geometry waiting to be uploaded. The model scene changes with the design; the /// highlight scene (hover, selection, the origin widget) changes often and is small. #[derive(Default, Clone)] @@ -39,6 +49,8 @@ pub struct GpuScene { /// Lines drawn first, under everything (the ground grid: the model hides it, it never /// hides the model). pub under: Vec, + /// Opaque bodies drawn from their plans on the GPU, like `tris`. + pub planned: Vec, } impl GpuScene { @@ -172,6 +184,12 @@ struct Batch { count: u32, } +/// A planned body on the GPU: its evaluated mesh and a colour per vertex. +struct PlannedDraw { + mesh: Arc, + colors: wgpu::Buffer, +} + #[derive(Default)] struct Batches { key: Option, @@ -187,14 +205,24 @@ struct Batches { ghost: Option, xray: Option, under: Option, + planned: Vec, + /// Evaluated plans by the plan's address, holding the plan so the address stays its own: + /// a body whose plan is unchanged is not evaluated again when the scene is rebuilt. + resident: std::collections::HashMap, Arc)>, } impl Batches { - fn take(&mut self, device: &wgpu::Device, key: u64, slot: &SceneSlot) { + /// Upload a new scene. Planned bodies not yet on the GPU are evaluated there: their + /// dispatches go into `encoder`, ahead of this frame's render pass. + fn take(&mut self, device: &wgpu::Device, key: u64, slot: &SceneSlot, tess: Option<(&solidcraft_gpu::TessGpu, &mut wgpu::CommandEncoder)>) { if self.key == Some(key) { return; } if let Some(sc) = slot.lock().ok().and_then(|mut s| s.take()) { + self.planned = match tess { + Some((tess, encoder)) => self.planned(device, tess, encoder, &sc.planned), + None => Vec::new(), + }; self.tris = upload(device, "sc_tris", &sc.tris, TRI_SIZE); self.trans = upload(device, "sc_trans", &sc.trans, TRI_SIZE); self.glass = upload_with(device, "sc_glass", &sc.glass, TRI_SIZE, wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST); @@ -209,10 +237,51 @@ impl Batches { self.key = Some(key); } } + + fn planned( + &mut self, + device: &wgpu::Device, + tess: &solidcraft_gpu::TessGpu, + encoder: &mut wgpu::CommandEncoder, + bodies: &[PlannedBody], + ) -> Vec { + let ptr = |b: &PlannedBody| Arc::as_ptr(&b.plan) as usize; + self.resident.retain(|k, _| bodies.iter().any(|b| ptr(b) == *k)); + let mut out = Vec::with_capacity(bodies.len()); + for b in bodies { + let mesh = match self.resident.get(&ptr(b)) { + Some((_, m)) => m.clone(), + None => match tess.prepare(&b.plan) { + Ok(m) => { + tess.record(encoder, &m); + let m = Arc::new(m); + self.resident.insert(ptr(b), (b.plan.clone(), m.clone())); + m + } + Err(e) => { + log::warn!("a body is not drawn on the GPU: {e}"); + continue; + } + }, + }; + let last = b.face_colors.last().copied().unwrap_or([200, 200, 200, 255]); + let bytes: Vec = mesh.vertex_face.iter().flat_map(|f| b.face_colors.get(*f as usize).copied().unwrap_or(last)).collect(); + if let Some(colors) = upload(device, "sc_planned_colors", &bytes, 4) { + out.push(PlannedDraw { mesh, colors: colors.buffer }); + } + } + out + } } struct Resources { tri: wgpu::RenderPipeline, + /// `tri` and `tri_cut` for planned bodies (positions and normals from the GPU evaluation, + /// colours per vertex, indexed). + tri_planned: wgpu::RenderPipeline, + tri_cut_planned: Option<(wgpu::RenderPipeline, wgpu::RenderPipeline)>, + /// The tessellation compute pipeline, where the device runs compute shaders. + tess: Option, /// While a section is shown (when the depth buffer has a stencil): `tri` marking where the /// cap shows, then the cap again at the section plane's depth, so it hides the inside. tri_cut: Option<(wgpu::RenderPipeline, wgpu::RenderPipeline)>, @@ -602,6 +671,22 @@ fn tri_layout() -> wgpu::VertexBufferLayout<'static> { wgpu::VertexBufferLayout { array_stride: TRI_SIZE as u64, step_mode: wgpu::VertexStepMode::Vertex, attributes: &A } } +/// A planned body's vertices as the GPU evaluation writes them: the position, then the normal +/// at [`solidcraft_gpu::NORMAL_OFFSET`] (each padded to four floats). +fn planned_layout() -> wgpu::VertexBufferLayout<'static> { + const A: [wgpu::VertexAttribute; 2] = [ + wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: 0, shader_location: 0 }, + wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: solidcraft_gpu::NORMAL_OFFSET, shader_location: 1 }, + ]; + wgpu::VertexBufferLayout { array_stride: solidcraft_gpu::VERTEX_STRIDE, step_mode: wgpu::VertexStepMode::Vertex, attributes: &A } +} + +/// A planned body's colours, one per vertex. +fn color_layout() -> wgpu::VertexBufferLayout<'static> { + const A: [wgpu::VertexAttribute; 1] = wgpu::vertex_attr_array![2 => Unorm8x4]; + wgpu::VertexBufferLayout { array_stride: 4, step_mode: wgpu::VertexStepMode::Vertex, attributes: &A } +} + fn line_layout() -> wgpu::VertexBufferLayout<'static> { const A: [wgpu::VertexAttribute; 4] = wgpu::vertex_attr_array![0 => Float32x3, 1 => Float32x3, 2 => Unorm8x4, 3 => Float32]; wgpu::VertexBufferLayout { array_stride: LINE_SIZE as u64, step_mode: wgpu::VertexStepMode::Instance, attributes: &A } @@ -613,10 +698,12 @@ pub struct GpuTarget { pub format: wgpu::TextureFormat, pub depth: Option, pub samples: u32, + /// The device evaluates display plans (compute shaders): bodies may be drawn from them. + pub tess: bool, } impl Resources { - fn new(device: &wgpu::Device, t: GpuTarget) -> Self { + fn new(device: &wgpu::Device, t: GpuTarget, tess: Option) -> Self { let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("sc_viewport"), source: wgpu::ShaderSource::Wgsl(SHADER.into()) }); let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { @@ -770,14 +857,16 @@ impl Resources { bind_group_layouts: &[Some(&bgl), Some(&access_bgl)], immediate_size: 0, }); - let solid = |label: &str, fs: &str, cull: Option, ds: Option| { + let built = [Some(tri_layout())]; + let evaluated = [Some(planned_layout()), Some(color_layout())]; + let solid = |label: &str, fs: &str, cull: Option, ds: Option, planned: bool| { device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { label: Some(label), layout: Some(&solid_layout), vertex: wgpu::VertexState { module: &module, entry_point: Some("vs_tri"), - buffers: &[Some(tri_layout())], + buffers: if planned { &evaluated[..] } else { &built[..] }, compilation_options: Default::default(), }, primitive: wgpu::PrimitiveState { topology: wgpu::PrimitiveTopology::TriangleList, cull_mode: cull, ..Default::default() }, @@ -793,21 +882,9 @@ impl Resources { cache: None, }) }; - let access_bind = access_texture(device, None, &access_bgl, 1, &[]); - let sampler = device.create_sampler(&wgpu::SamplerDescriptor { - label: Some("sc_image"), - mag_filter: wgpu::FilterMode::Linear, - min_filter: wgpu::FilterMode::Linear, - ..Default::default() - }); - Resources { - image, - image_bgl, - sampler, - textures: Default::default(), - quads: Vec::new(), - tri: solid("sc_tris", "fs_solid", None, depth(wgpu::CompareFunction::Less, true)), - tri_cut: t.depth.filter(|f| f.has_stencil_aspect()).map(|format| { + // While a section is shown: the stencil marking and the cap (see `Resources::tri_cut`). + let cut = |planned: bool| { + t.depth.filter(|f| f.has_stencil_aspect()).map(|format| { // The visible surface marks the stencil: 1 where it is a back face (the cap). use wgpu::StencilOperation::{Keep, Replace, Zero}; let face = |pass_op| wgpu::StencilFaceState { compare: wgpu::CompareFunction::Always, fail_op: Keep, depth_fail_op: Keep, pass_op }; @@ -827,8 +904,30 @@ impl Resources { stencil: wgpu::StencilState { front: only, back: only, read_mask: 0xff, write_mask: 0 }, bias: Default::default(), }; - (solid("sc_tris_mark", "fs_solid", None, Some(mark)), solid("sc_tris_cap", "fs_solid_cut", Some(wgpu::Face::Front), Some(cap))) - }), + ( + solid("sc_tris_mark", "fs_solid", None, Some(mark), planned), + solid("sc_tris_cap", "fs_solid_cut", Some(wgpu::Face::Front), Some(cap), planned), + ) + }) + }; + let access_bind = access_texture(device, None, &access_bgl, 1, &[]); + let sampler = device.create_sampler(&wgpu::SamplerDescriptor { + label: Some("sc_image"), + mag_filter: wgpu::FilterMode::Linear, + min_filter: wgpu::FilterMode::Linear, + ..Default::default() + }); + Resources { + image, + image_bgl, + sampler, + textures: Default::default(), + quads: Vec::new(), + tri: solid("sc_tris", "fs_solid", None, depth(wgpu::CompareFunction::Less, true), false), + tri_cut: cut(false), + tri_planned: solid("sc_tris_planned", "fs_solid", None, depth(wgpu::CompareFunction::Less, true), true), + tri_cut_planned: cut(true), + tess, access_bgl, access_bind, access_version: 0, @@ -894,8 +993,13 @@ impl Batches { /// Create the pipelines. Call once with the app's render state and its depth/stencil/MSAA /// settings. Without a stencil the section cap does not hide the inside of cut bodies. pub fn install(rs: &egui_wgpu::RenderState, depth_bits: u8, stencil_bits: u8, samples: u32) -> GpuTarget { - let t = GpuTarget { format: rs.target_format, depth: egui_wgpu::depth_format_from_bits(depth_bits, stencil_bits), samples }; - let res = Resources::new(&rs.device, t); + // Display plans are evaluated where the device runs compute shaders with the storage + // buffers the evaluation binds (not WebGL2): elsewhere bodies are built on the CPU. + let compute = rs.adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) + && rs.device.limits().max_storage_buffers_per_shader_stage >= 4; + let tess = compute.then(|| solidcraft_gpu::TessGpu::new(rs.adapter.get_info().name, rs.device.clone(), rs.queue.clone())); + let t = GpuTarget { format: rs.target_format, depth: egui_wgpu::depth_format_from_bits(depth_bits, stencil_bits), samples, tess: tess.is_some() }; + let res = Resources::new(&rs.device, t, tess); rs.renderer.write().callback_resources.insert(res); t } @@ -906,14 +1010,14 @@ impl egui_wgpu::CallbackTrait for ViewportCallback { device: &wgpu::Device, queue: &wgpu::Queue, _screen: &egui_wgpu::ScreenDescriptor, - _encoder: &mut wgpu::CommandEncoder, + encoder: &mut wgpu::CommandEncoder, resources: &mut egui_wgpu::CallbackResources, ) -> Vec { let Some(res) = resources.get_mut::() else { return Vec::new() }; - res.model.take(device, self.key, &self.slot); - res.sketch.take(device, self.sketch_key, &self.sketch_slot); - res.highlight.take(device, self.hl_key, &self.hl_slot); - res.preview.take(device, self.pv_key, &self.pv_slot); + res.model.take(device, self.key, &self.slot, res.tess.as_ref().map(|t| (t, &mut *encoder))); + res.sketch.take(device, self.sketch_key, &self.sketch_slot, None); + res.highlight.take(device, self.hl_key, &self.hl_slot, None); + res.preview.take(device, self.pv_key, &self.pv_slot, None); res.model.sort_glass(queue, self.back); queue.write_buffer(&res.uniform, 0, &uniform_bytes(self, self.size_px[0], self.size_px[1], res.linear_out)); if let Some(a) = &self.access @@ -1000,18 +1104,30 @@ impl egui_wgpu::CallbackTrait for ViewportCallback { lines(pass, &res.overlay, &m.under); lines(pass, &res.overlay, &sk.under); pass.set_bind_group(1, &res.access_bind, &[]); - let cut = res.tri_cut.as_ref().filter(|_| self.clip[..3].iter().map(|c| c * c).sum::() > 0.25); - match cut { - Some((mark, cap)) => { + let planned = |pass: &mut wgpu::RenderPass<'static>, p: &wgpu::RenderPipeline| { + for d in &m.planned { + pass.set_pipeline(p); + pass.set_vertex_buffer(0, d.mesh.vertices.slice(..)); + pass.set_vertex_buffer(1, d.colors.slice(..)); + pass.set_index_buffer(d.mesh.indices.slice(..), wgpu::IndexFormat::Uint32); + pass.draw_indexed(0..d.mesh.index_count, 0, 0..1); + } + }; + let sectioned = self.clip[..3].iter().map(|c| c * c).sum::() > 0.25; + match (res.tri_cut.as_ref().filter(|_| sectioned), res.tri_cut_planned.as_ref()) { + (Some((mark, cap)), Some((mark_planned, cap_planned))) => { pass.set_stencil_reference(1); tris(pass, mark, &m.tris); tris(pass, mark, &pv.tris); + planned(pass, mark_planned); tris(pass, cap, &m.tris); tris(pass, cap, &pv.tris); + planned(pass, cap_planned); } - None => { + _ => { tris(pass, &res.tri, &m.tris); tris(pass, &res.tri, &pv.tris); + planned(pass, &res.tri_planned); } } tris(pass, &res.tri_hl, &h.tris); @@ -1042,6 +1158,108 @@ impl egui_wgpu::CallbackTrait for ViewportCallback { } } +/// Render `scene` offscreen on a device of its own, through the viewport's real resources and +/// frame callback (`prepare`, then `paint`), and read the pixels back: straight RGBA, `size` × +/// `size`. `None` without a GPU adapter. With `tess`, planned bodies are evaluated on the GPU. +#[cfg(all(test, not(target_arch = "wasm32")))] +pub(crate) fn render_offscreen(scene: GpuScene, view_proj: [[f32; 4]; 4], back: [f32; 3], eye: [f32; 4], size: u32, tess: bool) -> Option> { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); + let adapter = pollster::block_on(instance.request_adapter(&Default::default())).ok()?; + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), + ..Default::default() + })) + .ok()?; + let (format, depth) = (wgpu::TextureFormat::Rgba8Unorm, wgpu::TextureFormat::Depth24PlusStencil8); + let target = GpuTarget { format, depth: Some(depth), samples: 1, tess }; + let compute = tess.then(|| solidcraft_gpu::TessGpu::new(adapter.get_info().name, device.clone(), queue.clone())); + let mut resources = egui_wgpu::CallbackResources::default(); + resources.insert(Resources::new(&device, target, compute)); + let none = || Arc::new(Mutex::new(None)); + let cb = ViewportCallback { + key: 1, + slot: Arc::new(Mutex::new(Some(scene))), + sketch_key: 0, + sketch_slot: none(), + hl_key: 0, + hl_slot: none(), + pv_key: 0, + pv_slot: none(), + view_proj, + back, + size_px: [size as f32; 2], + clip: [0.0; 4], + cap: [0.0; 4], + images: Vec::new(), + analysis: [0.0; 8], + access: None, + outline: false, + eye, + }; + let texture = |format, usage| { + device.create_texture(&wgpu::TextureDescriptor { + label: Some("offscreen"), + size: wgpu::Extent3d { width: size, height: size, depth_or_array_layers: 1 }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format, + usage, + view_formats: &[], + }) + }; + let color = texture(format, wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC); + let depth_tex = texture(depth, wgpu::TextureUsages::RENDER_ATTACHMENT); + let (color_view, depth_view) = (color.create_view(&Default::default()), depth_tex.create_view(&Default::default())); + let mut encoder = device.create_command_encoder(&Default::default()); + let screen = egui_wgpu::ScreenDescriptor { size_in_pixels: [size, size], pixels_per_point: 1.0 }; + let extra = egui_wgpu::CallbackTrait::prepare(&cb, &device, &queue, &screen, &mut encoder, &mut resources); + { + let mut pass = encoder + .begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("offscreen"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: &color_view, + depth_slice: None, + resolve_target: None, + ops: wgpu::Operations { load: wgpu::LoadOp::Clear(wgpu::Color::WHITE), store: wgpu::StoreOp::Store }, + })], + depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { + view: &depth_view, + depth_ops: Some(wgpu::Operations { load: wgpu::LoadOp::Clear(1.0), store: wgpu::StoreOp::Store }), + stencil_ops: Some(wgpu::Operations { load: wgpu::LoadOp::Clear(0), store: wgpu::StoreOp::Store }), + }), + timestamp_writes: None, + occlusion_query_set: None, + multiview_mask: None, + }) + .forget_lifetime(); + let rect = egui::Rect::from_min_size(egui::Pos2::ZERO, egui::vec2(size as f32, size as f32)); + let info = egui::PaintCallbackInfo { viewport: rect, clip_rect: rect, pixels_per_point: 1.0, screen_size_px: [size, size] }; + egui_wgpu::CallbackTrait::paint(&cb, info, &mut pass, &resources); + } + let read = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("offscreen read"), + size: u64::from(size * size * 4), + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + encoder.copy_texture_to_buffer( + color.as_image_copy(), + wgpu::TexelCopyBufferInfo { + buffer: &read, + layout: wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(size * 4), rows_per_image: Some(size) }, + }, + wgpu::Extent3d { width: size, height: size, depth_or_array_layers: 1 }, + ); + queue.submit(extra.into_iter().chain([encoder.finish()])); + let slice = read.slice(..); + slice.map_async(wgpu::MapMode::Read, |_| {}); + device.poll(wgpu::PollType::wait_indefinitely()).ok()?; + let bytes = slice.get_mapped_range().ok()?.to_vec(); + Some(bytes) +} + /// Uniform bytes for a frame. pub fn uniform_bytes(cb: &ViewportCallback, w: f32, h: f32, linear: bool) -> Vec { let mut v: Vec = cb.view_proj.iter().flatten().copied().collect(); diff --git a/crates/ui-egui/src/hole_placement_ui_tests.rs b/crates/ui-egui/src/hole_placement_ui_tests.rs index 849c61a8..d2c44212 100644 --- a/crates/ui-egui/src/hole_placement_ui_tests.rs +++ b/crates/ui-egui/src/hole_placement_ui_tests.rs @@ -435,14 +435,14 @@ fn cylinder_with_cluster(sketch: Value, at: [f64; 2], view: &str) -> Cluster { #[test] fn select_other_over_a_dense_cluster_looks_the_face_up_once() { - // The points on the cylinder's top face, a disc of a tenth of the body's triangles. Select + // The points on the cylinder's top face, a disc of a small part of the body's triangles. Select // Other looks the face under the cursor up once for all of them, not once per point (#718 // review): one planarity lookup, which tessellates the body, and one pass over the body's // triangles. Each point is then tested against the face's triangles near the points only, // not the face's or the body's; every point on the face is still counted. let mut c = cylinder_with_cluster(json!({"plane": {"face": [0, 0, 20]}}), [5.0, 4.0], "top"); let (many, body, face, near) = (c.many, c.body_tris, c.face_tris, c.near_tris); - assert!(body >= 8 * face && 4 * near <= face, "{body} triangles in the body, {face} in the face, {near} of them near the points"); + assert!(body >= 4 * face && 4 * near <= face, "{body} triangles in the body, {face} in the face, {near} of them near the points"); let (list, (lookups, tested)) = c.select_other(); assert_eq!((rows_named(&list, "Point"), rows_named(&list, "Face")), (crate::select_other::MAX_PER_KIND, 1), "{list}"); assert_eq!(list["total"], json!(many + 1), "every point on the face is counted: {list}"); diff --git a/crates/ui-egui/src/viewport.rs b/crates/ui-egui/src/viewport.rs index 432927d9..e2df0674 100644 --- a/crates/ui-egui/src/viewport.rs +++ b/crates/ui-egui/src/viewport.rs @@ -338,7 +338,24 @@ fn build_scene(app: &SolveApp) -> GpuScene { let opacity = s.doc.model_body_opacity(b); let face_looks: std::collections::HashMap = if shaded { s.doc.face_colors(b).into_iter().collect() } else { Default::default() }; - for (ti, t) in m.triangles.iter().enumerate().filter(|_| shaded) { + // Opaque bodies are drawn from their display plans on the GPU where the device evaluates + // them; faded and see-through ones (sorted and blended per triangle) are built here. + let opaque = !faded && opacity >= 0.999 && face_looks.values().all(|l| l.opacity >= 0.999); + let planned = (shaded && opaque && app.viewport.gpu.is_some_and(|g| g.tess)) + .then(|| b.display_plan()) + .flatten() + .filter(|plan| plan.vertex_count() <= solidcraft_gpu::MAX_BATCH); + let on_gpu = planned.is_some(); + if let Some(plan) = planned { + let face_colors = (0..b.body.face_count()) + .map(|f| match face_looks.get(&f) { + Some(l) => c4(Color32::from_rgb(l.color[0], l.color[1], l.color[2])), + None => col, + }) + .collect(); + sc.planned.push(crate::gpu::PlannedBody { plan, face_colors }); + } + for (ti, t) in m.triangles.iter().enumerate().filter(|_| shaded && !on_gpu) { let (mut col, a) = match m.tri_face.get(ti).and_then(|f| face_looks.get(&(*f as usize))) { Some(l) => { let c = Color32::from_rgb(l.color[0], l.color[1], l.color[2]); @@ -3189,6 +3206,44 @@ fn view_menu(app: &mut SolveApp, ui: &mut egui::Ui, button: Rect) -> bool { mod tests { use super::*; + /// Bodies drawn from their display plans, evaluated on the GPU, look like the same bodies + /// built into triangles on the CPU: the sample design rendered both ways through the + /// viewport's real pipelines, offscreen, on this machine's GPU (skipped without one). + #[cfg(not(target_arch = "wasm32"))] + #[test] + fn bodies_evaluated_on_the_gpu_look_like_bodies_built_on_the_cpu() { + let mut s = solidcraft_engine::Session::default(); + s.run_script(&solidcraft_engine::sample::script()).unwrap(); + let mut app = SolveApp::new(s, crate::Services::default()); + app.viewport.no_pixels = true; + app.fit_view(); + const SIZE: u32 = 256; + let rect = Rect::from_min_size(Pos2::ZERO, vec2(SIZE as f32, SIZE as f32)); + let mut render = |tess: bool| { + app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess }); + let scene = build_scene(&app); + let proj = projection(&app, rect); + let e = proj.cam.eye().to_f32(); + let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }]; + let planned = scene.planned.len(); + (planned, scene.tris.len(), crate::gpu::render_offscreen(scene, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, tess)) + }; + let (planned, gpu_tris, gpu) = render(true); + let (none, cpu_tris, cpu) = render(false); + let (Some(gpu), Some(cpu)) = (gpu, cpu) else { + eprintln!("skipping: no GPU adapter"); + return; + }; + assert!(planned > 0 && none == 0, "{planned} bodies planned"); + assert!(gpu_tris < cpu_tris, "the planned bodies' triangles are not built on the CPU ({gpu_tris} vs {cpu_tris} bytes)"); + let px = |img: &[u8], i: usize| [img[4 * i], img[4 * i + 1], img[4 * i + 2]]; + let n = (SIZE * SIZE) as usize; + let drawn = (0..n).filter(|&i| px(&cpu, i) != [255, 255, 255]).count(); + assert!(drawn > n / 20, "the model covers the view: {drawn} pixels"); + let differ = (0..n).filter(|&i| px(&gpu, i).iter().zip(px(&cpu, i)).any(|(a, b)| a.abs_diff(b) > 24)).count(); + assert!(differ * 100 < drawn, "{differ} of {drawn} model pixels differ"); + } + #[test] fn construction_drawing_and_hover_overlays_preserve_visible_gaps() { fn length(shape: &Shape) -> f32 { @@ -3383,7 +3438,7 @@ mod tests { app.viewport.no_pixels = true; if gpu { // The callback is queued by egui; no device is needed to inspect scene updates. - app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1 }); + app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess: false }); } app } diff --git a/docs/perf.md b/docs/perf.md index 99b78b92..bd92bb6c 100644 --- a/docs/perf.md +++ b/docs/perf.md @@ -40,6 +40,31 @@ and 444 MB either way, so two runs each do not resolve it: there the conversion dominates the per-item copy. An item that keeps every conversion - an assembly, or a file over the work budget - is unchanged. +## Display tessellation plans (#303, 2026-10-12) + +`Body::display_mesh` meshes from a `TessPlan` (`crates/kernel/src/tess`): every edge is +sampled once and both faces use those nodes exactly (watertight by construction); planar +faces are a constrained Delaunay triangulation of their edge nodes alone; faces bounded by +their parameter box are structured grids zipped to the edge nodes; other four-cornered faces +are Coons grids; only genuinely trimmed curved faces (holes, more than four corners) use +spade's constrained Delaunay with sag and Ruppert refinement. A body the planner cannot +handle is meshed by truck-meshalgo as before, which also still meshes for modelling +(booleans, measure, provenance). Evaluating a plan (surface points and normals) runs on the +CPU (`TessPlan::evaluate`) or in a compute shader (`solidcraft-gpu`). + +Release build, Apple M4 Max, medians of five. 400-hole plate (200 × 200 × 5 mm, 806 +faces), `cargo test --release -p solidcraft-gpu bench_plate -- --ignored --nocapture`: + +| Tolerance | Reference mesher | Plan | CPU evaluation | GPU evaluation | Triangles (reference / plan) | +|---|---:|---:|---:|---:|---:| +| 0.3 mm | 23.3 ms | 11.9 ms | 1.1 ms | 1.1 ms | 20,812 / 20,812 | +| 0.03 mm | 54.2 ms | 31.6 ms | 2.5 ms | 1.9 ms | 76,812 / 52,812 | +| 0.006 mm | 118.9 ms | 65.1 ms | 5.3 ms | 3.1 ms | 302,412 / 104,012 | + +Planning (edge sampling, parameter projection, triangulation) dominates; evaluation is a +few milliseconds either way, so GPU evaluation saves little until planning is cheaper. The +GPU timing includes upload and read-back. + ## Parallel STEP body preparation (2026-10-12) Profiling an original 48-cylinder STEP fixture before changing the importer measured diff --git a/xtask/src/layers.rs b/xtask/src/layers.rs index 3d5df575..3b482150 100644 --- a/xtask/src/layers.rs +++ b/xtask/src/layers.rs @@ -43,6 +43,7 @@ pub const TABLE: &[(&str, Class)] = &[ ("io", Class::Layer(3)), ("engine", Class::Layer(4)), ("ui-egui", Class::Layer(5)), + ("gpu", Class::Layer(5)), ("mcp", Class::Layer(5)), ("testkit", Class::Testkit), // apps and tooling @@ -53,7 +54,8 @@ pub const TABLE: &[(&str, Class)] = &[ ]; /// Explicit orderings *within* a layer (earlier may be used by later). -pub const INTRA_LAYER_ORDER: &[&[&str]] = &[]; +/// The viewport draws display plans the GPU crate evaluates. +pub const INTRA_LAYER_ORDER: &[&[&str]] = &[&["gpu", "ui-egui"]]; fn intra_layer_allowed(from: &str, to: &str) -> bool { let (from, to) = (short_name(from), short_name(to));