diff --git a/crates/geom/src/bvh.rs b/crates/geom/src/bvh.rs
new file mode 100644
index 00000000..24c5dc80
--- /dev/null
+++ b/crates/geom/src/bvh.rs
@@ -0,0 +1,270 @@
+//! A bounding volume hierarchy over boxes: the nearest hit along a ray, or every item whose box
+//! passes a test, without looking at each item. Built once (median split on the widest axis of
+//! the box centres, a few items per leaf) and only read afterwards. [`MeshIndex`] holds one over
+//! a mesh's triangles and one over its edge segments, for picking.
+
+use crate::{Aabb3, Mesh, Vec3};
+
+/// Most items in a leaf.
+const LEAF: usize = 4;
+
+#[derive(Clone, Copy, Debug)]
+struct Node {
+    bounds: Aabb3,
+    /// A leaf: the first of its `count` items in [`Bvh::order`]. Otherwise its children are
+    /// `first` and `first + 1`.
+    first: u32,
+    count: u32,
+}
+
+/// A hierarchy over boxes; items are the boxes' indices.
+#[derive(Clone, Debug, Default)]
+pub struct Bvh {
+    nodes: Vec<Node>,
+    order: Vec<u32>,
+}
+
+impl Bvh {
+    /// The hierarchy over `boxes` (empty boxes are left out).
+    pub fn new(boxes: &[Aabb3]) -> Bvh {
+        let mut order: Vec<u32> =
+            (0..boxes.len().min(u32::MAX as usize)).filter(|&i| boxes.get(i).is_some_and(|b| !b.is_empty() && finite(b))).map(|i| i as u32).collect();
+        if order.is_empty() {
+            return Bvh::default();
+        }
+        let item = |i: u32| boxes.get(i as usize).copied().unwrap_or(Aabb3::EMPTY);
+        let centre = |i: u32| item(i).center();
+        let mut nodes = vec![Node { bounds: Aabb3::EMPTY, first: 0, count: 0 }];
+        // (node, start, end) in `order` still to be filled in.
+        let mut todo = vec![(0usize, 0usize, order.len())];
+        while let Some((n, start, end)) = todo.pop() {
+            let Some(span) = order.get_mut(start..end) else { continue };
+            let mut bounds = Aabb3::EMPTY;
+            let mut centres = Aabb3::EMPTY;
+            for &i in span.iter() {
+                bounds = bounds.union(&item(i));
+                centres.add(centre(i));
+            }
+            let size = centres.size();
+            let axis = if size.x >= size.y && size.x >= size.z {
+                0
+            } else if size.y >= size.z {
+                1
+            } else {
+                2
+            };
+            let extent = [size.x, size.y, size.z][axis];
+            if span.len() <= LEAF || !(extent > 0.0) {
+                if let Some(node) = nodes.get_mut(n) {
+                    *node = Node { bounds, first: start as u32, count: span.len() as u32 };
+                }
+                continue;
+            }
+            let key = |i: &u32| {
+                let c = centre(*i);
+                [c.x, c.y, c.z][axis]
+            };
+            let mid = span.len() / 2;
+            span.select_nth_unstable_by(mid, |a, b| key(a).total_cmp(&key(b)));
+            let left = nodes.len();
+            nodes.push(Node { bounds: Aabb3::EMPTY, first: 0, count: 0 });
+            nodes.push(Node { bounds: Aabb3::EMPTY, first: 0, count: 0 });
+            if let Some(node) = nodes.get_mut(n) {
+                *node = Node { bounds, first: left as u32, count: 0 };
+            }
+            todo.push((left, start, start + mid));
+            todo.push((left + 1, start + mid, end));
+        }
+        Bvh { nodes, order }
+    }
+
+    /// The box around everything (empty for no items).
+    pub fn bounds(&self) -> Aabb3 {
+        self.nodes.first().map_or(Aabb3::EMPTY, |n| n.bounds)
+    }
+
+    /// Calls `item` with every item in a leaf whose box, and every box above it, `enter` accepts.
+    pub fn visit(&self, mut enter: impl FnMut(&Aabb3) -> bool, mut item: impl FnMut(u32)) {
+        let mut stack = vec![0usize];
+        while let Some(n) = stack.pop() {
+            let Some(node) = self.nodes.get(n) else { continue };
+            if !enter(&node.bounds) {
+                continue;
+            }
+            if node.count > 0 {
+                let (a, b) = (node.first as usize, node.first as usize + node.count as usize);
+                for &i in self.order.get(a..b).unwrap_or(&[]) {
+                    item(i);
+                }
+            } else {
+                stack.push(node.first as usize + 1);
+                stack.push(node.first as usize);
+            }
+        }
+    }
+
+    /// The nearest item along the ray from `o` along `d`: `hit` gives an item's distance along
+    /// `d` (`None`: missed). Equal distances go to the lowest item, as a scan in item order
+    /// keeping strictly nearer hits would answer. Boxes are entered nearest first, and one that
+    /// starts past the best hit so far is skipped.
+    pub fn nearest(&self, o: Vec3, d: Vec3, mut hit: impl FnMut(u32) -> Option<f64>) -> Option<(f64, u32)> {
+        if ![o.x, o.y, o.z, d.x, d.y, d.z].iter().all(|v| v.is_finite()) {
+            return None;
+        }
+        let inv = Vec3::new(1.0 / d.x, 1.0 / d.y, 1.0 / d.z);
+        let mut best: Option<(f64, u32)> = None;
+        let root = self.nodes.first().and_then(|n| ray_box(o, d, inv, &n.bounds, 0.0))?;
+        let mut stack = vec![(0usize, root)];
+        while let Some((n, t0)) = stack.pop() {
+            if best.is_some_and(|(bt, _)| t0 > bt) {
+                continue;
+            }
+            let Some(node) = self.nodes.get(n) else { continue };
+            if node.count > 0 {
+                let (a, b) = (node.first as usize, node.first as usize + node.count as usize);
+                for &i in self.order.get(a..b).unwrap_or(&[]) {
+                    if let Some(t) = hit(i)
+                        && best.is_none_or(|(bt, bi)| t < bt || (t == bt && i < bi))
+                    {
+                        best = Some((t, i));
+                    }
+                }
+                continue;
+            }
+            let (l, r) = (node.first as usize, node.first as usize + 1);
+            let tl = self.nodes.get(l).and_then(|c| ray_box(o, d, inv, &c.bounds, 0.0));
+            let tr = self.nodes.get(r).and_then(|c| ray_box(o, d, inv, &c.bounds, 0.0));
+            // The nearer child is popped first.
+            match (tl, tr) {
+                (Some(a), Some(b)) if a <= b => stack.extend([(r, b), (l, a)]),
+                (Some(a), Some(b)) => stack.extend([(l, a), (r, b)]),
+                (Some(a), None) => stack.push((l, a)),
+                (None, Some(b)) => stack.push((r, b)),
+                (None, None) => {}
+            }
+        }
+        best
+    }
+}
+
+fn finite(b: &Aabb3) -> bool {
+    [b.min.x, b.min.y, b.min.z, b.max.x, b.max.y, b.max.z].iter().all(|v| v.is_finite())
+}
+
+/// Where the ray from `o` along `d` (`inv` = 1 / `d` per axis) enters box `b` grown by `pad`,
+/// from 0 on; `None` when it misses or the box lies behind.
+fn ray_box(o: Vec3, d: Vec3, inv: Vec3, b: &Aabb3, pad: f64) -> Option<f64> {
+    let (mut t0, mut t1) = (0.0f64, f64::INFINITY);
+    for (o, d, inv, lo, hi) in [(o.x, d.x, inv.x, b.min.x, b.max.x), (o.y, d.y, inv.y, b.min.y, b.max.y), (o.z, d.z, inv.z, b.min.z, b.max.z)] {
+        let (lo, hi) = (lo - pad, hi + pad);
+        if d == 0.0 || !inv.is_finite() {
+            if o < lo || o > hi {
+                return None;
+            }
+            continue;
+        }
+        let (a, b) = ((lo - o) * inv, (hi - o) * inv);
+        let (a, b) = if a <= b { (a, b) } else { (b, a) };
+        t0 = t0.max(a);
+        t1 = t1.min(b);
+        if t0 > t1 {
+            return None;
+        }
+    }
+    (t0 <= t1).then_some(t0)
+}
+
+/// The farthest any corner of `b` lies along `d` from `o` (0 when the box is behind).
+fn farthest_along(o: Vec3, d: Vec3, b: &Aabb3) -> f64 {
+    let along = |lo: f64, hi: f64, o: f64, d: f64| if d >= 0.0 { (hi - o) * d } else { (lo - o) * d };
+    (along(b.min.x, b.max.x, o.x, d.x) + along(b.min.y, b.max.y, o.y, d.y) + along(b.min.z, b.max.z, o.z, d.z)).max(0.0)
+}
+
+/// A mesh's triangles and edge segments in hierarchies, for picking: the nearest triangle along
+/// a ray, and the edge segments near one. Built from one mesh and answers only for it.
+#[derive(Clone, Debug, Default)]
+pub struct MeshIndex {
+    triangles: Bvh,
+    triangle_count: usize,
+    segments: Bvh,
+    /// Per segment item: its edge, and its first point in that edge's polyline.
+    segment_of: Vec<(u32, u32)>,
+}
+
+impl MeshIndex {
+    /// The index of `mesh`.
+    pub fn new(mesh: &Mesh) -> MeshIndex {
+        let pad = |mut b: Aabb3| {
+            // Flat boxes (a triangle square to an axis) are grown a little, so a ray through
+            // them is never lost to rounding in the box test.
+            let scale = [b.min.x, b.min.y, b.min.z, b.max.x, b.max.y, b.max.z].iter().fold(1.0f64, |m, v| m.max(v.abs()));
+            let e = scale * 1e-9;
+            b.min -= Vec3::new(e, e, e);
+            b.max += Vec3::new(e, e, e);
+            b
+        };
+        let boxes: Vec<Aabb3> = mesh
+            .triangles
+            .iter()
+            .map(|t| {
+                let mut b = Aabb3::EMPTY;
+                for p in mesh.tri(t).into_iter().flatten() {
+                    b.add(p);
+                }
+                if b.is_empty() { b } else { pad(b) }
+            })
+            .collect();
+        let mut segment_of = Vec::new();
+        let mut seg_boxes = Vec::new();
+        for (ei, e) in mesh.edges.iter().enumerate() {
+            // A one-point edge still has its vertex: it stands as a segment from it to itself.
+            let n = if e.len() == 1 { 1 } else { e.len().saturating_sub(1) };
+            for si in 0..n {
+                let (Some(a), Some(b)) = (e.get(si), e.get(si + 1).or(e.get(si))) else { continue };
+                let mut bx = Aabb3::EMPTY;
+                bx.add(*a);
+                bx.add(*b);
+                seg_boxes.push(pad(bx));
+                segment_of.push((ei as u32, si as u32));
+            }
+        }
+        MeshIndex { triangles: Bvh::new(&boxes), triangle_count: mesh.triangles.len(), segments: Bvh::new(&seg_boxes), segment_of }
+    }
+
+    /// [`Mesh::raycast`] through the index: the same answer, including which triangle wins a tie.
+    /// For a mesh other than the one indexed, the plain scan answers.
+    pub fn raycast(&self, mesh: &Mesh, o: Vec3, dir: Vec3) -> Option<(f64, usize)> {
+        if mesh.triangles.len() != self.triangle_count {
+            return mesh.raycast(o, dir);
+        }
+        let frame = crate::mesh::RayFrame::new(o, dir)?;
+        self.triangles
+            .nearest(o, dir, |i| mesh.triangles.get(i as usize).and_then(|t| mesh.tri(t)).and_then(|p| frame.hit(p)))
+            .map(|(t, i)| (t, i as usize))
+    }
+
+    /// Calls `f(edge, first point)` for every edge segment that may come within `a + b·t` of the
+    /// ray from `o` along unit `d`, `t` being the distance along it (each at most once). A
+    /// constant `a` suits an orthographic view; a perspective view widens with `b`.
+    pub fn segments_near(&self, o: Vec3, d: Vec3, a: f64, b: f64, mut f: impl FnMut(usize, usize)) {
+        if ![a, b, o.x, o.y, o.z, d.x, d.y, d.z].iter().all(|v| v.is_finite()) {
+            return;
+        }
+        let inv = Vec3::new(1.0 / d.x, 1.0 / d.y, 1.0 / d.z);
+        self.segments.visit(
+            |bx| {
+                let reach = a.max(0.0) + b.max(0.0) * farthest_along(o, d, bx);
+                ray_box(o, d, inv, bx, reach).is_some()
+            },
+            |i| {
+                if let Some((e, s)) = self.segment_of.get(i as usize) {
+                    f(*e as usize, *s as usize);
+                }
+            },
+        );
+    }
+}
+
+#[cfg(test)]
+#[path = "bvh_tests.rs"]
+mod tests;
diff --git a/crates/geom/src/bvh_tests.rs b/crates/geom/src/bvh_tests.rs
new file mode 100644
index 00000000..f551e963
--- /dev/null
+++ b/crates/geom/src/bvh_tests.rs
@@ -0,0 +1,130 @@
+use super::*;
+
+/// A small deterministic generator (the tests need no crate for randomness).
+struct Lcg(u64);
+
+impl Lcg {
+    fn next(&mut self) -> f64 {
+        self.0 = self.0.wrapping_mul(6364136223846793005).wrapping_add(1442695040888963407);
+        (self.0 >> 11) as f64 / (1u64 << 53) as f64
+    }
+    fn range(&mut self, lo: f64, hi: f64) -> f64 {
+        lo + (hi - lo) * self.next()
+    }
+    fn vec(&mut self, r: f64) -> Vec3 {
+        Vec3::new(self.range(-r, r), self.range(-r, r), self.range(-r, r))
+    }
+}
+
+/// A grid of `n` × `n` squares of two triangles on z = 0, and `extra` random triangles above.
+fn mesh(n: u32, extra: usize, rng: &mut Lcg) -> Mesh {
+    let mut positions: Vec<Vec3> = (0..=n).flat_map(|j| (0..=n).map(move |i| Vec3::new(f64::from(i), f64::from(j), 0.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)]);
+        }
+    }
+    for _ in 0..extra {
+        let c = rng.vec(f64::from(n)) + Vec3::new(0.0, 0.0, f64::from(n));
+        let k = positions.len() as u32;
+        positions.extend([c + rng.vec(1.5), c + rng.vec(1.5), c + rng.vec(1.5)]);
+        triangles.push([k, k + 1, k + 2]);
+    }
+    let edges =
+        (0..n).map(|i| (0..=n).map(|j| Vec3::new(f64::from(i), f64::from(j), 0.0)).collect()).chain([vec![Vec3::new(0.5, 0.5, 2.0)]]).collect();
+    Mesh { tri_face: vec![0; triangles.len()], positions, triangles, edges, ..Default::default() }
+}
+
+#[test]
+fn indexed_raycast_answers_as_the_scan_does() {
+    let mut rng = Lcg(7);
+    let m = mesh(12, 400, &mut rng);
+    let index = MeshIndex::new(&m);
+    let mut hits = 0;
+    for k in 0..4000 {
+        // Every few rays go straight down through grid vertices and edges, where two or more
+        // triangles tie.
+        let (o, d) = if k % 4 == 0 {
+            let p = Vec3::new(f64::from(k % 13), f64::from((k / 13) % 13) * 0.5, 50.0);
+            (p, Vec3::new(0.0, 0.0, -1.0))
+        } else {
+            let o = rng.vec(40.0) + Vec3::new(6.0, 6.0, 0.0);
+            let target = Vec3::new(rng.range(-1.0, 13.0), rng.range(-1.0, 13.0), rng.range(-2.0, 14.0));
+            let Some(d) = (target - o).normalized() else { continue };
+            (o, d)
+        };
+        let want = m.raycast(o, d);
+        assert_eq!(index.raycast(&m, o, d), want, "ray {k} from {o:?} along {d:?}");
+        hits += usize::from(want.is_some());
+    }
+    assert!(hits > 1000, "the rays should mostly hit ({hits})");
+}
+
+#[test]
+fn indexed_raycast_of_another_mesh_falls_back_to_the_scan() {
+    let mut rng = Lcg(3);
+    let (a, b) = (mesh(2, 0, &mut rng), mesh(5, 0, &mut rng));
+    let index = MeshIndex::new(&a);
+    let (o, d) = (Vec3::new(4.5, 4.2, 9.0), Vec3::new(0.0, 0.0, -1.0));
+    assert!(b.raycast(o, d).is_some());
+    assert_eq!(index.raycast(&b, o, d), b.raycast(o, d));
+}
+
+/// Distance from point `p` to the ray from `o` along unit `d` (from its start on), and how far along.
+fn ray_point(o: Vec3, d: Vec3, p: Vec3) -> (f64, f64) {
+    let t = (p - o).dot(d).max(0.0);
+    ((o + d * t).dist(p), t)
+}
+
+#[test]
+fn segments_near_the_ray_are_all_reported() {
+    let mut rng = Lcg(11);
+    let m = mesh(16, 0, &mut rng);
+    let index = MeshIndex::new(&m);
+    for k in 0..500 {
+        let o = rng.vec(30.0) + Vec3::new(8.0, 8.0, 30.0);
+        let target = Vec3::new(rng.range(0.0, 16.0), rng.range(0.0, 16.0), 0.0);
+        let Some(d) = (target - o).normalized() else { continue };
+        let (a, b) = if k % 2 == 0 { (0.3, 0.0) } else { (0.0, 0.01) };
+        let mut got = std::collections::HashSet::new();
+        index.segments_near(o, d, a, b, |e, s| {
+            assert!(got.insert((e, s)), "segment {e}/{s} reported twice");
+        });
+        for (ei, e) in m.edges.iter().enumerate() {
+            for (si, w) in e.windows(2).enumerate() {
+                // Sampled along the segment: any point within reach must have its segment reported.
+                for f in 0..=8 {
+                    let p = w[0] + (w[1] - w[0]) * (f64::from(f) / 8.0);
+                    let (dist, t) = ray_point(o, d, p);
+                    if dist <= a + b * t {
+                        assert!(got.contains(&(ei, si)), "ray {k}: segment {ei}/{si} is within reach but was not reported");
+                    }
+                }
+            }
+        }
+        // The one-point edge stands as its own segment.
+        let last = m.edges.len() - 1;
+        let (dist, t) = ray_point(o, d, Vec3::new(0.5, 0.5, 2.0));
+        if dist <= a + b * t {
+            assert!(got.contains(&(last, 0)));
+        }
+    }
+}
+
+#[test]
+fn hostile_boxes_never_panic() {
+    let nan = Vec3::new(f64::NAN, 0.0, 0.0);
+    let boxes = [Aabb3::EMPTY, Aabb3 { min: nan, max: nan }, Aabb3 { min: Vec3::ZERO, max: Vec3::new(1.0, 1.0, 1.0) }];
+    let bvh = Bvh::new(&boxes);
+    assert_eq!(bvh.nearest(Vec3::new(0.5, 0.5, 5.0), Vec3::new(0.0, 0.0, -1.0), |i| Some(f64::from(i))), Some((2.0, 2)));
+    assert_eq!(bvh.nearest(Vec3::ZERO, Vec3::new(f64::NAN, 0.0, 0.0), |_| Some(1.0)), None);
+    let empty = Bvh::new(&[]);
+    assert!(empty.nearest(Vec3::ZERO, Vec3::new(0.0, 0.0, 1.0), |_| Some(1.0)).is_none());
+    empty.visit(|_| true, |_| panic!("an empty hierarchy has no items"));
+    let m = Mesh { positions: vec![Vec3::ZERO], triangles: vec![[0, 5, 9]], ..Default::default() };
+    assert!(MeshIndex::new(&m).raycast(&m, Vec3::new(0.0, 0.0, 1.0), Vec3::new(0.0, 0.0, -1.0)).is_none());
+    MeshIndex::new(&m).segments_near(Vec3::ZERO, Vec3::new(0.0, 0.0, 1.0), f64::NAN, 0.0, |_, _| panic!("no segments"));
+}
diff --git a/crates/geom/src/lib.rs b/crates/geom/src/lib.rs
index 347d150d..2ad4f3d8 100644
--- a/crates/geom/src/lib.rs
+++ b/crates/geom/src/lib.rs
@@ -12,6 +12,7 @@ mod arcfit;
 pub mod blend;
 pub mod brep;
 pub mod bspline;
+mod bvh;
 pub mod compare;
 pub mod env;
 pub mod fill;
@@ -25,6 +26,7 @@ mod vec;
 
 pub use arcfit::MAX_FIT_PIECES;
 pub use blend::hermite_bezier;
+pub use bvh::{Bvh, MeshIndex};
 pub use mesh::{Aabb3, BoxedTriangle, InsideIndex, MassProperties, Mesh, MeshMeasure, find_close_pair};
 pub use plane::Plane;
 pub use profile::{Loop2, MAX_BEZIER_CTRL, MAX_BEZIER_DEGREE, MAX_DERIVATIVE_ORDER, Region2, Seg2, bezier, regions_apart};
diff --git a/crates/geom/src/mesh.rs b/crates/geom/src/mesh.rs
index 8a9045d4..92339a0a 100644
--- a/crates/geom/src/mesh.rs
+++ b/crates/geom/src/mesh.rs
@@ -559,13 +559,39 @@ impl Mesh {
     /// 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 frame = RayFrame::new(o, dir)?;
+        let mut best: Option<(f64, usize)> = None;
+        for (i, t) in self.triangles.iter().enumerate() {
+            let Some(p) = self.tri(t) else { continue };
+            if let Some(t) = frame.hit(p)
+                && best.is_none_or(|(bt, _)| t < bt)
+            {
+                best = Some((t, i));
+            }
+        }
+        best
+    }
+}
+
+/// A ray in the frame of [`Mesh::raycast`]'s watertight test: its dominant axis becomes z; x and
+/// y follow it cyclically, swapped to keep the frame right-handed when the ray runs towards −z.
+pub(crate) struct RayFrame {
+    o: Vec3,
+    k: [usize; 3],
+    s: [f64; 3],
+}
+
+fn comp(v: Vec3, k: usize) -> f64 {
+    match k {
+        0 => v.x,
+        1 => v.y,
+        _ => v.z,
+    }
+}
+
+impl RayFrame {
+    /// `None` for a direction that is zero or not finite.
+    pub(crate) fn new(o: Vec3, dir: Vec3) -> Option<RayFrame> {
         let a = [dir.x.abs(), dir.y.abs(), dir.z.abs()];
         let kz = if a[0] >= a[1] && a[0] >= a[2] {
             0
@@ -582,33 +608,31 @@ impl Mesh {
         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([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 det = u + v + w;
-            if det == 0.0 {
-                continue;
-            }
-            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));
-            }
+        Some(RayFrame { o, k: [kx, ky, kz], s: [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.
+    fn frame(&self, p: Vec3) -> (f64, f64, f64) {
+        let q = p - self.o;
+        let z = comp(q, self.k[2]);
+        (comp(q, self.k[0]) - self.s[0] * z, comp(q, self.k[1]) - self.s[1] * z, self.s[2] * z)
+    }
+
+    /// The distance along the ray to triangle `[pa, pb, pc]` (`None`: missed, or not in front).
+    pub(crate) fn hit(&self, [pa, pb, pc]: [Vec3; 3]) -> Option<f64> {
+        let ((ax, ay, az), (bx, by, bz), (cx, cy, cz)) = (self.frame(pa), self.frame(pb), self.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) {
+            return None;
         }
-        best
+        let det = u + v + w;
+        if det == 0.0 {
+            return None;
+        }
+        let t = (u * az + v * bz + w * cz) / det;
+        (t.is_finite() && t > 1e-9).then_some(t)
     }
 }
 
diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs
index 98f4515a..85200480 100644
--- a/crates/gpu/src/lib.rs
+++ b/crates/gpu/src/lib.rs
@@ -21,6 +21,7 @@ use wgpu::util::DeviceExt;
 
 pub mod capture;
 mod pbr;
+pub mod pick;
 #[cfg(not(target_arch = "wasm32"))]
 pub mod select;
 pub mod view;
diff --git a/crates/gpu/src/pick.rs b/crates/gpu/src/pick.rs
new file mode 100644
index 00000000..c62d2c7c
--- /dev/null
+++ b/crates/gpu/src/pick.rs
@@ -0,0 +1,503 @@
+//! The pick buffer: which model edges are visible at each pixel of the viewport, for picking on a
+//! large model without testing every edge. Surfaces are drawn into the depth buffer only; then
+//! every edge segment is drawn as a line carrying its id, depth tested against them (lifted
+//! toward the eye by the picker's slack, the depth by which an edge may lie behind a face and
+//! still count as visible) and written into an `R32Uint` id texture (0: no edge).
+//!
+//! The pass is drawn once per view (camera, viewport size, model), and only small tiles around
+//! the cursor are read back, so a pointer moving over a still view reads a few kilobytes and
+//! draws nothing. Read-backs arrive a frame or two later ([`PickResults`]); until a tile is in,
+//! the picker answers on the CPU. The ids are hints: the picker still runs its exact tests on
+//! the edges a tile names.
+
+use std::sync::atomic::{AtomicU8, Ordering};
+use std::sync::{Arc, Mutex};
+
+use wgpu::util::DeviceExt;
+
+/// The side of a read-back tile, in pixels (a multiple of 64, so rows need no padding).
+pub const TILE: u32 = 128;
+/// Most tiles kept from one view.
+const KEEP: usize = 6;
+/// Most read-backs in flight.
+const IN_FLIGHT: usize = 3;
+/// Bytes per surface vertex (position) and per edge vertex (position, id).
+const FACE_VERTEX: u64 = 12;
+const EDGE_VERTEX: u64 = 16;
+const UNIFORM_SIZE: u64 = 96;
+
+const SHADER: &str = r#"
+struct U {
+    vp: mat4x4<f32>,
+    eye: vec4<f32>,  // eye position, w = 1 for perspective
+    back: vec4<f32>, // toward the eye (orthographic), w = how far edges are lifted toward it
+};
+@group(0) @binding(0) var<uniform> u: U;
+
+@vertex
+fn vs_face(@location(0) p: vec3<f32>) -> @builtin(position) vec4<f32> {
+    return u.vp * vec4<f32>(p, 1.0);
+}
+
+@fragment
+fn fs_face() -> @location(0) u32 {
+    return 0u;
+}
+
+struct EOut {
+    @builtin(position) pos: vec4<f32>,
+    @location(0) @interpolate(flat) id: u32,
+};
+
+@vertex
+fn vs_edge(@location(0) p: vec3<f32>, @location(1) id: u32) -> EOut {
+    var toward = u.back.xyz;
+    if (u.eye.w > 0.5) {
+        let v = u.eye.xyz - p;
+        let l = length(v);
+        if (l > 0.0) {
+            toward = v / l;
+        }
+    }
+    var o: EOut;
+    o.pos = u.vp * vec4<f32>(p + toward * u.back.w, 1.0);
+    o.id = id;
+    return o;
+}
+
+@fragment
+fn fs_edge(i: EOut) -> @location(0) u32 {
+    return i.id;
+}
+"#;
+
+/// What the pass draws: surfaces (for depth) and edge segments with their ids (from 1).
+#[derive(Default, Clone)]
+pub struct PickScene {
+    faces: Vec<u8>,
+    edges: Vec<u8>,
+}
+
+impl PickScene {
+    pub fn triangle(&mut self, t: [[f32; 3]; 3]) {
+        for p in t {
+            for c in p {
+                self.faces.extend_from_slice(&c.to_le_bytes());
+            }
+        }
+    }
+
+    /// Segment `a`–`b` of the edge `id` (0 is no edge: not drawn).
+    pub fn segment(&mut self, a: [f32; 3], b: [f32; 3], id: u32) {
+        if id == 0 {
+            return;
+        }
+        for p in [a, b] {
+            for c in p {
+                self.edges.extend_from_slice(&c.to_le_bytes());
+            }
+            self.edges.extend_from_slice(&id.to_le_bytes());
+        }
+    }
+
+    pub fn is_empty(&self) -> bool {
+        self.edges.is_empty()
+    }
+}
+
+/// The view the pass is drawn for.
+#[derive(Clone, Copy, Debug)]
+pub struct PickView {
+    /// Changes whenever anything below or the scene changes.
+    pub key: u64,
+    /// Column-major view-projection (as [`crate::view::Frame::view_proj`]).
+    pub view_proj: [[f32; 4]; 4],
+    /// Eye position, and 1 for a perspective camera (0: orthographic).
+    pub eye: [f32; 4],
+    /// Direction toward the eye (orthographic).
+    pub back: [f32; 3],
+    /// How far edges are lifted toward the eye (model units).
+    pub lift: f32,
+    /// The viewport in physical pixels.
+    pub size: [u32; 2],
+}
+
+/// Edge ids of a block of pixels of one view.
+#[derive(Clone, Debug, PartialEq, Eq)]
+pub struct PickTile {
+    pub key: u64,
+    pub origin: [u32; 2],
+    pub size: [u32; 2],
+    /// Row-major, `size[0] * size[1]`.
+    pub ids: Vec<u32>,
+}
+
+impl PickTile {
+    /// Whether the tile is of view `key` and holds every pixel of `[x0, y0, x1, y1]` (inclusive)
+    /// that lies inside a viewport of `view` pixels.
+    pub fn covers(&self, key: u64, view: [u32; 2], [x0, y0, x1, y1]: [i64; 4]) -> bool {
+        let clamp = |v: i64, n: u32| v.clamp(0, i64::from(n.max(1)) - 1);
+        let (x0, x1, y0, y1) = (clamp(x0, view[0]), clamp(x1, view[0]), clamp(y0, view[1]), clamp(y1, view[1]));
+        let [ox, oy] = self.origin.map(i64::from);
+        let [w, h] = self.size.map(i64::from);
+        self.key == key && x0 >= ox && y0 >= oy && x1 < ox + w && y1 < oy + h
+    }
+
+    /// The ids within `[x0, y0, x1, y1]` (inclusive; outside the tile: none), each once.
+    pub fn ids_in(&self, [x0, y0, x1, y1]: [i64; 4]) -> Vec<u32> {
+        let [ox, oy] = self.origin.map(i64::from);
+        let [w, h] = self.size.map(i64::from);
+        let mut out = Vec::new();
+        for y in y0.max(oy)..=y1.min(oy + h - 1) {
+            for x in x0.max(ox)..=x1.min(ox + w - 1) {
+                if let Some(&id) = self.ids.get(((y - oy) * w + (x - ox)) as usize)
+                    && id != 0
+                {
+                    out.push(id);
+                }
+            }
+        }
+        out.sort_unstable();
+        out.dedup();
+        out
+    }
+}
+
+/// The tiles read back so far (of the latest view only), newest last.
+#[derive(Default, Debug)]
+pub struct PickResults {
+    pub tiles: Vec<PickTile>,
+}
+
+impl PickResults {
+    /// The newest tile of view `key` covering `window`.
+    pub fn tile(&self, key: u64, view: [u32; 2], window: [i64; 4]) -> Option<&PickTile> {
+        self.tiles.iter().rev().find(|t| t.covers(key, view, window))
+    }
+}
+
+pub type PickOut = Arc<Mutex<PickResults>>;
+
+/// The tile (origin and size) around pixel `centre` of a viewport of `view` pixels.
+pub fn tile_around(centre: [u32; 2], view: [u32; 2]) -> ([u32; 2], [u32; 2]) {
+    let size = [TILE.min(view[0]), TILE.min(view[1])];
+    let origin = [0, 1].map(|k| centre[k].saturating_sub(size[k] / 2).min(view[k].saturating_sub(size[k])));
+    (origin, size)
+}
+
+const NEW: u8 = 0;
+const MAPPING: u8 = 1;
+const READY: u8 = 2;
+const FAILED: u8 = 3;
+
+/// A tile being copied out of the id texture.
+struct Pending {
+    key: u64,
+    origin: [u32; 2],
+    size: [u32; 2],
+    bytes_per_row: u32,
+    buffer: wgpu::Buffer,
+    /// The frame whose commands copy it (it is mapped once they are submitted).
+    frame: u64,
+    state: Arc<AtomicU8>,
+}
+
+struct Target {
+    size: [u32; 2],
+    ids: wgpu::Texture,
+    ids_view: wgpu::TextureView,
+    depth_view: wgpu::TextureView,
+}
+
+struct Uploaded {
+    key: u64,
+    faces: Option<(wgpu::Buffer, u32)>,
+    edges: Option<(wgpu::Buffer, u32)>,
+}
+
+/// The pass's pipelines, its textures and the read-backs in flight, on one device.
+pub struct PickBuffer {
+    face: wgpu::RenderPipeline,
+    edge: wgpu::RenderPipeline,
+    uniform: wgpu::Buffer,
+    bind: wgpu::BindGroup,
+    scene: Option<Uploaded>,
+    target: Option<Target>,
+    /// The view drawn into the textures.
+    drawn: Option<u64>,
+    pending: Vec<Pending>,
+    frame: u64,
+}
+
+const DEPTH: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float;
+const IDS: wgpu::TextureFormat = wgpu::TextureFormat::R32Uint;
+
+impl PickBuffer {
+    pub fn new(device: &wgpu::Device) -> PickBuffer {
+        let module =
+            device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("sc_pick"), source: wgpu::ShaderSource::Wgsl(SHADER.into()) });
+        let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("sc_pick"),
+            entries: &[wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::VERTEX,
+                ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
+                count: None,
+            }],
+        });
+        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("sc_pick"),
+            bind_group_layouts: &[Some(&bgl)],
+            immediate_size: 0,
+        });
+        let uniform = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("sc_pick_u"),
+            size: UNIFORM_SIZE,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+        let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("sc_pick"),
+            layout: &bgl,
+            entries: &[wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() }],
+        });
+        const FA: [wgpu::VertexAttribute; 1] = wgpu::vertex_attr_array![0 => Float32x3];
+        const EA: [wgpu::VertexAttribute; 2] = wgpu::vertex_attr_array![0 => Float32x3, 1 => Uint32];
+        let pipeline = |label: &str,
+                        vs: &str,
+                        fs: &str,
+                        stride: u64,
+                        attributes: &'static [wgpu::VertexAttribute],
+                        topology: wgpu::PrimitiveTopology,
+                        write: bool,
+                        mask: wgpu::ColorWrites| {
+            device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+                label: Some(label),
+                layout: Some(&layout),
+                vertex: wgpu::VertexState {
+                    module: &module,
+                    entry_point: Some(vs),
+                    buffers: &[Some(wgpu::VertexBufferLayout { array_stride: stride, step_mode: wgpu::VertexStepMode::Vertex, attributes })],
+                    compilation_options: Default::default(),
+                },
+                primitive: wgpu::PrimitiveState { topology, ..Default::default() },
+                depth_stencil: Some(wgpu::DepthStencilState {
+                    format: DEPTH,
+                    depth_write_enabled: Some(write),
+                    depth_compare: Some(wgpu::CompareFunction::LessEqual),
+                    stencil: Default::default(),
+                    bias: Default::default(),
+                }),
+                multisample: wgpu::MultisampleState::default(),
+                fragment: Some(wgpu::FragmentState {
+                    module: &module,
+                    entry_point: Some(fs),
+                    targets: &[Some(wgpu::ColorTargetState { format: IDS, blend: None, write_mask: mask })],
+                    compilation_options: Default::default(),
+                }),
+                multiview_mask: None,
+                cache: None,
+            })
+        };
+        let face =
+            pipeline("sc_pick_face", "vs_face", "fs_face", FACE_VERTEX, &FA, wgpu::PrimitiveTopology::TriangleList, true, wgpu::ColorWrites::empty());
+        let edge = pipeline("sc_pick_edge", "vs_edge", "fs_edge", EDGE_VERTEX, &EA, wgpu::PrimitiveTopology::LineList, false, wgpu::ColorWrites::ALL);
+        PickBuffer { face, edge, uniform, bind, scene: None, target: None, drawn: None, pending: Vec::new(), frame: 0 }
+    }
+
+    /// Before the frame is submitted: collect the tiles read back, then draw `view` of `scene`
+    /// (`scene_key` names it; it is uploaded when that changes) unless it is drawn already, and
+    /// start reading back the tile around pixel `centre` unless it is in or on its way. The tiles
+    /// land in `out` a frame or two later; keep frames coming until then.
+    #[allow(clippy::too_many_arguments)]
+    pub fn prepare(
+        &mut self,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+        encoder: &mut wgpu::CommandEncoder,
+        scene_key: u64,
+        scene: &PickScene,
+        view: &PickView,
+        centre: Option<[u32; 2]>,
+        out: &PickOut,
+    ) {
+        self.frame += 1;
+        self.collect(device, out);
+        let [w, h] = view.size;
+        let max = device.limits().max_texture_dimension_2d;
+        if w == 0 || h == 0 || w > max || h > max {
+            return;
+        }
+        if self.scene.as_ref().is_none_or(|s| s.key != scene_key) {
+            let limit = device.limits().max_buffer_size;
+            let upload = |label: &str, bytes: &[u8], stride: u64| {
+                let n = bytes.len() as u64 / stride;
+                (n > 0 && (bytes.len() as u64) <= limit && n <= u64::from(u32::MAX)).then(|| {
+                    (
+                        device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+                            label: Some(label),
+                            contents: bytes,
+                            usage: wgpu::BufferUsages::VERTEX,
+                        }),
+                        n as u32,
+                    )
+                })
+            };
+            self.scene = Some(Uploaded {
+                key: scene_key,
+                faces: upload("sc_pick_faces", &scene.faces, FACE_VERTEX),
+                edges: upload("sc_pick_edges", &scene.edges, EDGE_VERTEX),
+            });
+            self.drawn = None;
+        }
+        if self.target.as_ref().is_none_or(|t| t.size != view.size) {
+            let texture = |label: &str, format, usage| {
+                device.create_texture(&wgpu::TextureDescriptor {
+                    label: Some(label),
+                    size: wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
+                    mip_level_count: 1,
+                    sample_count: 1,
+                    dimension: wgpu::TextureDimension::D2,
+                    format,
+                    usage,
+                    view_formats: &[],
+                })
+            };
+            let ids = texture("sc_pick_ids", IDS, wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC);
+            let depth = texture("sc_pick_depth", DEPTH, wgpu::TextureUsages::RENDER_ATTACHMENT);
+            let (ids_view, depth_view) = (ids.create_view(&Default::default()), depth.create_view(&Default::default()));
+            self.target = Some(Target { size: view.size, ids, ids_view, depth_view });
+            self.drawn = None;
+        }
+        let (Some(target), Some(scene)) = (&self.target, &self.scene) else { return };
+        if self.drawn != Some(view.key) {
+            queue.write_buffer(&self.uniform, 0, &uniform_bytes(view));
+            let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                label: Some("sc_pick"),
+                color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                    view: &target.ids_view,
+                    depth_slice: None,
+                    resolve_target: None,
+                    ops: wgpu::Operations { load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), store: wgpu::StoreOp::Store },
+                })],
+                depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
+                    view: &target.depth_view,
+                    depth_ops: Some(wgpu::Operations { load: wgpu::LoadOp::Clear(1.0), store: wgpu::StoreOp::Discard }),
+                    stencil_ops: None,
+                }),
+                timestamp_writes: None,
+                occlusion_query_set: None,
+                multiview_mask: None,
+            });
+            pass.set_bind_group(0, &self.bind, &[]);
+            for (pipeline, part) in [(&self.face, &scene.faces), (&self.edge, &scene.edges)] {
+                if let Some((buffer, n)) = part {
+                    pass.set_pipeline(pipeline);
+                    pass.set_vertex_buffer(0, buffer.slice(..));
+                    pass.draw(0..*n, 0..1);
+                }
+            }
+            drop(pass);
+            self.drawn = Some(view.key);
+            // Tiles of other views no longer answer.
+            self.pending.retain(|p| p.key == view.key);
+            if let Ok(mut o) = out.lock() {
+                o.tiles.retain(|t| t.key == view.key);
+            }
+        }
+        let Some(centre) = centre else { return };
+        let (origin, size) = tile_around(centre, view.size);
+        let known = out.lock().is_ok_and(|o| o.tiles.iter().any(|t| t.key == view.key && t.origin == origin && t.size == size))
+            || self.pending.iter().any(|p| p.key == view.key && p.origin == origin && p.size == size);
+        if known || self.pending.len() >= IN_FLIGHT {
+            return;
+        }
+        let bytes_per_row = (size[0] * 4).div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
+        let buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("sc_pick_tile"),
+            size: u64::from(bytes_per_row) * u64::from(size[1]),
+            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+        encoder.copy_texture_to_buffer(
+            wgpu::TexelCopyTextureInfo {
+                texture: &target.ids,
+                mip_level: 0,
+                origin: wgpu::Origin3d { x: origin[0], y: origin[1], z: 0 },
+                aspect: wgpu::TextureAspect::All,
+            },
+            wgpu::TexelCopyBufferInfo {
+                buffer: &buffer,
+                layout: wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(bytes_per_row), rows_per_image: Some(size[1]) },
+            },
+            wgpu::Extent3d { width: size[0], height: size[1], depth_or_array_layers: 1 },
+        );
+        self.pending.push(Pending { key: view.key, origin, size, bytes_per_row, buffer, frame: self.frame, state: Arc::new(AtomicU8::new(NEW)) });
+    }
+
+    /// Whether read-backs are still on their way (keep frames coming).
+    pub fn waiting(&self) -> bool {
+        !self.pending.is_empty()
+    }
+
+    /// Maps the copies earlier frames submitted and moves the mapped ones into `out`.
+    fn collect(&mut self, device: &wgpu::Device, out: &PickOut) {
+        let frame = self.frame;
+        for p in &self.pending {
+            if p.frame < frame && p.state.load(Ordering::Acquire) == NEW {
+                p.state.store(MAPPING, Ordering::Release);
+                let state = p.state.clone();
+                p.buffer.slice(..).map_async(wgpu::MapMode::Read, move |r| state.store(if r.is_ok() { READY } else { FAILED }, Ordering::Release));
+            }
+        }
+        if self.pending.iter().any(|p| p.state.load(Ordering::Acquire) == MAPPING) {
+            let _ = device.poll(wgpu::PollType::Poll);
+        }
+        let mut done = Vec::new();
+        self.pending.retain(|p| match p.state.load(Ordering::Acquire) {
+            READY => {
+                if let Ok(bytes) = p.buffer.slice(..).get_mapped_range() {
+                    let [w, h] = p.size;
+                    let mut ids = Vec::with_capacity(w as usize * h as usize);
+                    for row in bytes.chunks(p.bytes_per_row as usize).take(h as usize) {
+                        ids.extend(row.as_chunks::<4>().0.iter().take(w as usize).map(|c| u32::from_le_bytes(*c)));
+                    }
+                    done.push(PickTile { key: p.key, origin: p.origin, size: p.size, ids });
+                }
+                p.buffer.unmap();
+                false
+            }
+            FAILED => false,
+            _ => true,
+        });
+        if !done.is_empty()
+            && let Ok(mut o) = out.lock()
+        {
+            for t in done {
+                if self.drawn == Some(t.key) && t.ids.len() == t.size[0] as usize * t.size[1] as usize {
+                    o.tiles.retain(|x| x.key == t.key);
+                    o.tiles.push(t);
+                }
+            }
+            let excess = o.tiles.len().saturating_sub(KEEP);
+            o.tiles.drain(..excess);
+        }
+    }
+}
+
+fn uniform_bytes(v: &PickView) -> Vec<u8> {
+    let mut b = Vec::with_capacity(UNIFORM_SIZE as usize);
+    for col in v.view_proj {
+        for x in col {
+            b.extend_from_slice(&x.to_le_bytes());
+        }
+    }
+    for x in v.eye.into_iter().chain(v.back).chain([v.lift]) {
+        b.extend_from_slice(&x.to_le_bytes());
+    }
+    b
+}
+
+#[cfg(test)]
+#[path = "pick_tests.rs"]
+mod tests;
diff --git a/crates/gpu/src/pick_tests.rs b/crates/gpu/src/pick_tests.rs
new file mode 100644
index 00000000..5ccb5e2b
--- /dev/null
+++ b/crates/gpu/src/pick_tests.rs
@@ -0,0 +1,79 @@
+use super::*;
+
+/// Looking down −z at x, y in −10..10 (orthographic): depth 0 at z = 50, 1 at z = −50.
+const VIEW: [[f32; 4]; 4] = [[0.1, 0.0, 0.0, 0.0], [0.0, 0.1, 0.0, 0.0], [0.0, 0.0, -0.01, 0.0], [0.0, 0.0, 0.5, 1.0]];
+const SIZE: u32 = 64;
+
+/// The pixel of model point (x, y).
+fn px(x: f32, y: f32) -> [i64; 2] {
+    [((x / 10.0 + 1.0) * 0.5 * SIZE as f32) as i64, ((1.0 - y / 10.0) * 0.5 * SIZE as f32) as i64]
+}
+
+/// A square face at z = 0 (x, y in −5..5), and three edges: 1 on it, 2 well behind it, 3 behind
+/// it by less than the lift.
+fn scene() -> PickScene {
+    let mut s = PickScene::default();
+    let (a, b, c, d) = ([-5.0, -5.0, 0.0], [5.0, -5.0, 0.0], [5.0, 5.0, 0.0], [-5.0, 5.0, 0.0]);
+    s.triangle([a, b, c]);
+    s.triangle([a, c, d]);
+    s.segment([-8.0, 0.0, 0.0], [8.0, 0.0, 0.0], 1);
+    s.segment([2.0, -8.0, -3.0], [2.0, 8.0, -3.0], 2);
+    s.segment([-2.0, -8.0, -0.2], [-2.0, 8.0, -0.2], 3);
+    s
+}
+
+#[test]
+fn the_pick_buffer_names_the_edges_seen_at_each_pixel() {
+    let (_, device, queue) = match crate::open_headless() {
+        Ok(d) => d,
+        Err(e) => {
+            assert!(std::env::var_os("SOLIDCRAFT_REQUIRE_GPU").is_none(), "a GPU is required here: {e}");
+            eprintln!("skipping: {e}");
+            return;
+        }
+    };
+    let mut pb = PickBuffer::new(&device);
+    let out = PickOut::default();
+    let view = PickView { key: 7, view_proj: VIEW, eye: [0.0, 0.0, 100.0, 0.0], back: [0.0, 0.0, 1.0], lift: 0.5, size: [SIZE, SIZE] };
+    let scene = scene();
+    let all = [0, 0, SIZE as i64 - 1, SIZE as i64 - 1];
+    for _ in 0..8 {
+        let mut encoder = device.create_command_encoder(&Default::default());
+        pb.prepare(&device, &queue, &mut encoder, 1, &scene, &view, Some([SIZE / 2, SIZE / 2]), &out);
+        queue.submit([encoder.finish()]);
+        let _ = device.poll(wgpu::PollType::wait_indefinitely());
+        if out.lock().unwrap().tile(7, [SIZE, SIZE], all).is_some() {
+            break;
+        }
+    }
+    let results = out.lock().unwrap();
+    let tile = results.tile(7, [SIZE, SIZE], all).expect("the tile is read back within a few frames");
+    let at = |x: f32, y: f32| {
+        let [px, py] = px(x, y);
+        tile.ids_in([px - 1, py - 1, px + 1, py + 1])
+    };
+    assert_eq!(at(0.0, 0.0), vec![1], "the edge on the face");
+    assert_eq!(at(2.0, 3.0), Vec::<u32>::new(), "the edge behind the face is hidden");
+    assert_eq!(at(2.0, 7.5), vec![2], "and seen where the face does not cover it");
+    assert_eq!(at(-2.0, 3.0), vec![3], "an edge behind the face by less than the lift counts as seen");
+    assert_eq!(at(3.5, 3.5), Vec::<u32>::new(), "the face itself names no edge");
+    assert!(results.tile(8, [SIZE, SIZE], all).is_none(), "another view's tile does not answer");
+    drop(results);
+    assert!(!pb.waiting());
+}
+
+#[test]
+fn tiles_cover_the_window_around_the_cursor_inside_the_view() {
+    assert_eq!(tile_around([10, 10], [1000, 800]), ([0, 0], [TILE, TILE]));
+    assert_eq!(tile_around([500, 400], [1000, 800]), ([500 - TILE / 2, 400 - TILE / 2], [TILE, TILE]));
+    assert_eq!(tile_around([999, 799], [1000, 800]), ([1000 - TILE, 800 - TILE], [TILE, TILE]));
+    assert_eq!(tile_around([30, 20], [50, 40]), ([0, 0], [50, 40]));
+    let t = PickTile { key: 1, origin: [100, 100], size: [TILE, TILE], ids: vec![0; (TILE * TILE) as usize] };
+    assert!(t.covers(1, [1000, 800], [110, 110, 120, 120]));
+    assert!(!t.covers(1, [1000, 800], [90, 110, 120, 120]));
+    assert!(!t.covers(2, [1000, 800], [110, 110, 120, 120]));
+    // A window past the view's border needs only the pixels inside it.
+    let corner = PickTile { key: 1, origin: [0, 0], size: [TILE, TILE], ids: vec![5; (TILE * TILE) as usize] };
+    assert!(corner.covers(1, [1000, 800], [-8, -8, 8, 8]));
+    assert_eq!(corner.ids_in([-8, -8, 8, 8]), vec![5]);
+}
diff --git a/crates/ui-egui/src/gpu.rs b/crates/ui-egui/src/gpu.rs
index 83e2e3f8..9c072c7e 100644
--- a/crates/ui-egui/src/gpu.rs
+++ b/crates/ui-egui/src/gpu.rs
@@ -136,7 +136,7 @@ pub(crate) fn offscreen_backends() -> wgpu::Backends {
 /// A device of its own for an offscreen render or check (`None` without a GPU adapter), on
 /// [`offscreen_backends`].
 #[cfg(all(test, not(target_arch = "wasm32")))]
-fn offscreen_device() -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
+pub(crate) fn offscreen_device() -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
     let mut desc = wgpu::InstanceDescriptor::new_without_display_handle_from_env();
     desc.backends = offscreen_backends();
     let instance = wgpu::Instance::new(desc);
diff --git a/crates/ui-egui/src/lib.rs b/crates/ui-egui/src/lib.rs
index a5e76ed5..906ef0d9 100644
--- a/crates/ui-egui/src/lib.rs
+++ b/crates/ui-egui/src/lib.rs
@@ -112,6 +112,8 @@ mod fillet_spline_handle_tests;
 pub(crate) mod navigation;
 pub mod palette;
 pub mod params_dialog;
+pub(crate) mod pick_buffer;
+pub(crate) mod pick_index;
 pub mod prefs;
 #[cfg(test)]
 mod preselect_tests;
diff --git a/crates/ui-egui/src/pick_buffer.rs b/crates/ui-egui/src/pick_buffer.rs
new file mode 100644
index 00000000..373018f2
--- /dev/null
+++ b/crates/ui-egui/src/pick_buffer.rs
@@ -0,0 +1,259 @@
+//! The GPU pick buffer in the viewport ([`solidcraft_gpu::pick`]): which model edges are seen at
+//! each pixel, so a pick over a large model tests only the edges near the cursor that are not
+//! hidden, instead of every edge the cursor's ray passes near.
+//!
+//! The pass is drawn for a view once it holds still (the same view two frames running), and a
+//! tile around the cursor is read back; the tile arrives a frame or two later, and until it
+//! does (and while the view moves) the picker answers from the CPU indexes
+//! ([`crate::pick_index`]). The tile only narrows the candidates: the picker's exact tests
+//! still decide. Not used for Select Other (every target along the ray), in wireframe (nothing
+//! hides an edge there), without a GPU, or for a model past [`MAX_TRIANGLES`].
+
+use std::collections::HashMap;
+use std::hash::{Hash, Hasher};
+use std::sync::{Arc, Mutex};
+
+use egui::Pos2;
+use solidcraft_engine::doc::ModelState;
+use solidcraft_gpu::pick::{PickBuffer, PickOut, PickScene, PickView};
+
+use crate::SolveApp;
+use crate::viewport::Proj;
+
+/// Most surface triangles the pass draws (beyond them the CPU indexes answer alone).
+pub const MAX_TRIANGLES: usize = 6_000_000;
+/// How near the cursor an edge's pixels are looked for (points): vertices pick within 7, edges 6.
+const WINDOW: f32 = 8.0;
+/// Frames kept coming after a tile is asked for, so its read-back lands.
+const FRAMES: u8 = 4;
+
+/// The pass's scene for one model state: the ids it draws stand for (body, edge).
+struct Scene {
+    state: Arc<ModelState>,
+    key: u64,
+    scene: Arc<PickScene>,
+    /// Per id − 1: (body index in the state, edge index).
+    edges: Vec<(u32, u32)>,
+    /// One-point edges (no line to draw): always candidates.
+    points: Vec<(u32, u32)>,
+}
+
+#[derive(Default)]
+pub(crate) struct State {
+    scene: Mutex<Option<Arc<Scene>>>,
+    out: PickOut,
+    /// A pick found no tile for its cursor: the view it was in and the pixel.
+    want: Mutex<Option<(u64, [u32; 2])>>,
+    /// The view of the last frame (the pass is drawn for a view that held still).
+    last_view: u64,
+    /// Frames still to request so read-backs land.
+    frames: u8,
+}
+
+/// The edges a tile names near the cursor, per body index of the world state.
+pub(crate) struct NearEdges {
+    by_body: HashMap<usize, Vec<usize>>,
+}
+
+impl NearEdges {
+    /// Every segment of the named edges of body `index` of `mesh`, as (edge, first point), in
+    /// edge order.
+    pub(crate) fn segments(&self, index: usize, mesh: &solidcraft_engine::geom::Mesh) -> Vec<(usize, usize)> {
+        let mut v = Vec::new();
+        for &ei in self.by_body.get(&index).map_or(&[][..], |e| &e[..]) {
+            let n = mesh.edges.get(ei).map_or(0, |e| e.len().max(2) - 1);
+            v.extend((0..n).map(|si| (ei, si)));
+        }
+        v
+    }
+}
+
+/// What the pass's scene depends on: the world state, and which bodies show.
+fn model_key(app: &SolveApp, st: &Arc<ModelState>) -> u64 {
+    let mut h = std::collections::hash_map::DefaultHasher::new();
+    (Arc::as_ptr(st) as usize).hash(&mut h);
+    app.session.revision.hash(&mut h);
+    app.session.visibility.hidden_bodies.hash(&mut h);
+    app.session.visibility.hidden_occurrences.hash(&mut h);
+    h.finish()
+}
+
+/// The view's pixels: the viewport's size in physical pixels.
+fn view_size(app: &SolveApp, proj: &Proj) -> [u32; 2] {
+    let ppp = app.viewport.pixels_per_point();
+    [(proj.rect.width() * ppp).round().max(0.0) as u32, (proj.rect.height() * ppp).round().max(0.0) as u32]
+}
+
+/// How far edges are lifted toward the eye: the picker's slack for an edge behind a face.
+fn lift(app: &SolveApp) -> f64 {
+    crate::viewport::scene_radius(app) * 0.02
+}
+
+/// Everything the drawn ids depend on.
+fn view_key(app: &SolveApp, proj: &Proj, model: u64) -> u64 {
+    let mut h = std::collections::hash_map::DefaultHasher::new();
+    model.hash(&mut h);
+    for col in proj.vp.to_f32() {
+        for x in col {
+            x.to_bits().hash(&mut h);
+        }
+    }
+    view_size(app, proj).hash(&mut h);
+    lift(app).to_bits().hash(&mut h);
+    h.finish()
+}
+
+fn usable(app: &SolveApp) -> bool {
+    app.viewport.gpu.is_some() && app.ui.visual_style != 2
+}
+
+/// The edges seen within reach of `pos`, when a tile of this view is in; otherwise `None`, and
+/// a tile around `pos` is asked for.
+pub(crate) fn edges_near(app: &SolveApp, proj: &Proj, pos: Pos2, st: &Arc<ModelState>) -> Option<NearEdges> {
+    if !usable(app) {
+        return None;
+    }
+    let model = model_key(app, st);
+    let scene = app.viewport.pick_buffer.scene.lock().ok()?.clone().filter(|s| s.key == model && Arc::ptr_eq(&s.state, st));
+    let key = view_key(app, proj, model);
+    let size = view_size(app, proj);
+    let ppp = app.viewport.pixels_per_point();
+    let (x, y) = ((pos.x - proj.rect.left()) * ppp, (pos.y - proj.rect.top()) * ppp);
+    if !(x >= 0.0 && y >= 0.0 && x < size[0] as f32 && y < size[1] as f32) {
+        return None;
+    }
+    let (x, y) = (x as i64, y as i64);
+    let r = (WINDOW * ppp).ceil() as i64;
+    let window = [x - r, y - r, x + r, y + r];
+    let ids = scene.as_ref().and_then(|_| app.viewport.pick_buffer.out.lock().ok()?.tile(key, size, window).map(|t| t.ids_in(window)));
+    let (Some(scene), Some(ids)) = (scene, ids) else {
+        if let Ok(mut w) = app.viewport.pick_buffer.want.lock() {
+            *w = Some((key, [x as u32, y as u32]));
+        }
+        return None;
+    };
+    let mut by_body: HashMap<usize, Vec<usize>> = HashMap::new();
+    for (b, e) in ids.iter().filter_map(|id| scene.edges.get((*id as usize).checked_sub(1)?)).chain(&scene.points) {
+        by_body.entry(*b as usize).or_default().push(*e as usize);
+    }
+    for v in by_body.values_mut() {
+        v.sort_unstable();
+        v.dedup();
+    }
+    Some(NearEdges { by_body })
+}
+
+/// The pass's scene for the world state: visible bodies' triangles, and their pickable edges.
+/// `None` past [`MAX_TRIANGLES`] or without edges.
+fn scene(app: &SolveApp, st: &Arc<ModelState>, key: u64) -> Option<Arc<Scene>> {
+    let f = |p: solidcraft_engine::geom::Vec3| [p.x as f32, p.y as f32, p.z as f32];
+    let mut scene = PickScene::default();
+    let (mut edges, mut points, mut triangles) = (Vec::new(), Vec::new(), 0usize);
+    for (bi, b) in st.bodies.iter().enumerate() {
+        if !app.session.body_visible(&b.name) {
+            continue;
+        }
+        let m = b.mesh();
+        triangles += m.triangles.len();
+        if triangles > MAX_TRIANGLES {
+            return None;
+        }
+        for t in &m.triangles {
+            if let Some([a, b, c]) = m.tri(t) {
+                scene.triangle([f(a), f(b), f(c)]);
+            }
+        }
+        for (ei, e) in m.edges.iter().enumerate() {
+            if !m.edge_pickable(ei) || edges.len() >= u32::MAX as usize - 1 {
+                continue;
+            }
+            if e.len() == 1 {
+                points.push((bi as u32, ei as u32));
+                continue;
+            }
+            edges.push((bi as u32, ei as u32));
+            let id = edges.len() as u32;
+            for w in e.windows(2) {
+                scene.segment(f(w[0]), f(w[1]), id);
+            }
+        }
+    }
+    (!scene.is_empty()).then(|| Arc::new(Scene { state: st.clone(), key, scene: Arc::new(scene), edges, points }))
+}
+
+/// In the frame's GPU paint: draw the pass for a view that held still when a pick asked for a
+/// tile, and keep frames coming until the read-back lands.
+pub(crate) fn paint(app: &mut SolveApp, ctx: &egui::Context, painter: &egui::Painter, proj: &Proj) {
+    if !usable(app) {
+        return;
+    }
+    let st = app.session.world_state();
+    let model = model_key(app, &st);
+    let key = view_key(app, proj, model);
+    let state = &mut app.viewport.pick_buffer;
+    let still = std::mem::replace(&mut state.last_view, key) == key;
+    let want = state.want.lock().ok().and_then(|mut w| w.take()).filter(|(k, _)| *k == key && still);
+    if want.is_none() && state.frames == 0 {
+        return;
+    }
+    state.frames = if want.is_some() { FRAMES } else { state.frames.saturating_sub(1) };
+    let cached = state.scene.lock().ok().and_then(|s| s.clone()).filter(|s| s.key == model && Arc::ptr_eq(&s.state, &st));
+    let scene = match cached {
+        Some(s) => s,
+        None => {
+            let Some(s) = scene(app, &st, model) else { return };
+            if let Ok(mut slot) = app.viewport.pick_buffer.scene.lock() {
+                *slot = Some(s.clone());
+            }
+            s
+        }
+    };
+    let e = proj.cam.eye().to_f32();
+    let view = PickView {
+        key,
+        view_proj: proj.vp.to_f32(),
+        eye: [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }],
+        back: proj.cam.back().to_f32(),
+        lift: lift(app) as f32,
+        size: view_size(app, proj),
+    };
+    let cb =
+        PickCallback { scene_key: scene.key, scene: scene.scene.clone(), view, centre: want.map(|w| w.1), out: app.viewport.pick_buffer.out.clone() };
+    painter.add(egui_wgpu::Callback::new_paint_callback(proj.rect, cb));
+    ctx.request_repaint();
+}
+
+/// The pass as an egui paint callback: all its work is in `prepare` (it draws into its own
+/// textures, not the frame).
+struct PickCallback {
+    scene_key: u64,
+    scene: Arc<PickScene>,
+    view: PickView,
+    centre: Option<[u32; 2]>,
+    out: PickOut,
+}
+
+impl egui_wgpu::CallbackTrait for PickCallback {
+    fn prepare(
+        &self,
+        device: &egui_wgpu::wgpu::Device,
+        queue: &egui_wgpu::wgpu::Queue,
+        _screen: &egui_wgpu::ScreenDescriptor,
+        encoder: &mut egui_wgpu::wgpu::CommandEncoder,
+        resources: &mut egui_wgpu::CallbackResources,
+    ) -> Vec<egui_wgpu::wgpu::CommandBuffer> {
+        if resources.get::<PickBuffer>().is_none() {
+            resources.insert(PickBuffer::new(device));
+        }
+        if let Some(pb) = resources.get_mut::<PickBuffer>() {
+            pb.prepare(device, queue, encoder, self.scene_key, &self.scene, &self.view, self.centre, &self.out);
+        }
+        Vec::new()
+    }
+
+    fn paint(&self, _info: egui::PaintCallbackInfo, _pass: &mut egui_wgpu::wgpu::RenderPass<'static>, _resources: &egui_wgpu::CallbackResources) {}
+}
+
+#[cfg(test)]
+#[path = "pick_buffer_tests.rs"]
+mod tests;
diff --git a/crates/ui-egui/src/pick_buffer_tests.rs b/crates/ui-egui/src/pick_buffer_tests.rs
new file mode 100644
index 00000000..393c9498
--- /dev/null
+++ b/crates/ui-egui/src/pick_buffer_tests.rs
@@ -0,0 +1,141 @@
+use egui::{Rect, pos2};
+use serde_json::json;
+use solidcraft_engine::geom::Vec3;
+use solidcraft_gpu::pick::{PickTile, TILE, tile_around};
+
+use super::*;
+use crate::viewport::{Hit, pick, projection};
+
+fn app() -> SolveApp {
+    let mut s = solidcraft_engine::Session::default();
+    s.execute("solid.box", &json!({"length": 40, "width": 30, "height": 20})).unwrap();
+    s.execute("solid.fillet", &json!({"edges": [[40, 15, 20]], "radius": 4})).unwrap();
+    s.execute("solid.cylinder", &json!({"radius": 8, "height": 35, "base": [10, 10, 0], "operation": "new_body"})).unwrap();
+    let mut app = SolveApp::new(s, crate::Services::default());
+    app.ui.show_origin = false;
+    app.cam.target = Vec3::new(20.0, 15.0, 12.0);
+    app.cam.distance = 110.0;
+    // A GPU viewport (no device is needed to hand the picker a tile).
+    app.viewport.gpu = Some(crate::gpu::GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess: false });
+    app
+}
+
+const RECT: Rect = Rect { min: pos2(0.0, 0.0), max: pos2(640.0, 480.0) };
+
+/// Installs the pass's scene, as `paint` makes it, and returns the view's key.
+fn install_scene(app: &SolveApp, proj: &Proj) -> u64 {
+    let st = app.session.world_state();
+    let model = model_key(app, &st);
+    let s = scene(app, &st, model).unwrap();
+    *app.viewport.pick_buffer.scene.lock().unwrap() = Some(s);
+    view_key(app, proj, model)
+}
+
+/// A tile only narrows the edges: the edge it names near the cursor is picked as the CPU picks
+/// it, an edge it does not name (hidden there) is not, and without a tile the CPU indexes answer
+/// and a tile is asked for.
+#[test]
+fn a_tile_names_the_edges_the_picker_tests() {
+    let app = app();
+    let proj = projection(&app, RECT);
+    // A point on the box's top front edge (y = 0, z = 20).
+    let at = proj.to_screen(Vec3::new(20.0, 0.0, 20.0)).unwrap();
+    let cpu = pick(&app, &proj, at);
+    let Some(Hit::Edge { index, body, .. }) = cpu.iter().find(|h| matches!(h, Hit::Edge { .. })).cloned() else { panic!("{cpu:?}") };
+    assert!(app.viewport.pick_buffer.want.lock().unwrap().is_some(), "no tile yet: one is asked for");
+    let key = install_scene(&app, &proj);
+    let st = app.session.world_state();
+    let bi = st.bodies.iter().position(|b| b.name == body).unwrap();
+    let s = app.viewport.pick_buffer.scene.lock().unwrap().clone().unwrap();
+    let id = s.edges.iter().position(|&(b, e)| b as usize == bi && e as usize == index).unwrap() as u32 + 1;
+    let size = view_size(&app, &proj);
+    let (origin, tsize) = tile_around([at.x as u32, at.y as u32], size);
+    let tile = |ids: Vec<u32>| PickTile { key, origin, size: tsize, ids };
+    let n = (tsize[0] * tsize[1]) as usize;
+    // The edge's id at the cursor's pixel.
+    let mut ids = vec![0; n];
+    ids[((at.y as u32 - origin[1]) * tsize[0] + (at.x as u32 - origin[0])) as usize] = id;
+    app.viewport.pick_buffer.out.lock().unwrap().tiles = vec![tile(ids)];
+    assert_eq!(pick(&app, &proj, at), cpu, "the named edge is picked as the CPU picks it");
+    // A tile naming no edge there: no edge or vertex is picked.
+    app.viewport.pick_buffer.out.lock().unwrap().tiles = vec![tile(vec![0; n])];
+    let none = pick(&app, &proj, at);
+    assert!(!none.iter().any(|h| matches!(h, Hit::Edge { .. } | Hit::Vertex { .. })), "{none:?}");
+    assert!(none.iter().any(|h| matches!(h, Hit::Face { .. })), "faces are still picked: {none:?}");
+    // A tile of another view does not answer.
+    app.viewport.pick_buffer.out.lock().unwrap().tiles = vec![PickTile { key: key ^ 1, ..tile(vec![0; n]) }];
+    assert_eq!(pick(&app, &proj, at), cpu);
+    // Wireframe: nothing hides an edge, and the buffer is not used.
+    let mut wire = app;
+    wire.ui.visual_style = 2;
+    wire.viewport.pick_buffer.out.lock().unwrap().tiles = vec![tile(vec![0; n])];
+    assert!(edges_near(&wire, &proj, at, &wire.session.world_state()).is_none());
+}
+
+/// On a real device: the pass drawn for the view, and picks over a grid of the view helped by its
+/// tiles answer as the CPU picks (an edge's pixels and the CPU's visibility test can disagree by
+/// a pixel at a silhouette, so a few may differ). Skipped without a GPU adapter.
+#[cfg(not(target_arch = "wasm32"))]
+#[test]
+fn picks_through_the_pass_match_the_cpu() {
+    let Some((_, device, queue)) = crate::gpu::offscreen_device() else {
+        assert!(std::env::var_os("SOLIDCRAFT_REQUIRE_GPU").is_none(), "a GPU is required here");
+        eprintln!("skipping: no GPU adapter");
+        return;
+    };
+    let mut app = app();
+    for (perspective, yaw, pitch) in [(false, 0.6, 0.5), (true, 2.4, 0.3)] {
+        app.ui.perspective = perspective;
+        app.cam.yaw = yaw;
+        app.cam.pitch = pitch;
+        let proj = projection(&app, RECT);
+        let key = install_scene(&app, &proj);
+        let s = app.viewport.pick_buffer.scene.lock().unwrap().clone().unwrap();
+        let e = proj.cam.eye().to_f32();
+        let view = PickView {
+            key,
+            view_proj: proj.vp.to_f32(),
+            eye: [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }],
+            back: proj.cam.back().to_f32(),
+            lift: lift(&app) as f32,
+            size: view_size(&app, &proj),
+        };
+        let mut pb = PickBuffer::new(&device);
+        let out = app.viewport.pick_buffer.out.clone();
+        // Tiles over the whole view.
+        let mut wanted = Vec::new();
+        for ty in (0..view.size[1]).step_by(TILE as usize - 32) {
+            for tx in (0..view.size[0]).step_by(TILE as usize - 32) {
+                wanted.push([tx + TILE / 2, ty + TILE / 2]);
+            }
+        }
+        let mut tiles = Vec::new();
+        for c in wanted {
+            for _ in 0..6 {
+                let mut encoder = device.create_command_encoder(&Default::default());
+                pb.prepare(&device, &queue, &mut encoder, s.key, &s.scene, &view, Some(c), &out);
+                queue.submit([encoder.finish()]);
+                let _ = device.poll(egui_wgpu::wgpu::PollType::wait_indefinitely());
+            }
+            tiles.append(&mut out.lock().unwrap().tiles);
+        }
+        let (mut same, mut differ, mut edges) = (0, 0, 0);
+        for j in 0..46 {
+            for i in 0..62 {
+                let p = pos2(10.0 + i as f32 * 10.0 + (j % 3) as f32, 10.0 + j as f32 * 10.0);
+                let r = 8i64;
+                let w = [p.x as i64 - r, p.y as i64 - r, p.x as i64 + r, p.y as i64 + r];
+                let Some(t) = tiles.iter().find(|t| t.covers(key, view.size, w)).cloned() else { continue };
+                out.lock().unwrap().tiles = vec![t];
+                let through = pick(&app, &proj, p);
+                out.lock().unwrap().tiles.clear();
+                let cpu = pick(&app, &proj, p);
+                edges += usize::from(cpu.iter().any(|h| matches!(h, Hit::Edge { .. })));
+                if through == cpu { same += 1 } else { differ += 1 }
+            }
+        }
+        eprintln!("perspective {perspective}: {same} same, {differ} differ, {edges} with an edge");
+        assert!(edges > 100, "the grid reaches edges ({edges})");
+        assert!(differ * 100 <= same + differ, "{differ} of {} picks differ", same + differ);
+    }
+}
diff --git a/crates/ui-egui/src/pick_index.rs b/crates/ui-egui/src/pick_index.rs
new file mode 100644
index 00000000..f19c44ea
--- /dev/null
+++ b/crates/ui-egui/src/pick_index.rs
@@ -0,0 +1,102 @@
+//! Picking without scanning every triangle and edge: each body mesh's [`MeshIndex`] (its
+//! triangles and edge segments in hierarchies), made the first time a pick reaches the body and
+//! kept while the mesh is. The viewport's picker asks it for the nearest triangle along the
+//! cursor's ray and for the edge segments near that ray, then runs its exact screen tests on
+//! those alone.
+
+use std::collections::HashMap;
+use std::sync::{Arc, Mutex};
+
+use egui::{Pos2, vec2};
+use solidcraft_engine::geom::{Mesh, MeshIndex, Vec3};
+
+use crate::viewport::Proj;
+
+/// Indexes by mesh address, each with its mesh held (so the address stays its own).
+#[derive(Default)]
+pub(crate) struct PickIndexes(Mutex<HashMap<usize, (Arc<Mesh>, Arc<MeshIndex>)>>);
+
+impl PickIndexes {
+    /// The index of `mesh`, made now when there is none yet.
+    pub(crate) fn of(&self, mesh: &Arc<Mesh>) -> Arc<MeshIndex> {
+        let key = Arc::as_ptr(mesh) as usize;
+        let Ok(mut map) = self.0.lock() else { return Arc::new(MeshIndex::new(mesh)) };
+        if let Some((_, index)) = map.get(&key) {
+            return index.clone();
+        }
+        let index = Arc::new(MeshIndex::new(mesh));
+        map.insert(key, (mesh.clone(), index.clone()));
+        index
+    }
+
+    /// Drops the indexes of meshes `keep` no longer names (a new model state).
+    pub(crate) fn retain(&self, keep: impl Fn(usize) -> bool) {
+        if let Ok(mut map) = self.0.lock() {
+            map.retain(|k, _| keep(*k));
+        }
+    }
+
+    #[cfg(test)]
+    pub(crate) fn len(&self) -> usize {
+        self.0.lock().map_or(0, |m| m.len())
+    }
+}
+
+/// How far from the ray through `pos` a point may lie and still come within `px` points of it on
+/// screen: `a + b·t` at distance `t` along the ray (constant for an orthographic view, widening
+/// with depth in perspective), with room to spare.
+pub(crate) fn reach(proj: &Proj, pos: Pos2, px: f32) -> (f64, f64) {
+    let (o, d) = proj.ray(pos);
+    let (mut a, mut b) = (0.0f64, 0.0f64);
+    for step in [vec2(px, 0.0), vec2(0.0, px), vec2(-px, 0.0), vec2(0.0, -px)] {
+        let (o2, d2) = proj.ray(pos + step);
+        a = a.max(o2.dist(o));
+        b = b.max(d2.dist(d));
+    }
+    // A screen offset of `px` along a diagonal, and rounding: half again.
+    (a * 1.5, b * 1.5)
+}
+
+/// The edge segments of `index` that may come within reach of the ray, as (edge, first point),
+/// in edge order and then along each edge: the order a scan of every edge would meet them in.
+pub(crate) fn segments_near(index: &MeshIndex, mesh: &Mesh, o: Vec3, d: Vec3, (a, b): (f64, f64)) -> Vec<(usize, usize)> {
+    let mut v = Vec::new();
+    if scanning() {
+        for (ei, e) in mesh.edges.iter().enumerate() {
+            v.extend((0..e.len().max(2) - 1).map(|si| (ei, si)));
+        }
+        return v;
+    }
+    index.segments_near(o, d, a, b, |e, s| v.push((e, s)));
+    v.sort_unstable();
+    v
+}
+
+/// The nearest triangle of `mesh` along the ray, through its index.
+pub(crate) fn raycast(index: &MeshIndex, mesh: &Mesh, o: Vec3, d: Vec3) -> Option<(f64, usize)> {
+    if scanning() {
+        return mesh.raycast(o, d);
+    }
+    index.raycast(mesh, o, d)
+}
+
+#[cfg(test)]
+thread_local! {
+    /// Tests: pick by scanning every triangle and edge, as before the indexes, to compare.
+    pub(crate) static SCAN: std::cell::Cell<bool> = const { std::cell::Cell::new(false) };
+}
+
+fn scanning() -> bool {
+    #[cfg(test)]
+    {
+        SCAN.with(|s| s.get())
+    }
+    #[cfg(not(test))]
+    {
+        false
+    }
+}
+
+#[cfg(test)]
+#[path = "pick_index_tests.rs"]
+mod tests;
diff --git a/crates/ui-egui/src/pick_index_tests.rs b/crates/ui-egui/src/pick_index_tests.rs
new file mode 100644
index 00000000..ed966233
--- /dev/null
+++ b/crates/ui-egui/src/pick_index_tests.rs
@@ -0,0 +1,122 @@
+use egui::{Rect, pos2, vec2};
+use serde_json::json;
+
+use crate::SolveApp;
+use crate::viewport::{Hit, pick, pick_all, projection};
+
+/// Overlapping bodies with flat and curved faces, filleted edges and seams.
+fn assembly() -> SolveApp {
+    let mut s = solidcraft_engine::Session::default();
+    s.execute("solid.box", &json!({"length": 40, "width": 30, "height": 20})).unwrap();
+    s.execute("solid.fillet", &json!({"edges": [[40, 15, 20]], "radius": 4})).unwrap();
+    s.execute("solid.cylinder", &json!({"radius": 8, "height": 35, "base": [10, 10, 0], "operation": "new_body"})).unwrap();
+    s.execute("solid.sphere", &json!({"radius": 9, "center": [30, 25, 22]})).unwrap();
+    let mut app = SolveApp::new(s, crate::Services::default());
+    app.ui.show_origin = false;
+    app
+}
+
+/// Every cursor position on a grid over the view, for each camera: the indexed pick answers what
+/// scanning every triangle and edge answers (the nearest of each kind, and Select Other's list).
+#[test]
+fn indexed_picks_match_the_full_scan() {
+    let mut app = assembly();
+    let rect = Rect::from_min_size(pos2(0.0, 0.0), vec2(640.0, 480.0));
+    let mut compared = (0, 0, 0);
+    for (perspective, yaw, pitch) in [(false, 0.6, 0.5), (true, 0.6, 0.5), (true, 2.4, -0.3), (false, 0.0, 1.5)] {
+        app.ui.perspective = perspective;
+        app.cam.yaw = yaw;
+        app.cam.pitch = pitch;
+        app.cam.target = solidcraft_engine::geom::Vec3::new(20.0, 15.0, 12.0);
+        app.cam.distance = 110.0;
+        let proj = projection(&app, rect);
+        for j in 0..48 {
+            for i in 0..64 {
+                let p = pos2(5.0 + i as f32 * 10.0 + (j % 3) as f32, 5.0 + j as f32 * 10.0);
+                let fast = (pick(&app, &proj, p), pick_all(&app, &proj, p));
+                super::SCAN.with(|s| s.set(true));
+                let scan = (pick(&app, &proj, p), pick_all(&app, &proj, p));
+                super::SCAN.with(|s| s.set(false));
+                assert_eq!(fast, scan, "at {p:?} (perspective {perspective}, yaw {yaw}, pitch {pitch})");
+                compared.0 += usize::from(fast.0.iter().any(|h| matches!(h, Hit::Face { .. })));
+                compared.1 += usize::from(fast.0.iter().any(|h| matches!(h, Hit::Edge { .. })));
+                compared.2 += usize::from(fast.0.iter().any(|h| matches!(h, Hit::Vertex { .. })));
+            }
+        }
+    }
+    // The grid reaches every kind many times over.
+    assert!(compared.0 > 2000 && compared.1 > 100 && compared.2 > 10, "{compared:?}");
+    assert_eq!(app.viewport.pick_index.len(), 3, "one index per body mesh");
+}
+
+/// A new model state lets go of the indexes of meshes it no longer has.
+#[test]
+fn indexes_follow_the_model_state() {
+    let mut app = assembly();
+    let rect = Rect::from_min_size(pos2(0.0, 0.0), vec2(640.0, 480.0));
+    app.cam.target = solidcraft_engine::geom::Vec3::new(20.0, 15.0, 12.0);
+    let proj = projection(&app, rect);
+    let centre = proj.to_screen(solidcraft_engine::geom::Vec3::new(20.0, 15.0, 10.0)).unwrap();
+    pick(&app, &proj, centre);
+    assert!(app.viewport.pick_index.len() >= 1);
+    app.session.execute("solid.box", &json!({"length": 5, "width": 5, "height": 5, "corner": [-20, -20, 0], "operation": "new_body"})).unwrap();
+    pick(&app, &proj, centre);
+    assert!(app.viewport.pick_index.len() <= 4, "indexes of meshes gone with the old state are dropped");
+    app.session.execute("solid.box", &json!({"length": 5, "width": 5, "height": 5, "corner": [-40, -20, 0], "operation": "new_body"})).unwrap();
+    pick(&app, &proj, centre);
+    assert!(app.viewport.pick_index.len() <= 5);
+}
+
+/// Timing, not a check (`cargo test -p solidcraft-ui-egui --lib pick_timing -- --ignored
+/// --nocapture`): picks over a grid of the view, scanning every triangle and edge versus through
+/// the indexes, on many curved bodies and on one body with many curved faces.
+#[test]
+#[ignore]
+fn pick_timing_on_curved_models() {
+    let mut many = solidcraft_engine::Session::default();
+    for i in 0..12 {
+        for j in 0..12 {
+            let (x, y) = (f64::from(i) * 30.0, f64::from(j) * 30.0);
+            many.execute("solid.cylinder", &json!({"radius": 9, "height": 30, "base": [x, y, 0], "operation": "new_body"})).unwrap();
+            many.execute("solid.sphere", &json!({"radius": 8, "center": [x, y, 40], "operation": "new_body"})).unwrap();
+        }
+    }
+    let mut one = solidcraft_engine::Session::default();
+    one.execute("solid.box", &json!({"length": 340, "width": 340, "height": 10, "corner": [-15, -15, -5]})).unwrap();
+    for i in 0..8 {
+        for j in 0..8 {
+            let (x, y) = (f64::from(i) * 42.0, f64::from(j) * 42.0);
+            one.execute("solid.cylinder", &json!({"radius": 14, "height": 30, "base": [x, y, 0], "operation": "join"})).unwrap();
+        }
+    }
+    for (name, s) in [("many bodies", many), ("one body", one)] {
+        let mut app = SolveApp::new(s, crate::Services::default());
+        app.ui.show_origin = false;
+        app.cam.target = solidcraft_engine::geom::Vec3::new(160.0, 160.0, 15.0);
+        app.cam.distance = 900.0;
+        let st = app.session.world_state();
+        let (tris, segs) = st
+            .bodies
+            .iter()
+            .map(|b| b.mesh())
+            .fold((0, 0), |(t, s), m| (t + m.triangles.len(), s + m.edges.iter().map(|e| e.len().saturating_sub(1)).sum::<usize>()));
+        let rect = Rect::from_min_size(pos2(0.0, 0.0), vec2(1400.0, 900.0));
+        let proj = projection(&app, rect);
+        let points: Vec<_> = (0..20).flat_map(|j| (0..30).map(move |i| pos2(30.0 + i as f32 * 45.0, 30.0 + j as f32 * 42.0))).collect();
+        // Indexes made once, before timing.
+        for p in &points {
+            pick(&app, &proj, *p);
+        }
+        let time = |scan: bool| {
+            super::SCAN.with(|s| s.set(scan));
+            let t = std::time::Instant::now();
+            for p in &points {
+                std::hint::black_box(pick(&app, &proj, *p));
+            }
+            super::SCAN.with(|s| s.set(false));
+            t.elapsed().as_secs_f64() * 1000.0 / points.len() as f64
+        };
+        let (scan, indexed) = (time(true), time(false));
+        eprintln!("{name}: {} bodies, {tris} triangles, {segs} edge segments: scan {scan:.3} ms per pick, indexed {indexed:.3} ms", st.bodies.len());
+    }
+}
diff --git a/crates/ui-egui/src/viewport.rs b/crates/ui-egui/src/viewport.rs
index 2dd0bc0d..3fb11a41 100644
--- a/crates/ui-egui/src/viewport.rs
+++ b/crates/ui-egui/src/viewport.rs
@@ -55,6 +55,11 @@ pub struct ViewportState {
     lod_view: u64,
     /// Each body mesh's box (the mesh held, so its address stays its own), from the last state.
     mesh_bounds: std::sync::Mutex<std::collections::HashMap<usize, (std::sync::Arc<solidcraft_engine::geom::Mesh>, solidcraft_engine::geom::Aabb3)>>,
+    /// Each body mesh's picking index (made the first time a pick reaches the body), kept while
+    /// the mesh is part of the world state.
+    pub(crate) pick_index: crate::pick_index::PickIndexes,
+    /// The GPU pick buffer's scene, read-backs and requests.
+    pub(crate) pick_buffer: crate::pick_buffer::State,
     /// The last click in the view was the middle button's (for the middle double-click).
     pub(crate) middle_clicked: bool,
     /// The last pick (its key and hits): reused while nothing it depends on changed.
@@ -1003,6 +1008,7 @@ fn body_bounds(
                 })
                 .collect::<Vec<_>>(),
         );
+        app.viewport.pick_index.retain(|k| now.contains_key(&k));
         if let Ok(mut k) = app.viewport.mesh_bounds.lock() {
             *k = now;
         }
@@ -1269,7 +1275,7 @@ fn pick_hits(app: &SolveApp, proj: &Proj, pos: Pos2, other_sketch_points: bool,
             }
         }
         let m = b.mesh();
-        if let Some((t, ti)) = m.raycast(o, d) {
+        if let Some((t, ti)) = crate::pick_index::raycast(&app.viewport.pick_index.of(&m), &m, o, d) {
             let fi = m.tri_face.get(ti).copied().unwrap_or(0) as usize;
             let hit = Hit::Face { body: b.name.clone(), index: fi, point: o + d * t };
             if bestf.as_ref().is_none_or(|(bt, _)| t < *bt) {
@@ -1284,7 +1290,11 @@ fn pick_hits(app: &SolveApp, proj: &Proj, pos: Pos2, other_sketch_points: bool,
     // Wireframe draws no faces, so nothing hides an edge or a vertex.
     let wireframe = app.ui.visual_style == 2;
     let visible = |p: Vec3| wireframe || bestf.as_ref().is_none_or(|(t, _)| (p - o).dot(d) <= *t + slack);
-    // Vertices (edge end points) and edges.
+    // Vertices (edge end points) and edges, within 7 and 6 points of the cursor.
+    let reach = crate::pick_index::reach(proj, pos, 8.0);
+    // The edges seen near the cursor, when the GPU pick buffer has them for this view (Select
+    // Other wants hidden ones too).
+    let seen = if all { None } else { crate::pick_buffer::edges_near(app, proj, pos, &st) };
     let mut bestv = Near::new(all);
     let mut beste = Near::new(all);
     for (i, b) in st.bodies.iter().enumerate() {
@@ -1298,26 +1308,37 @@ fn pick_hits(app: &SolveApp, proj: &Proj, pos: Pos2, other_sketch_points: bool,
             continue;
         }
         let m = b.mesh();
-        for (ei, e) in m.edges.iter().enumerate() {
+        // Only the segments that can come near the cursor: of the edges the pick buffer sees
+        // there, else from the index. Either way in the order a scan of every edge would meet
+        // them, so ties go the same way.
+        let near = match &seen {
+            Some(seen) => seen.segments(i, &m),
+            None => crate::pick_index::segments_near(&app.viewport.pick_index.of(&m), &m, o, d, reach),
+        };
+        let mut last_edge = None;
+        for (ei, si) in near {
+            let Some(e) = m.edges.get(ei) else { continue };
             if !m.edge_pickable(ei) {
                 continue;
             }
-            for v in [e.first(), e.last()].into_iter().flatten() {
-                if let Some(sp) = proj.to_screen(*v) {
-                    let dd = sp.distance(pos);
-                    if dd < 7.0 && visible(*v) && bestv.wants(dd) {
-                        bestv.offer(dd, Hit::Vertex { body: b.name.clone(), point: *v });
+            if last_edge != Some(ei) {
+                last_edge = Some(ei);
+                for v in [e.first(), e.last()].into_iter().flatten() {
+                    if let Some(sp) = proj.to_screen(*v) {
+                        let dd = sp.distance(pos);
+                        if dd < 7.0 && visible(*v) && bestv.wants(dd) {
+                            bestv.offer(dd, Hit::Vertex { body: b.name.clone(), point: *v });
+                        }
                     }
                 }
             }
-            for w in e.windows(2) {
-                let (Some(a), Some(c)) = (proj.to_screen(w[0]), proj.to_screen(w[1])) else { continue };
-                let dd = seg_dist(pos, a, c);
-                // Visible where the cursor is: the edge's point under the cursor, not an end or
-                // the middle, which can lie deeper than the face the cursor is on.
-                if dd < 6.0 && beste.wants(dd) && visible(ray_segment_closest(o, d, w[0], w[1])) {
-                    beste.offer(dd, Hit::Edge { body: b.name.clone(), index: ei, mid: polyline_mid(e) });
-                }
+            let (Some(p0), Some(p1)) = (e.get(si), e.get(si + 1)) else { continue };
+            let (Some(a), Some(c)) = (proj.to_screen(*p0), proj.to_screen(*p1)) else { continue };
+            let dd = seg_dist(pos, a, c);
+            // Visible where the cursor is: the edge's point under the cursor, not an end or
+            // the middle, which can lie deeper than the face the cursor is on.
+            if dd < 6.0 && beste.wants(dd) && visible(ray_segment_closest(o, d, *p0, *p1)) {
+                beste.offer(dd, Hit::Edge { body: b.name.clone(), index: ei, mid: polyline_mid(e) });
             }
         }
     }
@@ -2377,6 +2398,7 @@ pub fn show(app: &mut SolveApp, ui: &mut egui::Ui) {
             },
         });
         painter.add(egui_wgpu::Callback::new_paint_callback(rect, cb));
+        crate::pick_buffer::paint(app, ui.ctx(), &painter, &proj);
     }
     analysis_legend(app, &painter, rect);
     crate::frame_rate::paint(app, ui.ctx(), &painter, rect);