1b-agent/OneCAD
1
0 Fork 0
forked from 1b-agent/solidcraft-upstream
Code Pull Requests Activity

Display tessellation plans (3/4): display meshes from plans, GPU evaluation in the viewport (#805)

Browse Source 
Display meshes are built from tessellation plans, with GPU evaluation in the viewport.
This commit is contained in:
Irakli Kourtis authored and GitHub committed 2026-10-13 16:27:16 -04:00
1 parent 534393ed5a
commit ff7a347013
19 files changed
+1474 -77
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

crates/gpu/src/lib.rs
+442
View File
@@ -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<T> = std::result::Result<T, GpuError>;
/// 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<u32>,
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<String>, 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<Self> {
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<Mesh> {
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<GpuMesh> {
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::<Vec<u8>>();
let floats = |v: &[f32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::<Vec<u8>>();
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<u32> = 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<Vec<f32>> {
let d = &self.device;
let bytes_u32 = |v: &[u32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::<Vec<u8>>();
let bytes_f32 = |v: &[f32]| v.iter().flat_map(|x| x.to_le_bytes()).collect::<Vec<u8>>();
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<u32>,
data: Vec<f32>,
vertices: Vec<u32>,
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> {
u32::try_from(x).map_err(|_| GpuError::Failed("plan too large".into()))
}
fn push3(data: &mut Vec<f32>, v: Vec3) {
data.extend([v.x as f32, v.y as f32, v.z as f32]);
}
fn push_curve(data: &mut Vec<f32>, 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<Packed> {
let mut surfaces = Vec::with_capacity(plan.surfaces.len() * SURFACE_WORDS);
let mut data: Vec<f32> = 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;
crates/gpu/src/tess.wgsl
+289
View File
@@ -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<f32>,
};
struct Params {
count: u32,
_p0: u32,
_p1: u32,
_p2: u32,
};
@group(0) @binding(0) var<uniform> params: Params;
@group(0) @binding(1) var<storage, read> surfaces: array<Surface>;
@group(0) @binding(2) var<storage, read> data: array<f32>;
@group(0) @binding(3) var<storage, read> vertices: array<Vertex>;
// Out: position (xyz, w = 1 when evaluated) and normal (xyz) per vertex.
@group(0) @binding(4) var<storage, read_write> out: array<vec4<f32>>;
struct Derivs {
p: vec3<f32>,
du: vec3<f32>,
dv: vec3<f32>,
ok: bool,
};
struct Basis {
n0: array<f32, 8>,
n1: array<f32, 8>,
};
fn vec3_at(o: u32) -> vec3<f32> {
return vec3<f32>(data[o], data[o + 1u], data[o + 2u]);
}
fn ctrl_at(o: u32) -> vec4<f32> {
return vec4<f32>(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<array<f32, 8>, 8>;
var left: array<f32, 8>;
var right: array<f32, 8>;
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<f32>) -> vec3<f32> {
return a.xyz / a.w;
}
fn rational_der(a: vec4<f32>, da: vec4<f32>) -> vec3<f32> {
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<f32>(0.0);
var au = vec4<f32>(0.0);
var av = vec4<f32>(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<f32>, k: vec3<f32>, sn: f32, cs: f32) -> vec3<f32> {
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<f32>(0.0);
var dc = vec4<f32>(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<f32>(data[m], data[m + 1u], data[m + 2u], data[m + 3u]);
let row1 = vec4<f32>(data[m + 4u], data[m + 5u], data[m + 6u], data[m + 7u]);
let row2 = vec4<f32>(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<f32>(dot(row0.xyz, p) + row0.w, dot(row1.xyz, p) + row1.w, dot(row2.xyz, p) + row2.w);
let ldt = vec3<f32>(dot(row0.xyz, dt), dot(row1.xyz, dt), dot(row2.xyz, dt));
let lda = vec3<f32>(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<f32>) -> vec3<f32> {
if (s.mirror != 0u) {
return -n;
}
return n;
}
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
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<vec2<f32>, 4>(vec2<f32>(h, h), vec2<f32>(-h, h), vec2<f32>(h, -h), vec2<f32>(-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<f32>(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<f32>(p, 1.0);
out[2u * i + 1u] = vec4<f32>(nn, 0.0);
}
crates/gpu/src/tests.rs
+168
View File
@@ -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<TessGpu> {
match TessGpu::headless() {
Ok(g) => {
eprintln!("GPU: {}", g.adapter_name());
Some(g)
}
Err(e) => {
eprintln!("skipping the GPU comparison: {e}");
None
}
}
}
fn ok<T, E: std::fmt::Debug>(r: std::result::Result<T, E>) -> 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::<Vec<f32>>()
};
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<Loop2> = (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<f64>| {
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));
}
}