Release 0.10.1

* GPU tessellation: a pipeline the device cannot build falls back to CPU meshing instead of crashing (#1098)

Build the tessellation compute pipeline inside error scopes; TessGpu::new returns an error,
and the viewport logs it and meshes bodies on the CPU. A Windows CI job builds the pipeline
with D3D12's FXC on the WARP adapter.

Co-Authored-By: Claude Opus 5.5 <[email protected]>

* Tessellation shader: basis functions in one-dimensional local arrays, so D3D12's FXC compiles it (#1098)

FXC has to unroll loops that write an array inside a struct (Basis.n0/n1) or a 2-D array
(ndu) by a counter, and cannot with run-time trip counts (X3511). A2.2 basis functions in a
local 1-D array returned whole; derivatives from the degree p-1 functions; the pole step
offsets computed rather than indexed.

Co-Authored-By: Claude Opus 5.5 <[email protected]>

---------

Co-authored-by: Claude Opus 5.5 <[email protected]>

.github/workflows/ci.yml

@@ -286,9 +286,31 @@ jobs:
           name: oracle-report
           path: docs/oracle.md
 
+  # D3D12's older shader compiler (FXC, used where dxcompiler.dll is absent) is stricter than
+  # DXC, Vulkan and Metal: a shader it cannot compile once crashed the app at startup. Build the
+  # GPU tessellation pipeline with FXC on the runner's WARP software adapter and compare its
+  # results with the CPU.
+  windows-fxc:
+    name: Windows D3D12 FXC shaders
+    needs: changes
+    if: needs.changes.outputs.native == 'true'
+    runs-on: windows-latest
+    timeout-minutes: 45
+    env:
+      WGPU_BACKEND: dx12
+      WGPU_DX12_COMPILER: fxc
+      SOLIDCRAFT_REQUIRE_GPU: "1"
+    steps:
+      - uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1
+      - uses: dtolnay/rust-toolchain@stable
+      - uses: Swatinem/rust-cache@6323deb102c322ba6fcbdcafc7e3dddab59af2b6 # v2.9.2
+        with:
+          key: windows-fxc
+      - run: cargo test -p solidcraft-gpu --profile ci --locked
+
   # Preserve the existing check name for branch protection, including docs-only PRs.
   ci:
-    needs: [changes, checks, clippy, build-tests, tests, doc-tests, oracle]
+    needs: [changes, checks, clippy, build-tests, tests, doc-tests, oracle, windows-fxc]
     if: always()
     runs-on: ubuntu-24.04
     steps:
@@ -302,6 +324,7 @@ jobs:
           TESTS: ${{ needs.tests.result }}
           DOCS: ${{ needs.doc-tests.result }}
           ORACLE: ${{ needs.oracle.result }}
+          WINDOWS_FXC: ${{ needs.windows-fxc.result }}
         run: |
           test "$CHANGES" = success
           test "$CHECKS" = success
@@ -311,6 +334,7 @@ jobs:
             test "$TESTS" = success
             test "$DOCS" = success
             test "$ORACLE" = success
+            test "$WINDOWS_FXC" = success
           fi
 
   wasm:

Cargo.lock

@@ -3281,7 +3281,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "eframe",
  "egui",
@@ -3299,7 +3299,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-cli"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "serde_json",
  "solidcraft-engine",
@@ -3308,7 +3308,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-doc"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "log",
  "serde",
@@ -3323,7 +3323,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-engine"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "log",
  "serde",
@@ -3339,7 +3339,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-geom"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "serde",
  "serde_json",
@@ -3348,9 +3348,10 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-gpu"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "log",
+ "naga",
  "pollster 0.4.0",
  "solidcraft-geom",
  "solidcraft-kernel",
@@ -3360,7 +3361,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-io"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "quick-xml",
  "serde",
@@ -3375,7 +3376,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-kernel"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "log",
  "rayon",
@@ -3393,7 +3394,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-mcp"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "serde_json",
  "solidcraft-engine",
@@ -3401,7 +3402,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-render"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "image",
  "serde",
@@ -3411,7 +3412,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-sketch"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "serde",
  "serde_json",
@@ -3422,7 +3423,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-ui-egui"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "egui",
  "egui-wgpu",
@@ -3439,7 +3440,7 @@ dependencies = [
 
 [[package]]
 name = "solidcraft-web"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "eframe",
  "egui",
@@ -4862,7 +4863,7 @@ checksum = "e450f9b2ed1dff33c94c12589a87338689467b9c4f5d8a5710bd09a847d2c8a7"
 
 [[package]]
 name = "xtask"
-version = "0.10.0"
+version = "0.10.1"
 dependencies = [
  "serde_json",
  "sha2",

Cargo.toml

@@ -5,7 +5,7 @@ members = ["crates/*", "apps/*", "xtask"]
 exclude = ["vendor"]
 
 [workspace.package]
-version = "0.10.0"
+version = "0.10.1"
 edition = "2024"
 license = "MIT OR Apache-2.0"
 rust-version = "1.95"
@@ -61,6 +61,7 @@ web-time = "1"
 egui-wgpu = "0.36"
 # the wgpu egui-wgpu uses, for compute (crates/gpu); backends come from the apps' features
 wgpu = { version = "30", default-features = false, features = ["std", "wgsl"] }
+naga = { version = "30", default-features = false }
 pollster = "0.4"
 rfd = "0.17"
 

crates/doc/src/eval.rs

@@ -31,6 +31,9 @@ pub struct ModelBody {
     /// The display plan (none when the body has no plan: a mesh body, or one the planner
     /// refuses), shared like the mesh. A GPU viewport draws it.
     plan: Arc<OnceLock<Option<Arc<TessPlan>>>>,
+    /// The plan evaluated on the CPU (none without a plan), shared like the mesh: the triangles a
+    /// GPU viewport draws, so a highlight over a face lies exactly on it.
+    plan_mesh: Arc<OnceLock<Option<Arc<Mesh>>>>,
     /// Identity of this body's content for the evaluation cache: the key of the evaluation
     /// that produced it (kept while later features pass it through unchanged).
     key: u64,
@@ -59,6 +62,7 @@ impl ModelBody {
             feature,
             mesh: Arc::new(OnceLock::new()),
             plan: Arc::new(OnceLock::new()),
+            plan_mesh: Arc::new(OnceLock::new()),
             key: 0,
             names: Default::default(),
             made: Default::default(),
@@ -98,6 +102,14 @@ impl ModelBody {
             .clone()
     }
 
+    /// The display plan's mesh, evaluated on the CPU once: the same triangles a GPU viewport
+    /// evaluates and draws for this body. Highlights over a face drawn from the plan use it, since
+    /// the reference mesh ([`ModelBody::mesh`]) triangulates the face differently and the two
+    /// would fight at equal depth. `None` when the body has no plan.
+    pub fn plan_mesh(&self) -> Option<Arc<Mesh>> {
+        self.plan_mesh.get_or_init(|| self.display_plan().and_then(|plan| self.body.display_mesh_of(&plan).ok()).map(Arc::new)).clone()
+    }
+
     /// The display mesh's chord tolerance: 1/1000 of the body's size.
     fn display_tol(&self) -> f64 {
         (self.body.size() * 1e-3).max(1e-3)

crates/gpu/Cargo.toml

@@ -21,3 +21,5 @@ wgpu = { workspace = true, features = ["metal", "vulkan", "dx12", "gles"] }
 
 [dev-dependencies]
 solidcraft-kernel = { workspace = true }
+# The HLSL the D3D12 backend compiles (tests check it suits FXC).
+naga = { workspace = true, features = ["wgsl-in", "hlsl-out"] }

crates/gpu/src/lib.rs

@@ -17,6 +17,37 @@ use solidcraft_geom::Vec3;
 use solidcraft_geom::tess::{EvalCurve, EvalSurface, TessPlan};
 use wgpu::util::DeviceExt;
 
+/// The tessellation compute shader (WGSL).
+pub const SHADER: &str = include_str!("tess.wgsl");
+
+/// The error a popped error scope caught. Native devices answer at once; a browser's answer
+/// comes later, and is not waited for (no error is assumed).
+fn popped(f: impl std::future::Future<Output = Option<wgpu::Error>>) -> Option<wgpu::Error> {
+    let mut f = std::pin::pin!(f);
+    match f.as_mut().poll(&mut std::task::Context::from_waker(std::task::Waker::noop())) {
+        std::task::Poll::Ready(e) => e,
+        std::task::Poll::Pending => None,
+    }
+}
+
+/// An adapter and device without a window (native only). The environment may pick the backend
+/// and the D3D12 shader compiler (`WGPU_BACKEND`, `WGPU_DX12_COMPILER`).
+#[cfg(not(target_arch = "wasm32"))]
+fn open_headless() -> Result<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
+    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle_from_env());
+    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((adapter, device, queue))
+}
+
 /// 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;
@@ -29,6 +60,8 @@ pub enum GpuError {
     Device(String),
     #[error("the GPU evaluation failed: {0}")]
     Failed(String),
+    #[error("the GPU tessellation pipeline could not be created: {0}")]
+    Pipeline(String),
     #[error("not supported yet: {0}")]
     Unsupported(String),
 }
@@ -78,11 +111,47 @@ pub struct TessGpu {
 }
 
 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 {
+    /// The pipeline on an existing device (`adapter`: its name, for reports). An error (never
+    /// a panic) when the device cannot build it, for example where the D3D12 driver's older shader
+    /// compiler (FXC) rejects the shader: the caller then builds bodies on the CPU.
+    pub fn new(adapter: impl Into<String>, device: wgpu::Device, queue: wgpu::Queue) -> Result<Self> {
+        Self::with_shader(adapter, device, queue, SHADER)
+    }
+
+    /// [`TessGpu::new`] with another shader source: tests use it to make pipeline creation fail
+    /// as it does on a device whose compiler rejects the real shader.
+    #[doc(hidden)]
+    pub fn with_shader(adapter: impl Into<String>, device: wgpu::Device, queue: wgpu::Queue, shader: &str) -> Result<Self> {
+        // Errors here would go to the device's uncaptured-error handler, which panics by default:
+        // capture them, so a device that cannot run the shader only loses GPU tessellation.
+        let scopes = [wgpu::ErrorFilter::OutOfMemory, wgpu::ErrorFilter::Internal, wgpu::ErrorFilter::Validation].map(|f| device.push_error_scope(f));
+        let (pipeline, layout) = Self::build(&device, shader);
+        let mut failed = None;
+        for scope in scopes.into_iter().rev() {
+            if let Some(e) = popped(scope.pop()) {
+                failed.get_or_insert(e.to_string());
+            }
+        }
+        if let Some(e) = failed {
+            return Err(GpuError::Pipeline(e));
+        }
+        // A viewport records the evaluation into its own frame (see `TessGpu::record`).
+        #[cfg(target_arch = "wasm32")]
+        let _ = queue;
+        Ok(TessGpu {
+            adapter: adapter.into(),
+            device,
+            #[cfg(not(target_arch = "wasm32"))]
+            queue,
+            pipeline,
+            layout,
+        })
+    }
+
+    fn build(device: &wgpu::Device, shader: &str) -> (wgpu::ComputePipeline, wgpu::BindGroupLayout) {
         let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("solidcraft tessellation"),
-            source: wgpu::ShaderSource::Wgsl(include_str!("tess.wgsl").into()),
+            source: wgpu::ShaderSource::Wgsl(shader.into()),
         });
         let entry = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry {
             binding,
@@ -114,34 +183,14 @@ impl TessGpu {
             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,
-        }
+        (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))
+        let (adapter, device, queue) = open_headless()?;
+        Self::new(adapter.get_info().name, device, queue)
     }
 
     /// The adapter's name, for reports.

crates/gpu/src/tess.wgsl

@@ -1,7 +1,7 @@
 // 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
+// the same algorithms (The NURBS Book A2.1 span search, A2.2 basis functions, A2.3 first
 // derivatives, the rational quotient rule), in f32.
 
 const PLANE: u32 = 0u;
@@ -60,11 +60,6 @@ struct Derivs {
     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]);
 }
@@ -101,42 +96,58 @@ fn span_of(k0: u32, p: u32, n: u32, t: f32) -> u32 {
     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;
+// The arrays below are one-dimensional locals, written by loop counters and returned whole:
+// D3D12's older shader compiler (FXC) must unroll a loop that writes an array inside a struct
+// (or a two-dimensional one) by a counter, and it cannot unroll loops whose trip count is only
+// known at run time, so such a loop fails to compile there ("unable to unroll loop").
+
+// The non-zero basis functions of degree p on span `span` (The NURBS Book A2.2).
+fn basis_funs(k0: u32, span: u32, p: u32, t: f32) -> array<f32, W> {
+    var n: array<f32, W>;
+    n[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];
+            // right[r + 1] and left[j - r] of A2.2 (span ≥ p ≥ j, so no underflow).
+            let right = data[k0 + span + r + 1u] - t;
+            let left = t - data[k0 + span + 1u + r - j];
+            let den = right + left;
             var temp = 0.0;
-            if (ndu[j][r] != 0.0) {
-                temp = ndu[r][j - 1u] / ndu[j][r];
+            if (den != 0.0) {
+                temp = n[r] / den;
             }
-            ndu[r][j] = saved + right[r + 1u] * temp;
-            saved = left[j - r] * temp;
+            n[r] = saved + right * temp;
+            saved = left * temp;
         }
-        ndu[j][j] = saved;
+        n[j] = saved;
     }
-    var b: Basis;
-    for (var j = 0u; j <= p; j = j + 1u) {
-        b.n0[j] = ndu[j][p];
+    return n;
+}
+
+// Their first derivatives, from the degree p - 1 functions on the same span (A2.3, n = 1).
+fn basis_ders(k0: u32, span: u32, p: u32, t: f32) -> array<f32, W> {
+    var d: array<f32, W>;
+    if (p == 0u) {
+        return d;
     }
+    let lower = basis_funs(k0, span, p - 1u, t);
     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];
+        var x = 0.0;
+        if (r >= 1u) {
+            let den = data[k0 + span + r] - data[k0 + span + r - p];
+            if (den != 0.0) {
+                x = x + lower[r - 1u] / den;
+            }
         }
-        if (r < p && ndu[p][r] != 0.0) {
-            d = d - ndu[r][p - 1u] / ndu[p][r];
+        if (r < p) {
+            let den = data[k0 + span + r + 1u] - data[k0 + span + r + 1u - p];
+            if (den != 0.0) {
+                x = x - lower[r] / den;
+            }
         }
-        b.n1[r] = d * f32(p);
+        d[r] = x * f32(p);
     }
-    return b;
+    return d;
 }
 
 // Point and derivative of a rational function from its homogeneous value and derivative.
@@ -152,17 +163,19 @@ 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);
+    let bu0 = basis_funs(s.knots_u, su, s.pu, u);
+    let bu1 = basis_ders(s.knots_u, su, s.pu, u);
+    let bv0 = basis_funs(s.knots_v, sv, s.pv, v);
+    let bv1 = basis_ders(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;
+            a = a + bu0[i] * bv0[j] * c;
+            au = au + bu1[i] * bv0[j] * c;
+            av = av + bu0[i] * bv1[j] * c;
         }
     }
     r.ok = abs(a.w) > 0.0;
@@ -189,13 +202,14 @@ fn revolved(s: Surface, u: f32, v: f32) -> Derivs {
     }
     // 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);
+    let b0 = basis_funs(s.knots_u, span, s.pu, t);
+    let b1 = basis_ders(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;
+        c = c + b0[i] * h;
+        dc = dc + b1[i] * h;
     }
     r.ok = abs(c.w) > 0.0;
     if (!r.ok) {
@@ -263,9 +277,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     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));
+        // Steps (h, h), (-h, h), (h, -h), (-h, -h), without an array indexed by k (FXC).
         for (var k = 0u; k < 4u; k = k + 1u) {
-            let e = derivs(s, vert.u + steps[k].x, vert.v + steps[k].y);
+            let du = select(h, -h, (k & 1u) != 0u);
+            let dv = select(h, -h, (k & 2u) != 0u);
+            let e = derivs(s, vert.u + du, vert.v + dv);
             let m = oriented(s, cross(e.du, e.dv));
             if (e.ok && length(m) > 0.0) {
                 n = m;

crates/gpu/src/tests.rs

@@ -4,12 +4,17 @@ 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).
+/// A device that cannot build the pipeline fails the test: that is a shader the device rejects
+/// (as D3D12's FXC once did), not a missing GPU. `SOLIDCRAFT_REQUIRE_GPU=1` (the Windows CI job,
+/// on the WARP software adapter) fails without an adapter too.
 fn gpu() -> Option<TessGpu> {
     match TessGpu::headless() {
         Ok(g) => {
             eprintln!("GPU: {}", g.adapter_name());
             Some(g)
         }
+        Err(e @ GpuError::Pipeline(_)) => panic!("{e}"),
+        Err(e) if std::env::var_os("SOLIDCRAFT_REQUIRE_GPU").is_some() => panic!("a GPU is required here: {e}"),
         Err(e) => {
             eprintln!("skipping the GPU comparison: {e}");
             None
@@ -17,6 +22,80 @@ fn gpu() -> Option<TessGpu> {
     }
 }
 
+/// A shader the device rejects when the pipeline is created (its binding 1 is writable, the
+/// layout's read-only), standing in for one a device's compiler cannot build.
+const REJECTED: &str = "
+@group(0) @binding(1) var<storage, read_write> surfaces: array<u32>;
+@group(0) @binding(4) var<storage, read_write> out: array<vec4<f32>>;
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    surfaces[gid.x] = 1u;
+    out[gid.x] = vec4<f32>(0.0);
+}
+";
+
+/// A pipeline the device cannot build is an error the caller can fall back from (to the CPU
+/// mesh), not a panic in the device's uncaptured-error handler; and the device stays usable.
+#[test]
+fn a_pipeline_the_device_rejects_is_an_error_not_a_crash() {
+    let (adapter, device, queue) = match 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 name = adapter.get_info().name;
+    match TessGpu::with_shader(name.clone(), device.clone(), queue.clone(), REJECTED) {
+        Err(GpuError::Pipeline(e)) => eprintln!("rejected as expected: {e}"),
+        Err(e) => panic!("unexpected error: {e}"),
+        Ok(_) => panic!("the rejected shader built a pipeline"),
+    }
+    // Not wgsl at all: also an error, not a panic.
+    assert!(matches!(TessGpu::with_shader(name.clone(), device.clone(), queue.clone(), "fn main( {"), Err(GpuError::Pipeline(_))));
+    // The same device then builds the real pipeline and evaluates on it.
+    let gpu = ok(TessGpu::new(name, device, queue));
+    let body = ok(solidcraft_kernel::torus(Vec3::ZERO, 6.0, 1.5));
+    let plan = ok(body.display_plan(0.01));
+    assert_eq!(ok(gpu.evaluate(&plan)).positions.len(), plan.evaluate().positions.len());
+}
+
+/// The shader as the D3D12 backend translates it (naga's HLSL for shader model 5.1, the FXC
+/// target, with wgpu's options) avoids what FXC cannot compile without unrolling a loop:
+/// arrays inside structs and two-dimensional arrays written by a loop counter. Runs everywhere;
+/// the Windows CI job compiles the real thing with FXC.
+#[test]
+fn the_hlsl_for_fxc_has_no_arrays_fxc_must_unroll_loops_for() {
+    use naga::back::hlsl;
+    let module = ok(naga::front::wgsl::parse_str(SHADER));
+    let info = ok(naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::default()).validate(&module));
+    let options = hlsl::Options {
+        shader_model: hlsl::ShaderModel::V5_1,
+        force_loop_bounding: true,
+        restrict_indexing: true,
+        fake_missing_bindings: true,
+        ..Default::default()
+    };
+    let pipeline = hlsl::PipelineOptions::default();
+    let mut out = String::new();
+    ok(hlsl::Writer::new(&mut out, &options, &pipeline).write(&module, &info, None));
+    assert!(out.contains("void main("), "{out}");
+    // Struct members that are arrays: `float n0_[8];` inside a `struct ... { };`.
+    let mut in_struct = false;
+    for line in out.lines() {
+        let l = line.trim();
+        if l.starts_with("struct ") && !l.ends_with(';') {
+            in_struct = true;
+        } else if l.starts_with("};") {
+            in_struct = false;
+        } else if in_struct {
+            assert!(!l.contains('['), "an array inside a struct: {l}");
+        }
+        assert!(!l.contains("]["), "a two-dimensional array: {l}");
+    }
+}
+
 fn ok<T, E: std::fmt::Debug>(r: std::result::Result<T, E>) -> T {
     match r {
         Ok(v) => v,

crates/ui-egui/src/gpu.rs

@@ -642,7 +642,7 @@ impl Resources {
             access_bgl,
             access_bind,
             access_version: 0,
-            tri_hl: pipeline("sc_tris_hl", "vs_tri", "fs_tri", tri_layout(), depth(wgpu::CompareFunction::LessEqual, false), alpha),
+            tri_hl: pipeline("sc_tris_hl", "vs_tri", "fs_tri", tri_layout(), highlight_depth(t.depth), alpha),
             trans: pipeline("sc_trans", "vs_tri", "fs_tri", tri_layout(), depth(wgpu::CompareFunction::LessEqual, false), alpha),
             glass_far: glass("sc_glass_far", wgpu::Face::Front),
             glass_near: glass("sc_glass_near", wgpu::Face::Back),
@@ -662,6 +662,24 @@ impl Resources {
     }
 }
 
+/// Depth state of faces lit over the model (hover, selection): tested, not written, and pulled a
+/// little toward the eye. They lie on the model's own faces, but the GPU evaluates a body's plan
+/// in f32 while the highlight's copy of the same triangles is evaluated on the CPU, so the two
+/// depths differ by rounding and an unbiased equal test lights the face in speckles and wedges.
+/// The bias is a few depth steps plus the face's own slope: far less than any real gap.
+fn highlight_depth(format: Option<wgpu::TextureFormat>) -> Option<wgpu::DepthStencilState> {
+    format.map(|format| wgpu::DepthStencilState {
+        format,
+        depth_write_enabled: Some(false),
+        depth_compare: Some(wgpu::CompareFunction::LessEqual),
+        stencil: Default::default(),
+        bias: HIGHLIGHT_BIAS,
+    })
+}
+
+/// See [`highlight_depth`]. Negative: toward the eye (depth grows away from it).
+const HIGHLIGHT_BIAS: wgpu::DepthBiasState = wgpu::DepthBiasState { constant: -8, slope_scale: -1.0, clamp: 0.0 };
+
 fn upload(device: &wgpu::Device, label: &str, bytes: &[u8], stride: usize) -> Option<Batch> {
     upload_with(device, label, bytes, stride, wgpu::BufferUsages::VERTEX)
 }
@@ -718,14 +736,30 @@ impl Batches {
     }
 }
 
+/// The tessellation pipeline (`shader`: its source), where display plans are evaluated: on a
+/// device that runs compute shaders with the storage buffers the evaluation binds (not WebGL2)
+/// and builds the pipeline. Elsewhere `None`, and bodies are built on the CPU: a device whose
+/// shader compiler rejects the shader (D3D12's FXC once did) only loses GPU evaluation.
+fn tessellation(adapter: &wgpu::Adapter, device: &wgpu::Device, queue: &wgpu::Queue, shader: &str) -> Option<solidcraft_gpu::TessGpu> {
+    let compute = adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS)
+        && device.limits().max_storage_buffers_per_shader_stage >= 4;
+    if !compute {
+        return None;
+    }
+    let info = adapter.get_info();
+    match solidcraft_gpu::TessGpu::with_shader(info.name.clone(), device.clone(), queue.clone(), shader) {
+        Ok(t) => Some(t),
+        Err(e) => {
+            log::warn!("bodies are built on the CPU on {} ({:?}): {e}", info.name, info.backend);
+            None
+        }
+    }
+}
+
 /// Create the pipelines. Call once with the app's render state and its depth/stencil/MSAA
 /// settings. Without a stencil the section cap does not hide the inside of cut bodies.
 pub fn install(rs: &egui_wgpu::RenderState, depth_bits: u8, stencil_bits: u8, samples: u32) -> GpuTarget {
-    // Display plans are evaluated where the device runs compute shaders with the storage
-    // buffers the evaluation binds (not WebGL2): elsewhere bodies are built on the CPU.
-    let compute = rs.adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS)
-        && rs.device.limits().max_storage_buffers_per_shader_stage >= 4;
-    let tess = compute.then(|| solidcraft_gpu::TessGpu::new(rs.adapter.get_info().name, rs.device.clone(), rs.queue.clone()));
+    let tess = tessellation(&rs.adapter, &rs.device, &rs.queue, solidcraft_gpu::SHADER);
     let t = GpuTarget { format: rs.target_format, depth: egui_wgpu::depth_format_from_bits(depth_bits, stencil_bits), samples, tess: tess.is_some() };
     let res = Resources::new(&rs.device, t, tess);
     rs.renderer.write().callback_resources.insert(res);
@@ -894,19 +928,43 @@ impl egui_wgpu::CallbackTrait for ViewportCallback {
 
 /// Render `scene` offscreen on a device of its own, through the viewport's real resources and
 /// frame callback (`prepare`, then `paint`), and read the pixels back: straight RGBA, `size` ×
-/// `size`. `None` without a GPU adapter. With `tess`, planned bodies are evaluated on the GPU.
+/// `size`. `None` without a GPU adapter. With `tess` (the tessellation shader), planned bodies
+/// are evaluated on the GPU where the device builds that pipeline, as [`install`] sets it up.
+/// `highlight` is drawn as the viewport's highlight scene (hover and selection).
 #[cfg(all(test, not(target_arch = "wasm32")))]
-pub(crate) fn render_offscreen(scene: GpuScene, view_proj: [[f32; 4]; 4], back: [f32; 3], eye: [f32; 4], size: u32, tess: bool) -> Option<Vec<u8>> {
-    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle());
+fn offscreen_device() -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
+    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle_from_env());
     let adapter = pollster::block_on(instance.request_adapter(&Default::default())).ok()?;
     let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
         required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()),
         ..Default::default()
     }))
     .ok()?;
+    Some((adapter, device, queue))
+}
+
+/// Whether [`install`] would evaluate plans on the GPU with `shader`, on an offscreen device
+/// (`None` without a GPU adapter).
+#[cfg(all(test, not(target_arch = "wasm32")))]
+pub(crate) fn offscreen_tessellates(shader: &str) -> Option<bool> {
+    let (adapter, device, queue) = offscreen_device()?;
+    Some(tessellation(&adapter, &device, &queue, shader).is_some())
+}
+
+#[cfg(all(test, not(target_arch = "wasm32")))]
+pub(crate) fn render_offscreen(
+    scene: GpuScene,
+    highlight: Option<GpuScene>,
+    view_proj: [[f32; 4]; 4],
+    back: [f32; 3],
+    eye: [f32; 4],
+    size: u32,
+    tess: Option<&str>,
+) -> Option<Vec<u8>> {
+    let (adapter, device, queue) = offscreen_device()?;
     let (format, depth) = (wgpu::TextureFormat::Rgba8Unorm, wgpu::TextureFormat::Depth24PlusStencil8);
-    let target = GpuTarget { format, depth: Some(depth), samples: 1, tess };
-    let compute = tess.then(|| solidcraft_gpu::TessGpu::new(adapter.get_info().name, device.clone(), queue.clone()));
+    let compute = tess.and_then(|shader| tessellation(&adapter, &device, &queue, shader));
+    let target = GpuTarget { format, depth: Some(depth), samples: 1, tess: compute.is_some() };
     let mut resources = egui_wgpu::CallbackResources::default();
     resources.insert(Resources::new(&device, target, compute));
     let none = || Arc::new(Mutex::new(None));
@@ -915,8 +973,8 @@ pub(crate) fn render_offscreen(scene: GpuScene, view_proj: [[f32; 4]; 4], back:
         slot: Arc::new(Mutex::new(Some(scene))),
         sketch_key: 0,
         sketch_slot: none(),
-        hl_key: 0,
-        hl_slot: none(),
+        hl_key: u64::from(highlight.is_some()),
+        hl_slot: Arc::new(Mutex::new(highlight)),
         pv_key: 0,
         pv_slot: none(),
         view_proj,

crates/ui-egui/src/viewport.rs

@@ -305,6 +305,51 @@ fn sketch_color(t: &Tokens, c: Rgb) -> [u8; 4] {
     }
 }
 
+/// With a component active, bodies outside it are drawn faded.
+fn body_faded(app: &SolveApp, b: &solidcraft_engine::doc::ModelBody, active_bodies: &Option<Vec<String>>) -> bool {
+    let s = &app.session;
+    if let Some(names) = active_bodies {
+        !names.contains(&b.name)
+    } else {
+        s.active_component != 0 && !s.doc.component_within(b.component, s.active_component)
+    }
+}
+
+/// The display plan the model scene draws `b` from on the GPU, or `None` when its faces are built
+/// here from the reference mesh. Opaque shaded bodies (`opaque`: not faded, and no see-through
+/// body or face) are drawn from their plans where the device evaluates them; faded and
+/// see-through ones (sorted and blended per triangle) are built here.
+fn drawn_plan(
+    app: &SolveApp,
+    b: &solidcraft_engine::doc::ModelBody,
+    opaque: bool,
+) -> Option<std::sync::Arc<solidcraft_engine::geom::tess::TessPlan>> {
+    let shaded = app.ui.visual_style != 2;
+    (shaded && opaque && app.viewport.gpu.is_some_and(|g| g.tess))
+        .then(|| b.display_plan())
+        .flatten()
+        .filter(|plan| plan.vertex_count() <= solidcraft_gpu::MAX_BATCH)
+}
+
+/// The mesh whose triangles the model scene shows for `b`'s faces: the plan's (evaluated on the
+/// CPU, the same triangles the GPU draws) or the reference mesh. A face highlighted over the model
+/// must use the same triangles: another triangulation of the same face fights it at equal depth
+/// and shimmers.
+fn shown_mesh(
+    app: &SolveApp,
+    b: &solidcraft_engine::doc::ModelBody,
+    active_bodies: &Option<Vec<String>>,
+) -> std::sync::Arc<solidcraft_engine::geom::Mesh> {
+    let doc = &app.session.doc;
+    let opaque =
+        || !body_faded(app, b, active_bodies) && doc.model_body_opacity(b) >= 0.999 && doc.face_colors(b).iter().all(|(_, l)| l.opacity >= 0.999);
+    (app.viewport.gpu.is_some_and(|g| g.tess) && opaque())
+        .then(|| drawn_plan(app, b, true))
+        .flatten()
+        .and_then(|_| b.plan_mesh())
+        .unwrap_or_else(|| b.mesh())
+}
+
 fn build_scene(app: &SolveApp) -> GpuScene {
     let tk = Tokens::get();
     let mut sc = GpuScene::default();
@@ -316,12 +361,7 @@ fn build_scene(app: &SolveApp) -> GpuScene {
         if !app.session.body_visible(&b.name) || app.preview.replaced.contains(&b.name) {
             continue;
         }
-        // With a component active, bodies outside it are drawn faded.
-        let faded = if let Some(names) = &active_bodies {
-            !names.contains(&b.name)
-        } else {
-            s.active_component != 0 && !s.doc.component_within(b.component, s.active_component)
-        };
+        let faded = body_faded(app, b, &active_bodies);
         let base = body_color(&s.doc, b).unwrap_or(tk.body);
         let col = if faded { c4(base.gamma_multiply(0.35)) } else { c4(base) };
         let m = b.mesh();
@@ -338,13 +378,8 @@ fn build_scene(app: &SolveApp) -> GpuScene {
         let opacity = s.doc.model_body_opacity(b);
         let face_looks: std::collections::HashMap<usize, solidcraft_engine::doc::appearance::Look> =
             if shaded { s.doc.face_colors(b).into_iter().collect() } else { Default::default() };
-        // Opaque bodies are drawn from their display plans on the GPU where the device evaluates
-        // them; faded and see-through ones (sorted and blended per triangle) are built here.
         let opaque = !faded && opacity >= 0.999 && face_looks.values().all(|l| l.opacity >= 0.999);
-        let planned = (shaded && opaque && app.viewport.gpu.is_some_and(|g| g.tess))
-            .then(|| b.display_plan())
-            .flatten()
-            .filter(|plan| plan.vertex_count() <= solidcraft_gpu::MAX_BATCH);
+        let planned = drawn_plan(app, b, opaque);
         let on_gpu = planned.is_some();
         if let Some(plan) = planned {
             let face_colors = (0..b.body.face_count())
@@ -1400,6 +1435,10 @@ fn highlight_key(app: &SolveApp) -> u64 {
     format!("{:?}", app.viewport.hover).hash(&mut h);
     app.viewport.hover_feature.hash(&mut h);
     app.viewport.hover_bodies.hash(&mut h);
+    // Which mesh lights a face follows how the model draws its body (see `shown_mesh`).
+    app.session.active_component.hash(&mut h);
+    app.session.active_occurrence_path.hash(&mut h);
+    app.ui.visual_style.hash(&mut h);
     app.ui.pick_bodies.hash(&mut h);
     crate::theme::is_dark().hash(&mut h);
     app.preview.replaced.hash(&mut h);
@@ -1418,6 +1457,10 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
     let t = Tokens::get();
     let mut sc = GpuScene::default();
     let st = app.session.world_state();
+    let s = &app.session;
+    let active_bodies = s.active_occurrence_path.as_ref().and_then(|p| s.occurrence_bodies(p, true).ok());
+    // Faces lit over the model use the triangles the model shows them with.
+    let shown = |b: &solidcraft_engine::doc::ModelBody| shown_mesh(app, b, &active_bodies);
     let mut sel = Vec::new();
     for item in app.highlighted() {
         for piece in displayed_selection(app, &item) {
@@ -1491,7 +1534,7 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
             Sel::Face { body, index, .. } => {
                 // A live preview stands in for the body: its old faces would fight with it.
                 if let Some(b) = st.body(body).filter(|_| app.session.body_visible(body) && !app.preview.replaced.contains(body)) {
-                    face_tris(&mut sc, &b.mesh(), *index, c4(t.sel_face), false);
+                    face_tris(&mut sc, &shown(b), *index, c4(t.sel_face), false);
                     hidden_face_outline(&mut sc, &b.mesh(), *index, t.sel_edge);
                 }
             }
@@ -1499,7 +1542,7 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
                 if let Some(b) = st.body(name)
                     && app.session.body_visible(name)
                 {
-                    let m = b.mesh();
+                    let m = shown(b);
                     for f in 0..b.body.face_count() {
                         face_tris(&mut sc, &m, f, c4(t.body_selected), true);
                     }
@@ -1573,11 +1616,12 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
         let k = t.hover_face_lift;
         let c = body_rgb(app, &b.name);
         let m = b.mesh();
+        let lit = shown(b);
         for &f in &faces {
             if !sel.iter().any(|x| {
                 matches!(x, Sel::Face { body, index, .. } if body == &b.name && *index == f) || matches!(x, Sel::Body { name } if name == &b.name)
             }) {
-                face_tris(&mut sc, &m, f, [c.0.saturating_add(k), c.1.saturating_add(k), c.2.saturating_add(k), 255], true);
+                face_tris(&mut sc, &lit, f, [c.0.saturating_add(k), c.1.saturating_add(k), c.2.saturating_add(k), 255], true);
             }
         }
         for (ei, e) in m.edges.iter().enumerate() {
@@ -1593,7 +1637,7 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
             if let Some(b) = st.body(body) {
                 let k = t.hover_face_lift;
                 let c = body_rgb(app, body);
-                let m = b.mesh();
+                let m = shown(b);
                 for f in 0..b.body.face_count() {
                     face_tris(&mut sc, &m, f, [c.0.saturating_add(k), c.1.saturating_add(k), c.2.saturating_add(k), 255], true);
                 }
@@ -1608,8 +1652,9 @@ fn build_highlight(app: &SolveApp) -> GpuScene {
                 let c = body_rgb(app, body);
                 let m = b.mesh();
                 let patch = m.surface_patch(*index);
+                let lit = shown(b);
                 for &f in &patch {
-                    face_tris(&mut sc, &m, f, [c.0.saturating_add(k), c.1.saturating_add(k), c.2.saturating_add(k), 255], true);
+                    face_tris(&mut sc, &lit, f, [c.0.saturating_add(k), c.1.saturating_add(k), c.2.saturating_add(k), 255], true);
                 }
                 for (e, _) in m.edge_faces.iter().enumerate().filter(|(_, fs)| fs.iter().any(|f| patch.contains(&(*f as usize)))) {
                     if let Some(p) = m.edges.get(e)
@@ -3271,8 +3316,8 @@ mod tests {
             let proj = projection(&app, rect);
             let e = proj.cam.eye().to_f32();
             let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }];
-            let planned = scene.planned.len();
-            (planned, scene.tris.len(), crate::gpu::render_offscreen(scene, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, tess))
+            let (planned, shader) = (scene.planned.len(), tess.then_some(solidcraft_gpu::SHADER));
+            (planned, scene.tris.len(), crate::gpu::render_offscreen(scene, None, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, shader))
         };
         let (planned, gpu_tris, gpu) = render(true);
         let (none, cpu_tris, cpu) = render(false);
@@ -3290,6 +3335,102 @@ mod tests {
         assert!(differ * 100 < drawn, "{differ} of {drawn} model pixels differ");
     }
 
+    /// A device that cannot build the tessellation pipeline (D3D12's FXC rejected the shader
+    /// at startup and the app panicked) still shows the model: the failure is caught, logged,
+    /// and bodies are built on the CPU and drawn. A shader the device rejects stands in for FXC.
+    /// Skipped without a GPU adapter.
+    #[cfg(not(target_arch = "wasm32"))]
+    #[test]
+    fn bodies_are_built_on_the_cpu_when_the_tessellation_pipeline_fails() {
+        const REJECTED: &str = "
+            @group(0) @binding(1) var<storage, read_write> surfaces: array<u32>;
+            @compute @workgroup_size(64)
+            fn main(@builtin(global_invocation_id) gid: vec3<u32>) { surfaces[gid.x] = 1u; }
+        ";
+        let Some(tess) = crate::gpu::offscreen_tessellates(REJECTED) else {
+            eprintln!("skipping: no GPU adapter");
+            return;
+        };
+        assert!(!tess, "the rejected shader built a pipeline");
+        let mut s = solidcraft_engine::Session::default();
+        s.run_script(&solidcraft_engine::sample::script()).unwrap();
+        let mut app = SolveApp::new(s, crate::Services::default());
+        app.viewport.no_pixels = true;
+        app.fit_view();
+        app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess });
+        const SIZE: u32 = 256;
+        let rect = Rect::from_min_size(Pos2::ZERO, vec2(SIZE as f32, SIZE as f32));
+        let scene = build_scene(&app);
+        assert!(scene.planned.is_empty(), "no body waits for the GPU evaluation");
+        assert!(!scene.tris.is_empty(), "the bodies are meshed on the CPU");
+        let proj = projection(&app, rect);
+        let e = proj.cam.eye().to_f32();
+        let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }];
+        let Some(img) = crate::gpu::render_offscreen(scene, None, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, Some(REJECTED)) else {
+            return;
+        };
+        let n = (SIZE * SIZE) as usize;
+        let drawn = img.chunks(4).filter(|p| p[..3] != [255, 255, 255]).count();
+        assert!(drawn > n / 20, "the model is drawn: {drawn} pixels");
+    }
+
+    /// A hovered face lights solidly over a body drawn from its display plan on the GPU: the
+    /// highlight uses the plan's triangles, not the reference mesh's other triangulation of the
+    /// same face, which fought the model at equal depth and shimmered in wedges. Every face of a
+    /// box with a pocket and a round boss, hovered, changes the same pixels as on the CPU path (where model and
+    /// highlight share the reference mesh). Skipped without a GPU adapter.
+    #[cfg(not(target_arch = "wasm32"))]
+    #[test]
+    fn a_hovered_face_lights_whole_over_a_body_drawn_from_its_plan() {
+        let mut s = solidcraft_engine::Session::default();
+        s.execute("solid.box", &json!({"length": 40, "width": 40, "height": 40, "corner": [0, 0, 0], "operation": "new"})).unwrap();
+        s.execute("solid.box", &json!({"length": 16, "width": 16, "height": 20, "corner": [12, 12, 20], "operation": "cut"})).unwrap();
+        s.execute("solid.cylinder", &json!({"base": [33, 33, 40], "diameter": 10, "height": 10, "operation": "join"})).unwrap();
+        let mut app = SolveApp::new(s, crate::Services::default());
+        app.viewport.no_pixels = true;
+        app.fit_view();
+        const SIZE: u32 = 256;
+        let rect = Rect::from_min_size(Pos2::ZERO, vec2(SIZE as f32, SIZE as f32));
+        let body = app.session.world_state().bodies[0].clone();
+        let faces = body.body.face_count();
+        let mut render = |tess: bool, hover: Option<usize>| {
+            app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess });
+            app.viewport.hover = hover.map(|index| Hit::Face { body: body.name.clone(), index, point: Vec3::ZERO });
+            let (scene, highlight) = (build_scene(&app), build_highlight(&app));
+            assert_eq!(!scene.planned.is_empty(), tess, "the box is drawn from its plan exactly on the GPU path");
+            let proj = projection(&app, rect);
+            let e = proj.cam.eye().to_f32();
+            let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }];
+            let shader = tess.then_some(solidcraft_gpu::SHADER);
+            crate::gpu::render_offscreen(scene, Some(highlight), proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, shader)
+        };
+        let lit = |plain: &[u8], hovered: &[u8]| -> Vec<bool> { plain.chunks(4).zip(hovered.chunks(4)).map(|(a, b)| a != b).collect() };
+        let (Some(gpu_plain), Some(cpu_plain)) = (render(true, None), render(false, None)) else {
+            eprintln!("skipping: no GPU adapter");
+            return;
+        };
+        let mut checked = 0;
+        for f in 0..faces {
+            let (Some(gpu), Some(cpu)) = (render(true, Some(f)), render(false, Some(f))) else { return };
+            let (on_gpu, on_cpu) = (lit(&gpu_plain, &gpu), lit(&cpu_plain, &cpu));
+            let area = on_cpu.iter().filter(|x| **x).count();
+            if area < 400 {
+                continue; // facing away or hidden
+            }
+            // Inside the face, two pixels in from its outline: a curved face's outline follows
+            // its own facets, which the plan and the reference mesh cut differently.
+            let n = SIZE as i32;
+            let inside = |i: usize| {
+                let (x, y) = (i as i32 % n, i as i32 / n);
+                (-2..=2).all(|dy| (-2..=2).all(|dx| on_cpu.get(((y + dy).clamp(0, n - 1) * n + (x + dx).clamp(0, n - 1)) as usize) == Some(&true)))
+            };
+            let missed = (0..on_cpu.len()).filter(|&i| inside(i) && !on_gpu[i]).count();
+            assert!(missed * 200 < area, "face {f}: {missed} of its {area} hovered pixels do not light (the highlight fights the model)");
+            checked += 1;
+        }
+        assert!(checked >= 3, "{checked} faces seen lit");
+    }
+
     #[test]
     fn construction_drawing_and_hover_overlays_preserve_visible_gaps() {
         fn length(shape: &Shape) -> f32 {