git.lucas.co / cce-designer
graphic design tool
git clone https://git.lucas.co/cce-designer.git

commit4bcecc01f2c66e20c652513619cbd08c46332d89
parent035a33f5e0
authorLucas Galante <[email protected]>
date2026-09-24 11:22
feat: retire OpenCL and both kernel backends

The second half of Phase 7 step 3 (shapeshifter.md). With Sphere, Box,
Plane and Extrude native, nothing shipped used a kernel, and what the
runtime cost was out of proportion to a serial C interpreter: two
backends held to each other by a cross-check, a JIT compile and a
synchronous buffer round trip per evaluation, and an ICD that closed
file descriptors it did not own — the flake that made CCE_KERNEL_CPU=1
the only reliable way to run the suite.

Gone: the opencl node and its template, kernel_cpu.rs (1.8k lines), the
launcher and preprocessor in geometry.rs (~940 lines), the Update
Parameters button, the opencl3 dependency, CCE_KERNEL_CPU, the soup weld
that only the generator path used, and sixteen tests that ran kernels.

Kept: an `opencl` node in an old save passes its input through and
reports "OpenCL nodes are retired; rewrite the kernel as a wrangle"
through the error slot — visible on the status line and in the CLI
warning, never silently dropped. The type stays in is_geometry_node_type
for that one arm, and a_retired_opencl_node_passes_its_input_through_and_says_so
holds it to that.

The suite runs with no GPU, no OpenCL and no environment variable: 315.
CLAUDE.md's ICD section is replaced by a record of what left and why,
with the commit to read if the pattern recurs under another driver.

Co-Authored-By: Claude Fable 5.1 <[email protected]>

 CLAUDE.md         |  129 ++--
 Cargo.lock        |   31 -
 Cargo.toml        |    1 -
 nodes/opencl.json |   15 -
 shapeshifter.md   |   18 +-
 src/app.rs        |   34 -
 src/command.rs    |    1 -
 src/detail.rs     |    2 +-
 src/export_cli.rs |    2 +-
 src/geometry.rs   | 1355 +--------------------------------------
 src/kernel_cpu.rs | 1809 -----------------------------------------------------
 src/main.rs       |  101 ++-
 src/render.rs     |   21 +-
 src/thumbnail.rs  |    6 +-
 src/wrangle.rs    |    7 +-
 15 files changed, 132 insertions(+), 3400 deletions(-)

diff --git a/CLAUDE.md b/CLAUDE.md
index e8233bf..5923c79 100644
--- a/CLAUDE.md
+++ b/CLAUDE.md
@@ -5,8 +5,8 @@ This file provides guidance to Claude Code (claude.ai/code) when working with co
 ## What this is
 
 `cce-designer` is a node-based procedural 3D design app (Houdini-style) for the cce
-desktop environment: a node graph is evaluated into geometry — OpenCL kernels do the
-generation — and displayed in a 3D viewport with both a raster pass and a path-traced
+desktop environment: a node graph is evaluated into geometry — native Rust
+operators, plus a Rhai wrangle for per-element scripting — and displayed in a 3D viewport with both a raster pass and a path-traced
 (RT) preview mode.
 
 This crate is one member of the multi-repo `cce` Cargo workspace; workspace-wide rules
@@ -27,11 +27,11 @@ Two binaries: `cce-designer` (the app) and `vk-smoke` (`src/vk_smoke.rs`) — a
 standalone renderer smoke test that opens its own window; run it inside a Wayland
 session with `cargo run -p cce-designer --bin vk-smoke`.
 
-Some tests (e.g. `test_sphere_subnet_geometry_generation`) execute real OpenCL
-kernels THROUGH the OpenCL runtime when one exists; without one they exercise
-the CPU reference backend via the automatic fallback, so the suite is green
-headless. `kernel_cpu`'s template tests always run the CPU side; the
-cross-validation test compares backends and skips silently with no platform.
+The suite is pure CPU and needs no GPU, no OpenCL and no Wayland. (Until
+2026-09-24 it ran node kernels through the OpenCL runtime when one existed,
+with a CPU interpreter as the headless fallback, and `CCE_KERNEL_CPU=1` was
+the reliable way to run it — see "OpenCL is retired" below for why that is
+gone.)
 
 ### CLI modes
 
@@ -139,18 +139,6 @@ gone from cce-ui with the wgpu path).
   order over the widget slots, circular-pane clipping via `PaintItem::clip_circle`,
   network fade via text alpha. Rebuilt every drawn frame; the engine tessellates,
   shapes, and draws it.
-- `src/kernel_cpu.rs` — the CPU reference backend for node kernels: a
-  tree-walking interpreter for the C subset the kernels use (scalars, arrays,
-  user functions with pointer params, casts, the positional buffer ABI),
-  running the exact launcher contract of `run_opencl_kernel_with_params`. It is
-  the SEMANTIC REFERENCE — `cpu_matches_opencl_on_every_shipped_kernel`
-  compares the two backends vertex-by-vertex on every template kernel when a
-  platform exists, and the absolute template tests keep kernel coverage green
-  headless. Selected automatically when there is no OpenCL platform (one
-  stderr note), or forced with `CCE_KERNEL_CPU=1`. A kernel that fails ON a
-  present platform does not fall back — the error is in the kernel, and the
-  GPU diagnostics should surface. Step-budgeted so a non-terminating kernel is
-  an error, not a UI freeze. No vector types / barriers / local memory.
 - `src/geometry.rs` — node-graph evaluation. Every evaluator threads an
   `EvalSim` (current frame + `SimCache` + feedback stack) alongside the error
   slot. The `simnet` node type iterates: the chain between its `input` and
@@ -173,11 +161,9 @@ gone from cce-ui with the wgpu path).
   displayed level `input`/`output` children draw their resolved geometry (top
   level of the walk only, so outer views don't draw subnet chains twice).
   Frame changes invalidate the scene only when the
-  graph `contains_simnet`. Each OpenCL node's kernel code is
-  preprocessed: `chf("name", default)` / `chi` / `chv` calls are parsed into dynamic
-  UI parameters (`parse_dynamic_params`) and rewritten to `param_values[i]` reads
-  (`preprocess_opencl_code`). `network_sphere_vertices_with_errors` walks the graph
-  from output nodes; OpenCL failures are collected, not fatal.
+  graph `contains_simnet`. `network_sphere_vertices_with_errors` walks the
+  graph from output nodes; node failures are collected into one error slot
+  (still named `ocl_error` from the days it held OpenCL's), not fatal.
 - `src/viewport_3d.rs` — app-owned `Viewport3D` widget (camera orbit/zoom, inertial
   scroll, `rt_mode` flag switching the pane to the `cce_ui::vk` compute path tracer).
 - `src/viewer_state.rs` — the **viewer-state framework**: interactive viewport
@@ -414,54 +400,34 @@ off both by cce-ui default and in practice. Config the suite still reads is
 cosmetic in the same way (colors, fonts, plate radii), and no test asserts on
 it; the empty-`$XDG_CONFIG_HOME` run is how to check that claim again.
 
-### The OpenCL ICD corrupts unrelated file descriptors
-
-**Mesa's Rusticl ICD closes a file descriptor it does not own**, somewhere
-under `clGetPlatformIDs`. Caught under `strace -k` on 2026-09-23, the stack
-reading `__close` <- `libRusticlOpenCL.so` <- `clIcdGetPlatformIDsKHR` <-
-`clGetPlatformIDs`. By the time it closes, that descriptor number has been
-recycled to whatever another thread opened a moment earlier, so that thread's
-next `read` comes back **EBADF on a file nothing is wrong with**.
-
-It surfaced as a flake with no apparent connection to any of this: roughly one
-`cargo test` run in eight failed somewhere unrelated, most often
-`test_every_template_names_a_type_something_resolves` reporting "load_fs_tree
-returned 50 of 51 templates". The victim is whichever file lost the race —
-`sphere.json`, `output.json`, `hull.json`, a different one each time — and the
-tests that then failed were simply the ones that needed it.
-
-Three things follow, and the order they were tried in is worth keeping,
-because two of the three plausible fixes did nothing:
-
-- **Probing less often does NOT help.** Four tests each called
-  `get_platforms()` to decide whether to skip; caching the answer
-  (`has_opencl_platform`) and serializing every enumeration behind one mutex
-  (`probe_platforms`) took a process from dozens of overlapping enumerations
-  to two that cannot overlap — and the failure rate did not move (9 in 80,
-  against 10 in 60 before). The dangerous window is opening the ICD at all,
-  once per process, not how many times we ask afterwards. Both are kept
-  anyway: they are right on their own terms and cost nothing.
-- **Forcing the CPU backend did not help either, until it actually meant it.**
-  `CCE_KERNEL_CPU=1` selected the CPU kernel path but everything still
-  *probed*, and `cpu_matches_opencl_on_every_shipped_kernel` called straight
-  into `run_opencl_kernel_with_params`. That function now refuses at the top
-  when the backend is forced — one gate, reported as the no-platform error
-  every caller already handles — so forced-CPU never loads the ICD.
-- **What actually works is not loading the ICD.** `CCE_KERNEL_CPU=1 cargo
-  test` is the reliable way to run the suite, and `OCL_ICD_VENDORS=<empty
-  dir>` is the sharper instrument for confirming the ICD is the cause: with no
-  vendor to load, the EBADF failures disappear outright.
-
-The bug is in the ICD and cannot be fixed from here. What can be fixed is
-never being silent about the damage: `load_fs_tree` used to drop a template it
-could not read or parse inside an `if let Ok` pair, so the node just left the
-palette — which looks nothing like an I/O error and nothing like a parse error
-either. It says which file and why now, on stderr. That one change is what
-turned an afternoon of bisecting into a diagnosis.
-
-The app is far less exposed than the suite: it reads its templates at startup,
-on one thread, before OpenCL is in play. The suite is exposed because libtest
-runs its tests in parallel.
+### OpenCL is retired (2026-09-24)
+
+There is no OpenCL in this crate any more: no `opencl` node, no
+`kernel_cpu.rs`, no launcher, no `opencl3` dependency, no
+`CCE_KERNEL_CPU`. Phase 7 of `shapeshifter.md` is where the decision is
+argued; the short form is that the only scripting surface was a C subset
+carried by two backends that had to agree, every shipped kernel was serial
+(`if (id == 0)`), and the four templates that used them are native nodes
+now. Per-element scripting is the `wrangle` node (Rhai, CPU). GPU
+parallelism, when a solver needs it, comes back as WGSL compute through
+cce-ui's renderer — step 4 of the same phase — not as OpenCL.
+
+**An `opencl` node in an old save is not dropped.** `retired_opencl_node`
+passes its input through and reports `<name>: OpenCL nodes are retired;
+rewrite the kernel as a wrangle` through the error slot, so the status line
+says what happened and the fix is one rewrite. The type stays in
+`is_geometry_node_type` for exactly that arm.
+
+**What left with it, for the record.** Mesa's Rusticl ICD closed a file
+descriptor it did not own under `clGetPlatformIDs` (caught under `strace
+-k` on 2026-09-23), which had the suite failing one run in eight on
+whichever template file lost the race, and made `CCE_KERNEL_CPU=1` the only
+reliable way to run it. Probing less often did not help; forcing the CPU
+backend did not help until it also stopped loading the ICD; what worked was
+not loading it. The retirement is the final form of that fix. The
+diagnosis is in the git history of this section (commit `8fd0c29`) if the
+pattern ever recurs with another driver: a `read` returning EBADF on a file
+nothing is wrong with, in a process that has loaded a vendor ICD.
 
 ### Conditional parameter rows
 
@@ -687,10 +653,9 @@ name, position, flag and values stay, the children go, and a parameter the
 native template lacks goes with them. Only when the `opencl` child is
 actually there, so a subnet someone built by hand and called "sphere2"
 keeps what is inside it. The bundled `default_project.json` and
-`project.json` were converted in place. The `opencl` node itself still
-ships and still runs; retiring it and both kernel backends is the rest of
-step 3, and `test_loader_merges_new_template_params` is the migration's
-test.
+`project.json` were converted in place, and
+`test_loader_merges_new_template_params` is the migration's test. The
+`opencl` node and both kernel backends were retired the same day (above).
 
 ### The Embryo node is a template of nodes
 
@@ -750,7 +715,7 @@ itself would fail rather than recurse forever.
 
 `src/wrangle.rs` is a script run once per element, on Rhai — Phase 7 step 1
 of `shapeshifter.md`, and the app's scripting surface for per-element work
-where the `opencl` node used to be the only one. The engine is a dependency;
+where the retired `opencl` node used to be the only one. The engine is a dependency;
 what the module owns is the BINDING to the `Detail`, and it is VEX-shaped on
 purpose so `@P.y += sin(@P.x) * 0.1;` reads as it does there.
 
@@ -793,13 +758,13 @@ build geometry from no input at all — a wrangle with nothing wired still runs.
 Ints and floats mix (`@P.y * 2` works), which Rhai does not do on its own;
 the mixed arithmetic and comparison operators are registered by hand, as are
 `vec3`'s. Two budgets: `OPS_PER_ELEMENT` operations per element, which is
-`kernel_cpu`'s step budget as a setting rather than a hand-rolled counter,
+the retired `kernel_cpu`'s step budget as a setting rather than a hand-rolled counter,
 and `RUN_BUDGET` seconds of wall clock for the whole run, checked in
 `on_progress` every few thousand operations. Any failure — syntax, a runtime
 error on an element, a budget — fails the WHOLE run, named by node and
 element (`wrangle1: point 4: …`), and the input passes through untouched: a
 half-wrangled geometry is not a result. Compiled scripts cache by desugared
-source in a thread-local, as `OPENCL_CACHE` keys kernels.
+source in a thread-local, as the retired launcher cached kernels.
 
 CPU only, deliberately: an interpreter is an order of magnitude or more
 below native Rust, which is fine for tens of thousands of elements per edit
@@ -1579,10 +1544,10 @@ where the template puts them (after the last template param the instance
 already has — so the Sphere's Method lands above Radius in an old save, not
 below Color), existing ones keep their value but take the template's UI
 metadata, and a subnet template (the Embryo, since the four kernel subnets
-went native) refreshes its children's params and any kernel child's `Code`
+went native) refreshes its children's params and any child's `Code`
 outright — **the template owns the surface and implementation, the
-instance owns its values.** A kernel hand-edited inside a template
-instance reverts on load; custom kernels belong in bare OpenCL nodes,
+instance owns its values.** A script hand-edited inside a template
+instance reverts on load; custom scripts belong in bare wrangle nodes,
 which the merge never touches.
 Native nodes match their template by type, subnet instances by name
 ("sphere3" → "Sphere", case-insensitively) plus a full child name/type match; the merge never
diff --git a/Cargo.lock b/Cargo.lock
index c8eae52..df82ddc 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -423,7 +423,6 @@ dependencies = [
  "env_logger",
  "glam",
  "log",
- "opencl3",
  "png",
  "rhai",
  "serde",
@@ -479,17 +478,6 @@ version = "0.2.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "f079e83a288787bcd14a6aea84cee5c87a67c5a3e660c30f557a3d24761b3527"
 
-[[package]]
-name = "cl3"
-version = "0.9.5"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b823f24e72fa0c68aa14a250ae1c0848e68d4ae188b71c3972343e45b46f8644"
-dependencies = [
- "libc",
- "opencl-sys",
- "thiserror 1.0.69",
-]
-
 [[package]]
 name = "codespan-reporting"
 version = "0.11.1"
@@ -1472,25 +1460,6 @@ version = "1.70.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "384b8ab6d37215f3c5301a95a4accb5d64aa607f1fcb26a11b5303878451b4fe"
 
-[[package]]
-name = "opencl-sys"
-version = "0.2.9"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "de15dd01496ae90c5799f5266184ab020082b4065800ff0b732f489371d0e5cf"
-dependencies = [
- "libc",
-]
-
-[[package]]
-name = "opencl3"
-version = "0.9.5"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "26ab4a90cb496f787d3934deb0c54fa9d65e7bed710c10071234aab0196fba04"
-dependencies = [
- "cl3",
- "libc",
-]
-
 [[package]]
 name = "ordered-stream"
 version = "0.2.0"
diff --git a/Cargo.toml b/Cargo.toml
index 508a1c6..e25adf8 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -18,7 +18,6 @@ glam = "0.29"
 serde = { version = "1", features = ["derive"] }
 serde_json = "1"
 toml = "0.8"
-opencl3 = "0.9"
 log = "0.4"
 env_logger = "0.11"
 png = "0.17"
diff --git a/nodes/opencl.json b/nodes/opencl.json
deleted file mode 100644
index 8d913f8..0000000
--- a/nodes/opencl.json
+++ /dev/null
@@ -1,15 +0,0 @@
-{
-  "name": "OpenCL",
-  "type": "opencl",
-  "inputs": 1,
-  "outputs": 1,
-  "params": [
-    { "name": "Input", "default": "", "type": "text" },
-    {
-      "name": "Code",
-      "type": "code",
-      "default": "__kernel void process(__global float* pos, __global float* col, int count) {\n    int id = get_global_id(0);\n    if (id < count) {\n        // Default deformation kernel (sine wave deforming Y based on X)\n        pos[id * 3 + 1] += sin(pos[id * 3] * 4.0f) * 0.15f;\n    }\n}"
-    },
-    { "name": "Update Parameters", "type": "button", "default": "" }
-  ]
-}
diff --git a/shapeshifter.md b/shapeshifter.md
index 0733ae1..6f89b1e 100644
--- a/shapeshifter.md
+++ b/shapeshifter.md
@@ -660,14 +660,16 @@ step 1 resolves through `expr.rs`'s scope, so a referenced parameter that is
 itself an expression evaluates before the wrangle sees it, and neither
 language has to know the other exists.
 
-> **Started (2026-09-24): the four are native.** `src/shapes.rs` — Sphere
-> (all three methods, Cube as quads), Box (with a Center, without its dead
-> Input), Plane, and Extrude as a WHOLE (walls on boundary edges only,
-> where the kernel walled every interior edge). Saved kernel subnets
-> migrate on load through `nativize_kernel_subnets`; the bundled project
-> files were converted in place. The `opencl` node, `kernel_cpu.rs`, the
-> launcher and `opencl3` still ship — the retirement is the half of this
-> step still to do.
+> **Landed (2026-09-24).** `src/shapes.rs` — Sphere (all three methods,
+> Cube as quads), Box (with a Center, without its dead Input), Plane, and
+> Extrude as a WHOLE (walls on boundary edges only, where the kernel walled
+> every interior edge). Saved kernel subnets migrate on load through
+> `nativize_kernel_subnets`; the bundled project files were converted in
+> place. Then the retirement: the `opencl` node, `kernel_cpu.rs`, the
+> launcher and preprocessor, `opencl3`, `CCE_KERNEL_CPU` and the ICD
+> hazard are gone. An `opencl` node in an old save passes its input through
+> and reports itself. The suite runs with no GPU, no OpenCL and no
+> environment variable.
 
 **Step 3 — port the four kernel templates native, then retire OpenCL.**
 Sphere, Box, Plane and Extrude are the only kernels that ship. A native
diff --git a/src/app.rs b/src/app.rs
index 04e0629..a9daff5 100644
--- a/src/app.rs
+++ b/src/app.rs
@@ -3123,40 +3123,6 @@ impl State {
             "Export" => {
                 self.run_export();
             }
-            "Update Parameters" => {
-                if let Some(slot_idx) = self.graph().selected_node() {
-                    let dir = self.current_dir_mut();
-                    if let Some(child) = dir.children.get_mut(slot_idx) {
-                        if child.node_type == "opencl" {
-                            let code_val = child.params.iter()
-                                .find(|p| p.name == "Code")
-                                .map(|p| p.default.clone())
-                                .unwrap_or_default();
-                            let parsed_params = crate::geometry::parse_dynamic_params(&code_val);
-                            let mut new_params = Vec::new();
-                            for base_name in &["Input", "Code", "Update Parameters"] {
-                                if let Some(p) = child.params.iter().find(|p| p.name == *base_name) {
-                                    new_params.push(p.clone());
-                                }
-                            }
-                            for mut parsed in parsed_params {
-                                if let Some(existing) = child.params.iter().find(|p| p.name == parsed.name) {
-                                    parsed.default = existing.default.clone();
-                                }
-                                new_params.push(parsed);
-                            }
-                            child.params = new_params;
-                            let updated_disp = param_display(&child.params);
-                            self.param_mut().set_display_params(&updated_disp);
-                            self.rebuild_scene_geometry();
-                            self.sync_nodes();
-                        }
-                    }
-                }
-            }
-            "New Project" | "New" => {
-                self.new_project();
-            }
             "Set As Default" => {
                 self.set_current_as_default();
             }
diff --git a/src/command.rs b/src/command.rs
index 0ca75c9..12c373e 100644
--- a/src/command.rs
+++ b/src/command.rs
@@ -203,7 +203,6 @@ pub const COMMANDS: &[Command] = &[
 
     // --- Parameters ---
     Command { id: "export", label: "Export", context: Context::Parameters, run: Run::Menu("Export"), default_chord: None },
-    Command { id: "update_parameters", label: "Update Parameters", context: Context::Parameters, run: Run::Menu("Update Parameters"), default_chord: None },
 
     // --- Playbar ---
     Command { id: "play_pause", label: "Play / Pause", context: Context::Playbar, run: Run::Key(Action::PlayPause), default_chord: Some("Up") },
diff --git a/src/detail.rs b/src/detail.rs
index 5b3ab99..df2f1bf 100644
--- a/src/detail.rs
+++ b/src/detail.rs
@@ -1738,7 +1738,7 @@ impl Detail {
     /// The point behind every corner [`Detail::triangulate`] emits, in the
     /// same order.
     ///
-    /// Lets a caller that had to flatten to triangles — the OpenCL launcher,
+    /// Lets a caller that had to flatten to triangles — the OpenCL launcher once,
     /// until the Phase 1 ABI binds attributes directly — put results back on
     /// the points they came from instead of welding the output and losing
     /// every identity.
diff --git a/src/export_cli.rs b/src/export_cli.rs
index b27c5a9..7a0ba11 100644
--- a/src/export_cli.rs
+++ b/src/export_cli.rs
@@ -90,7 +90,7 @@ pub fn run(
     if let Some(e) = ocl_error {
         // Non-fatal, like the thumbnail: the rest of the scene still exports,
         // and a silent partial file would be worse than a warning.
-        eprintln!("cce-designer --export: OpenCL error (geometry partially skipped): {e}");
+        eprintln!("cce-designer --export: node error (geometry partially skipped): {e}");
     }
     if geom.num_prims() == 0 {
         return Err("the geometry has no primitives".to_string());
diff --git a/src/geometry.rs b/src/geometry.rs
index 37cdd51..ec8fbdd 100644
--- a/src/geometry.rs
+++ b/src/geometry.rs
@@ -1,13 +1,5 @@
 use crate::app::{FsNode, ParamDef};
 use std::collections::HashMap;
-use opencl3::platform::get_platforms;
-use opencl3::device::{Device, CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_CPU};
-use opencl3::context::Context;
-use opencl3::command_queue::CommandQueue;
-use opencl3::program::Program;
-use opencl3::kernel::{Kernel, ExecuteKernel};
-use opencl3::memory::{Buffer as ClBuffer, CL_MEM_READ_WRITE};
-use opencl3::types::{cl_float, cl_int, CL_TRUE};
 use glam::Vec3;
 use crate::detail::{AttribData, AttribValue, Detail, CD};
 
@@ -180,60 +172,6 @@ pub fn detail_to_soup(d: &Detail) -> Geometry {
     Geometry { vertices }
 }
 
-/// Weld a soup back into a [`Detail`], carrying its per-corner attributes onto
-/// the points they welded into (first corner wins).
-///
-/// The inverse of [`detail_to_soup`], and the other half of the migration
-/// bridge. Attribute transfer is not incidental: the kernel launcher stamps a
-/// default `Norm` and `UV` on every vertex it generates, and a weld that only
-/// took positions and colors would quietly drop them — which is exactly what
-/// the parameter pane's attribute picker reads.
-pub fn soup_to_detail(soup: &Geometry) -> Detail {
-    let positions: Vec<[f32; 3]> = soup.vertices.iter().map(|v| v.pos).collect();
-    let colors: Vec<[f32; 3]> = soup.vertices.iter().map(|v| v.col).collect();
-    let (mut d, point_of) = Detail::from_triangle_soup_with_map(&positions, &colors);
-
-    let mut names: Vec<&str> = Vec::new();
-    for v in &soup.vertices {
-        for k in v.attributes.keys() {
-            if !names.contains(&k.as_str()) {
-                names.push(k);
-            }
-        }
-    }
-    names.sort_unstable();
-
-    for name in names {
-        let mut written = vec![false; d.num_points()];
-        let mut data: Option<AttribData> = None;
-        for (corner, v) in soup.vertices.iter().enumerate() {
-            let (Some(&p), Some(val)) = (point_of.get(corner), v.attributes.get(name)) else {
-                continue;
-            };
-            let p = p as usize;
-            if std::mem::replace(&mut written[p], true) {
-                continue;
-            }
-            let value = detail_attr(val);
-            let arr = data.get_or_insert_with(|| AttribData::zeroed(value.ty(), d.num_points()));
-            let _ = arr.set(p, value);
-        }
-        if let Some(arr) = data {
-            let _ = d.points_mut().insert(name, arr);
-        }
-    }
-    d
-}
-
-fn detail_attr(v: &GAttribute) -> AttribValue {
-    match *v {
-        GAttribute::Float(x) => AttribValue::Float(x),
-        GAttribute::Float2(x) => AttribValue::Float2(x),
-        GAttribute::Float3(x) => AttribValue::Float3(x),
-        GAttribute::Float4(x) => AttribValue::Float4(x),
-    }
-}
-
 fn soup_attr(v: AttribValue) -> GAttribute {
     match v {
         AttribValue::Float(x) => GAttribute::Float(x),
@@ -1201,7 +1139,7 @@ pub fn generate_single_node_geometry_with_errors(
     } else if target.node_type.eq_ignore_ascii_case("attribute") {
         resolve_attribute_geometry_with_errors(root, target, visited, ocl_error, sim)
     } else if target.node_type.eq_ignore_ascii_case("opencl") {
-        resolve_opencl_geometry_with_errors(root, target, visited, ocl_error, sim)
+        retired_opencl_node(root, target, visited, ocl_error, sim)
     } else if target.node_type.eq_ignore_ascii_case("simnet") {
         resolve_simnet_geometry_with_errors(root, target, visited, ocl_error, sim)
     } else if target.node_type.eq_ignore_ascii_case("node") {
@@ -2325,6 +2263,25 @@ pub fn resolve_extrude_geometry_with_errors(
     Some(crate::shapes::extrude_detail(&input, distance, keep_base))
 }
 
+/// An `opencl` node from a save older than 2026-09-24, when the node and
+/// the runtime behind it were retired (`shapeshifter.md`, Phase 7 step 3).
+/// It passes its input through unchanged and reports itself through the
+/// error slot — visible, not silently dropped — so the fix is one rewrite
+/// as a wrangle rather than a hunt for geometry that stopped appearing.
+pub fn retired_opencl_node(
+    root: &FsNode,
+    target: &FsNode,
+    visited: &mut Vec<String>,
+    ocl_error: &mut Option<String>,
+    sim: &mut EvalSim,
+) -> Option<Detail> {
+    if ocl_error.is_none() {
+        *ocl_error = Some(format!("{}: OpenCL nodes are retired; rewrite the kernel as a wrangle", target.name));
+    }
+    let input_node = find_input_node(root, target, &node_param_str(target, "Input", ""))?;
+    generate_single_node_geometry_with_errors(root, input_node, visited, ocl_error, sim)
+}
+
 /// The Switch node: one of up to four inputs, chosen by Index.
 ///
 /// What a composed subnet puts behind a choice: the Embryo's Source is a
@@ -4643,943 +4600,6 @@ pub fn resolve_scatter_geometry_with_errors(
     Some(res)
 }
 
-/// The point attributes a kernel names, in first-use order — one buffer each,
-/// bound after `param_values`.
-///
-/// A kernel reaches an attribute the same way it reaches a parameter: by name,
-/// through a call the preprocessor rewrites. `attrf("mass", i)` reads point
-/// `i`'s float attribute and `setattrf("mass", i, v)` writes it. The name is
-/// the declaration — an attribute the geometry does not carry is created,
-/// zeroed, so a kernel can produce one.
-pub fn parse_attr_refs(code: &str) -> Vec<String> {
-    let mut names: Vec<String> = Vec::new();
-    for call in ["attrf(", "setattrf("] {
-        let mut from = 0;
-        while let Some(pos) = code[from..].find(call) {
-            let at = from + pos;
-            // `setattrf(` also contains `attrf(`; only count the outer one.
-            let is_inner = call == "attrf(" && at >= 3 && &code[at - 3..at] == "set";
-            from = at + call.len();
-            if is_inner {
-                continue;
-            }
-            let Some(name) = quoted_arg(&code[from..]) else { continue };
-            if !names.contains(&name) {
-                names.push(name);
-            }
-        }
-    }
-    names
-}
-
-/// The first single- or double-quoted string in an argument list.
-fn quoted_arg(args: &str) -> Option<String> {
-    let q = args.find(|c| c == '"' || c == '\'')?;
-    let quote = args.as_bytes()[q] as char;
-    let rest = &args[q + 1..];
-    let end = rest.find(quote)?;
-    let name = &rest[..end];
-    (!name.is_empty()).then(|| name.to_string())
-}
-
-/// Append arguments to `process`'s parameter list, in binding order.
-fn rewrite_kernel_signature_with(code: &str, extra: &[String]) -> String {
-    if extra.is_empty() {
-        return code.to_string();
-    }
-    let bytes = code.as_bytes();
-    if let Some(process_idx) = code.find("process") {
-        let mut idx = process_idx + "process".len();
-        while idx < bytes.len() && (bytes[idx] as char).is_whitespace() {
-            idx += 1;
-        }
-        if idx < bytes.len() && bytes[idx] == b'(' {
-            let start_args = idx + 1;
-            let mut paren_count = 1;
-            let mut end_args = start_args;
-            while end_args < bytes.len() && paren_count > 0 {
-                if bytes[end_args] == b'(' {
-                    paren_count += 1;
-                } else if bytes[end_args] == b')' {
-                    paren_count -= 1;
-                }
-                end_args += 1;
-            }
-            if paren_count == 0 {
-                let closing_paren_idx = end_args - 1;
-                let before = &code[..closing_paren_idx];
-                let after = &code[closing_paren_idx..];
-                let args_str = code[start_args..closing_paren_idx].trim();
-                let sep = if args_str.is_empty() { "" } else { ", " };
-                return format!("{}{}{}{}", before, sep, extra.join(", "), after);
-            }
-        }
-    }
-    code.to_string()
-}
-
-pub fn preprocess_opencl_code(code: &str) -> String {
-    let parsed_params = parse_dynamic_params(code);
-    let attrs = parse_attr_refs(code);
-    if parsed_params.is_empty() && attrs.is_empty() {
-        return code.to_string();
-    }
-
-    let mut param_indices = std::collections::HashMap::new();
-    let mut flat_idx = 0;
-    for p in &parsed_params {
-        param_indices.insert(p.name.clone(), flat_idx);
-        if p.param_type == "float3" {
-            flat_idx += 3;
-        } else {
-            flat_idx += 1;
-        }
-    }
-
-    // Binding order, and therefore signature order: the fixed arguments, then
-    // `param_values` if the kernel names any parameter, then one buffer per
-    // attribute it names. Both launchers bind positionally against exactly
-    // this list.
-    let mut extra: Vec<String> = Vec::new();
-    if !parsed_params.is_empty() {
-        extra.push("__global const float* param_values".to_string());
-    }
-    for i in 0..attrs.len() {
-        extra.push(format!("__global float* attr_{}", i));
-    }
-    let mut processed = rewrite_kernel_signature_with(code, &extra);
-
-    // Attribute access rewrites before parameter ones: `attrf("mass", chi("k"))`
-    // is legal, and the parameter pass would otherwise rewrite inside a call
-    // this pass still needs to find by name.
-    for (slot, name) in attrs.iter().enumerate() {
-        processed = rewrite_attr_calls(&processed, name, slot);
-    }
-
-    let prefixes = [("chf", "slider"), ("chi", "spinbox"), ("chv", "float3"), ("chb", "toggle")];
-    for &(prefix, _) in &prefixes {
-        let pattern = format!("{}(", prefix);
-        while let Some(pos) = processed.find(&pattern) {
-            let start_idx = pos + pattern.len();
-            let mut paren_count = 1;
-            let mut end_pos = start_idx;
-            let bytes = processed.as_bytes();
-            while end_pos < bytes.len() && paren_count > 0 {
-                if bytes[end_pos] == b'(' {
-                    paren_count += 1;
-                } else if bytes[end_pos] == b')' {
-                    paren_count -= 1;
-                }
-                end_pos += 1;
-            }
-            if paren_count == 0 {
-                let full_match = &processed[pos..end_pos];
-                let args_str = &processed[start_idx..end_pos - 1];
-                let mut replacement = None;
-                if let Some(name) = quoted_arg(args_str) {
-                    if let Some(&flat_idx) = param_indices.get(&name) {
-                        match prefix {
-                            "chf" => {
-                                replacement = Some(format!("param_values[{}]", flat_idx));
-                            }
-                            "chi" | "chb" => {
-                                replacement = Some(format!("((int)param_values[{}])", flat_idx));
-                            }
-                            "chv" => {
-                                replacement = Some(format!(
-                                    "(float3)(param_values[{}], param_values[{}], param_values[{}])",
-                                    flat_idx, flat_idx + 1, flat_idx + 2
-                                ));
-                            }
-                            _ => {}
-                        }
-                    }
-                }
-                if let Some(rep) = replacement {
-                    processed = processed.replace(full_match, &rep);
-                } else {
-                    processed = processed.replace(full_match, "0");
-                }
-            } else {
-                break;
-            }
-        }
-    }
-    processed
-}
-
-/// Rewrite one attribute's reads and writes into indexing on its buffer.
-///
-/// `attrf("mass", E)` becomes `attr_N[E]` and `setattrf("mass", I, V)` becomes
-/// `attr_N[I] = (V)` — an assignment expression, so it reads as a statement
-/// where the kernel wrote one and still composes where it did not.
-fn rewrite_attr_calls(code: &str, name: &str, slot: usize) -> String {
-    let mut out = code.to_string();
-    for setter in [true, false] {
-        let call = if setter { "setattrf(" } else { "attrf(" };
-        let mut from = 0;
-        loop {
-            let Some(pos) = out[from..].find(call) else { break };
-            let at = from + pos;
-            if !setter && at >= 3 && &out[at - 3..at] == "set" {
-                from = at + call.len();
-                continue;
-            }
-            let args_start = at + call.len();
-            let Some(args_end) = matching_paren(&out, args_start) else { break };
-            let args = &out[args_start..args_end];
-            let Some(found) = quoted_arg(args) else {
-                from = args_end + 1;
-                continue;
-            };
-            if found != name {
-                from = args_end + 1;
-                continue;
-            }
-            // Arguments after the name, split at the top level so an index
-            // expression containing a comma inside parentheses stays whole.
-            let after_name = match args.find(',') {
-                Some(c) => &args[c + 1..],
-                None => "",
-            };
-            let parts = split_top_level(after_name);
-            let replacement = if setter {
-                match (parts.first(), parts.get(1)) {
-                    (Some(i), Some(v)) => format!("attr_{}[{}] = ({})", slot, i.trim(), v.trim()),
-                    _ => "0".to_string(),
-                }
-            } else {
-                match parts.first() {
-                    Some(i) => format!("attr_{}[{}]", slot, i.trim()),
-                    None => "0".to_string(),
-                }
-            };
-            out.replace_range(at..args_end + 1, &replacement);
-            from = at + replacement.len();
-        }
-    }
-    out
-}
-
-/// Index just past the `(` at `open`, of its matching `)`.
-fn matching_paren(s: &str, open: usize) -> Option<usize> {
-    let bytes = s.as_bytes();
-    let mut depth = 1;
-    let mut i = open;
-    while i < bytes.len() {
-        match bytes[i] {
-            b'(' => depth += 1,
-            b')' => {
-                depth -= 1;
-                if depth == 0 {
-                    return Some(i);
-                }
-            }
-            _ => {}
-        }
-        i += 1;
-    }
-    None
-}
-
-/// Split on commas that are not inside parentheses or brackets.
-fn split_top_level(s: &str) -> Vec<&str> {
-    let mut parts = Vec::new();
-    let (mut depth, mut start) = (0i32, 0usize);
-    for (i, c) in s.char_indices() {
-        match c {
-            '(' | '[' => depth += 1,
-            ')' | ']' => depth -= 1,
-            ',' if depth == 0 => {
-                parts.push(&s[start..i]);
-                start = i + 1;
-            }
-            _ => {}
-        }
-    }
-    if start <= s.len() {
-        parts.push(&s[start..]);
-    }
-    parts.retain(|p| !p.trim().is_empty());
-    parts
-}
-
-pub fn parse_dynamic_params(code: &str) -> Vec<ParamDef> {
-    let mut parsed = Vec::new();
-    let prefixes = [("chf", "slider"), ("chi", "spinbox"), ("chv", "float3"), ("chb", "toggle")];
-    for &(prefix, ptype) in &prefixes {
-        let pattern = format!("{}(", prefix);
-        let mut start_idx = 0;
-        while let Some(pos) = code[start_idx..].find(&pattern) {
-            let actual_pos = start_idx + pos;
-            start_idx = actual_pos + pattern.len();
-            let mut paren_count = 1;
-            let mut end_pos = start_idx;
-            let code_bytes = code.as_bytes();
-            while end_pos < code_bytes.len() && paren_count > 0 {
-                if code_bytes[end_pos] == b'(' {
-                    paren_count += 1;
-                } else if code_bytes[end_pos] == b')' {
-                    paren_count -= 1;
-                }
-                end_pos += 1;
-            }
-            if paren_count == 0 {
-                let args_str = &code[start_idx..end_pos - 1];
-                if let Some(first_quote_pos) = args_str.find(|c| c == '"' || c == '\'') {
-                    let quote_char = args_str.chars().nth(first_quote_pos).unwrap();
-                    if let Some(second_quote_pos) = args_str[first_quote_pos + 1..].find(quote_char) {
-                        let name = &args_str[first_quote_pos + 1..first_quote_pos + 1 + second_quote_pos];
-                        if !name.is_empty() {
-                            let mut default_val = match prefix {
-                                "chf" => "0.5".to_string(),
-                                "chi" => "0".to_string(),
-                                "chb" => "false".to_string(),
-                                "chv" => "0.00:0.00:0.00".to_string(),
-                                _ => "".to_string(),
-                            };
-                            let rest = &args_str[first_quote_pos + 1 + second_quote_pos + 1..];
-                            if let Some(comma_pos) = rest.find(',') {
-                                let val_part = rest[comma_pos + 1..].trim();
-                                if !val_part.is_empty() {
-                                    let mut clean_val = val_part.to_string();
-                                    if clean_val.ends_with('f') {
-                                        clean_val.pop();
-                                    }
-                                    if clean_val.ends_with("f32") {
-                                        clean_val.truncate(clean_val.len() - 3);
-                                    }
-                                    let clean_val = clean_val.trim();
-                                    if prefix == "chv" {
-                                        let parts: Vec<String> = val_part.split(',')
-                                            .map(|p| {
-                                                let mut s = p.trim().to_string();
-                                                if s.ends_with('f') { s.pop(); }
-                                                if s.ends_with("f32") { s.truncate(s.len() - 3); }
-                                                s.trim().to_string()
-                                            })
-                                            .collect();
-                                        if parts.len() >= 3 {
-                                            if let (Ok(x), Ok(y), Ok(z)) = (parts[0].parse::<f32>(), parts[1].parse::<f32>(), parts[2].parse::<f32>()) {
-                                                default_val = format!("{:.2}:{:.2}:{:.2}", x, y, z);
-                                            }
-                                        } else {
-                                            if let Ok(val) = clean_val.parse::<f32>() {
-                                                default_val = format!("{:.2}:{:.2}:{:.2}", val, val, val);
-                                            }
-                                        }
-                                    } else {
-                                        default_val = clean_val.to_string();
-                                    }
-                                }
-                            }
-                            if !parsed.iter().any(|p: &ParamDef| p.name == name) {
-                                let (min, max, step) = match prefix {
-                                    "chf" => (Some(0.0), Some(2.0), Some(0.01)),
-                                    "chi" => (Some(0.0), Some(1000.0), Some(1.0)),
-                                    _ => (None, None, None),
-                                };
-                                parsed.push(ParamDef {
-                                    name: name.to_string(),
-                                    label: String::new(),
-                                    param_type: ptype.to_string(),
-                                    default: default_val,
-                                    options: Vec::new(),
-                                    min,
-                                    max,
-                                    step,
-                                    show_when: String::new(), expr: false,
-                                });
-                            }
-                        }
-                    }
-                }
-            }
-        }
-    }
-    parsed
-}
-
-pub fn resolve_opencl_geometry_with_errors(
-    root: &FsNode,
-    target: &FsNode,
-    visited: &mut Vec<String>,
-    ocl_error: &mut Option<String>,
-    sim: &mut EvalSim,
-) -> Option<Detail> {
-    let input_name = node_param_str(target, "Input", "");
-    let mut input = if !input_name.is_empty() {
-        // Siblings first, exactly like the output type's lookup: subnet
-        // templates (Extrude) wire their inner opencl to a child named
-        // "input1", and a global-first search would resolve to the FIRST
-        // subnet's child once two instances exist.
-        let sibling = find_parent_node(root, &target.id)
-            .and_then(|p| p.children.iter().find(|c| c.name == input_name || c.id == input_name));
-        if let Some(input_node) = sibling.or_else(|| find_input_node(root, target, &input_name)) {
-            generate_single_node_geometry_with_errors(root, input_node, visited, ocl_error, sim)
-                .unwrap_or_default()
-        } else {
-            Detail::new()
-        }
-    } else {
-        Detail::new()
-    };
-
-    let code = node_param_str(target, "Code", "");
-    if !code.is_empty() {
-        let parsed_params = parse_dynamic_params(&code);
-        let mut flat_values = Vec::new();
-        // The kernel's own parameter, else the enclosing subnet's — the
-        // parent's value RESOLVED: a Sphere instance whose Radius is
-        // ch("Radius") hands its kernel the subnet's number, not the
-        // reference string (which parses as nothing and left the kernel at
-        // its default). The DEFINITION is looked up rather than the value
-        // string, because a choice is worth its option index and only the
-        // definition knows the options.
-        let resolved_parent = find_parent_node(root, &target.id)
-            .map(|parent| resolve_param_refs(root, parent, sim.frame, ocl_error).unwrap_or_else(|| parent.clone()));
-        let find_def = |name: &str| -> Option<&ParamDef> {
-            target.params.iter().find(|d| d.name.eq_ignore_ascii_case(name)).or_else(|| {
-                resolved_parent.as_ref().and_then(|parent| parent.params.iter().find(|d| d.name.eq_ignore_ascii_case(name)))
-            })
-        };
-        for p in &parsed_params {
-            let def = find_def(&p.name);
-            if p.param_type == "float3" {
-                let val_str = def.map_or(p.default.as_str(), |d| d.default.as_str());
-                let parts: Vec<&str> = val_str.split(':').collect();
-                let (x, y, z) = if parts.len() >= 3 {
-                    (parts[0].parse::<f32>().unwrap_or(0.0), parts[1].parse::<f32>().unwrap_or(0.0), parts[2].parse::<f32>().unwrap_or(0.0))
-                } else {
-                    (0.0, 0.0, 0.0)
-                };
-                flat_values.push(x);
-                flat_values.push(y);
-                flat_values.push(z);
-            } else {
-                flat_values.push(def.map_or_else(|| number_of_str(&p.default), param_number));
-            }
-        }
-
-        let processed_code = preprocess_opencl_code(&code);
-        let attr_names = parse_attr_refs(&code);
-        let is_generator = code.contains("out_count");
-
-        // A DEFORMER runs over POINTS. It never sees a triangle, so topology,
-        // groups and point identities pass through untouched — the flatten,
-        // weld and "first corner wins" reconciliation this used to need are
-        // all gone. Inside a simnet that is the difference between a solver
-        // that can follow a point across frames and one that cannot.
-        if !is_generator {
-            let result = run_kernel_on_detail(&processed_code, &attr_names, &mut input, &flat_values);
-            if let Err(e) = result {
-                if ocl_error.is_none() {
-                    *ocl_error = Some(e);
-                }
-            }
-            return Some(input);
-        }
-
-        // A GENERATOR builds a new corner list, so it still speaks soup and
-        // its output welds into fresh geometry with fresh identities. Widening
-        // the generator ABI to emit points and primitives directly is the
-        // piece of Phase 1 still outstanding.
-        let mut geom = detail_to_soup(&input);
-        // CPU reference backend (kernel_cpu): forced via CCE_KERNEL_CPU=1, and
-        // the automatic fallback when there is no OpenCL platform at all — the
-        // state this machine reached silently when nvidia-open fell out of
-        // kernel lockstep, which used to mean every kernel node produced
-        // empty geometry. A kernel that FAILS on a present platform (compile
-        // error, bad code) does NOT fall back: the two backends share the
-        // language, so the error is almost certainly in the kernel, and
-        // hiding the GPU diagnostics behind a second attempt would obscure it.
-        let result = if crate::kernel_cpu::forced() {
-            crate::kernel_cpu::run_kernel_cpu(&processed_code, &mut geom, &flat_values)
-        } else {
-            match run_opencl_kernel_with_params(&processed_code, &mut geom, &flat_values) {
-                Err(e) if e.contains("No OpenCL platforms/devices found") => {
-                    note_cpu_fallback_once();
-                    crate::kernel_cpu::run_kernel_cpu(&processed_code, &mut geom, &flat_values)
-                }
-                r => r,
-            }
-        };
-        if let Err(e) = result {
-            if ocl_error.is_none() {
-                *ocl_error = Some(e);
-            }
-        }
-        return Some(soup_to_detail(&geom));
-    }
-
-    Some(input)
-}
-
-
-/// One stderr note per process when kernels silently move to the CPU
-/// reference — the sphere disappearing taught us "silently" is the problem.
-fn note_cpu_fallback_once() {
-    static ONCE: std::sync::Once = std::sync::Once::new();
-    ONCE.call_once(|| {
-        eprintln!("cce-designer: no OpenCL platform — node kernels running on the CPU reference backend");
-    });
-}
-
-struct OpenClCache {
-    device: opencl3::device::Device,
-    context: opencl3::context::Context,
-    queue: opencl3::command_queue::CommandQueue,
-    kernels: std::collections::HashMap<String, opencl3::kernel::Kernel>,
-}
-
-static OPENCL_CACHE: std::sync::OnceLock<std::sync::Mutex<Option<OpenClCache>>> = std::sync::OnceLock::new();
-
-/// Every `clGetPlatformIDs` in this process goes through here, one at a time.
-///
-/// **Enumerating OpenCL platforms is not safe to call concurrently on this
-/// stack.** Mesa's Rusticl ICD closes a file descriptor it does not own while
-/// enumerating — caught under `strace -k` on 2026-09-23, the stack reading
-/// `__close` <- libRusticlOpenCL <- `clIcdGetPlatformIDsKHR` <-
-/// `clGetPlatformIDs`. The fd it closes has already been recycled to whatever
-/// another thread opened a moment earlier, so that thread's next `read` comes
-/// back `EBADF` on a file nothing was wrong with.
-///
-/// What it looked like: roughly one suite run in eight failed somewhere
-/// unrelated, most often `test_every_template_names_a_type_something_resolves`
-/// reporting "load_fs_tree returned 50 of 51 templates" — a node template
-/// whose `fs::read_to_string` had been handed EBADF and which `load_fs_tree`
-/// then dropped without a word. Four tests probed the platform independently
-/// to decide whether to skip, and with the suite running its tests in
-/// parallel those probes overlapped each other and everything else.
-///
-/// Serializing is this side of the fix and [`has_opencl_platform`] is the
-/// other: together they take a process from dozens of overlapping
-/// enumerations to two that cannot overlap. The bug is in the ICD and cannot
-/// be fixed from here — what can be fixed is how often, and how
-/// concurrently, we ask.
-fn probe_platforms() -> Vec<opencl3::platform::Platform> {
-    static PROBE: std::sync::Mutex<()> = std::sync::Mutex::new(());
-    // A poisoned probe lock carries no state worth protecting.
-    let _serialize = PROBE.lock().unwrap_or_else(|e| e.into_inner());
-    get_platforms().unwrap_or_default()
-}
-
-/// Whether this machine has an OpenCL platform at all, asked ONCE per process.
-///
-/// The answer cannot change while the process runs, and asking costs an
-/// enumeration — which, per [`probe_platforms`], is the thing that corrupts
-/// unrelated file descriptors. Every test that skips itself without a
-/// platform reads this rather than probing for itself.
-pub(crate) fn has_opencl_platform() -> bool {
-    static HAS: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
-    *HAS.get_or_init(|| {
-        // `CCE_KERNEL_CPU=1` means the CPU backend, so there is nothing to
-        // ask — and asking is what loads the ICD and costs a descriptor.
-        // Answering without probing is what makes the variable a usable
-        // workaround for the bug above rather than half of one: it took the
-        // suite from roughly one run in eight to none.
-        !crate::kernel_cpu::forced() && !probe_platforms().is_empty()
-    })
-}
-
-fn init_opencl() -> Option<OpenClCache> {
-    let platforms = probe_platforms();
-    if platforms.is_empty() {
-        return None;
-    }
-    let mut device_id = None;
-    for platform in &platforms {
-        if let Ok(devices) = platform.get_devices(CL_DEVICE_TYPE_GPU) {
-            if !devices.is_empty() {
-                device_id = Some(devices[0]);
-                break;
-            }
-        }
-    }
-    if device_id.is_none() {
-        for platform in &platforms {
-            if let Ok(devices) = platform.get_devices(CL_DEVICE_TYPE_CPU) {
-                if !devices.is_empty() {
-                    device_id = Some(devices[0]);
-                    break;
-                }
-            }
-        }
-    }
-    let device_id = device_id?;
-    let device = Device::new(device_id);
-    let context = Context::from_device(&device).ok()?;
-    let queue = unsafe { CommandQueue::create_with_properties(&context, device_id, 0, 0) }.ok()?;
-    Some(OpenClCache {
-        device,
-        context,
-        queue,
-        kernels: std::collections::HashMap::new(),
-    })
-}
-
-pub fn run_opencl_kernel(code: &str, geom: &mut Geometry) -> Result<(), String> {
-    run_opencl_kernel_with_params(code, geom, &[])
-}
-
-pub fn run_opencl_kernel_with_params(code: &str, geom: &mut Geometry, params: &[f32]) -> Result<(), String> {
-    // One gate, so no path can load the ICD once the CPU backend is forced —
-    // not the fallback in `apply_opencl`, not the cross-backend test that
-    // calls straight in here. Reported as the no-platform error every caller
-    // already handles, because "you told me not to use OpenCL" and "there is
-    // no OpenCL" want the same response from all of them.
-    if crate::kernel_cpu::forced() {
-        return Err("No OpenCL platforms/devices found (CCE_KERNEL_CPU is set)".to_string());
-    }
-    let is_generator = code.contains("out_count");
-    if geom.vertices.is_empty() && !is_generator {
-        return Ok(());
-    }
-
-    let mut cache_guard = OPENCL_CACHE
-        .get_or_init(|| std::sync::Mutex::new(init_opencl()))
-        .lock()
-        .map_err(|e| format!("Failed to lock OpenCL cache: {:?}", e))?;
-
-    let cache = cache_guard.as_mut().ok_or_else(|| "No OpenCL platforms/devices found".to_string())?;
-
-    if !cache.kernels.contains_key(code) {
-        let mut program = Program::create_from_source(&cache.context, code)
-            .map_err(|e| format!("Failed to create Program: {:?}", e))?;
-        if let Err(e) = program.build(&[cache.device.id()], "") {
-            let log = program.get_build_log(cache.device.id()).unwrap_or_else(|_| "Failed to retrieve build log".to_string());
-            return Err(format!("OpenCL JIT compilation error: {}\nLog:\n{}", e, log));
-        }
-        let kernel = Kernel::create(&program, "process")
-            .map_err(|e| format!("Failed to create kernel 'process': {:?}", e))?;
-        cache.kernels.insert(code.to_string(), kernel);
-    }
-    let kernel = cache.kernels.get(code).unwrap();
-    let context = &cache.context;
-    let queue = &cache.queue;
-
-    // Prepare parameter values buffer
-    let mut param_values_data = params.to_vec();
-    if param_values_data.is_empty() {
-        param_values_data.push(0.0);
-    }
-    let mut param_values_buf = unsafe {
-        ClBuffer::<cl_float>::create(&context, CL_MEM_READ_WRITE, param_values_data.len(), std::ptr::null_mut())
-            .map_err(|e| format!("Failed to create param_values buffer: {:?}", e))?
-    };
-    let _write_param_event = unsafe {
-        queue.enqueue_write_buffer(&mut param_values_buf, CL_TRUE, 0, &param_values_data, &[])
-            .map_err(|e| format!("Failed to write param_values buffer: {:?}", e))?
-    };
-
-    if is_generator {
-        let in_count = geom.vertices.len();
-        let max_vertices = 200_000;
-
-        // Prepare flat input position and color data
-        let mut in_pos_data: Vec<cl_float> = Vec::with_capacity(in_count * 3);
-        let mut in_col_data: Vec<cl_float> = Vec::with_capacity(in_count * 3);
-        for v in &geom.vertices {
-            in_pos_data.extend_from_slice(&v.pos);
-            in_col_data.extend_from_slice(&v.col);
-        }
-
-        // Create GPU buffers for inputs
-        let mut in_pos_buf = unsafe {
-            ClBuffer::<cl_float>::create(&context, CL_MEM_READ_WRITE, (in_count * 3).max(1), std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create input positions buffer: {:?}", e))?
-        };
-        let mut in_col_buf = unsafe {
-            ClBuffer::<cl_float>::create(&context, CL_MEM_READ_WRITE, (in_count * 3).max(1), std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create input colors buffer: {:?}", e))?
-        };
-
-        // Write input data to GPU
-        let _write_pos_event = unsafe {
-            let write_data = if in_pos_data.is_empty() { &[0.0f32] } else { &in_pos_data[..] };
-            queue.enqueue_write_buffer(&mut in_pos_buf, CL_TRUE, 0, write_data, &[])
-                .map_err(|e| format!("Failed to write input positions buffer: {:?}", e))?
-        };
-        let _write_col_event = unsafe {
-            let write_data = if in_col_data.is_empty() { &[0.0f32] } else { &in_col_data[..] };
-            queue.enqueue_write_buffer(&mut in_col_buf, CL_TRUE, 0, write_data, &[])
-                .map_err(|e| format!("Failed to write input colors buffer: {:?}", e))?
-        };
-
-        // Create GPU buffers for outputs
-        let out_pos_buf = unsafe {
-            ClBuffer::<cl_float>::create(&context, CL_MEM_READ_WRITE, max_vertices * 3, std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create output positions buffer: {:?}", e))?
-        };
-        let out_col_buf = unsafe {
-            ClBuffer::<cl_float>::create(&context, CL_MEM_READ_WRITE, max_vertices * 3, std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create output colors buffer: {:?}", e))?
-        };
-
-        // Create output count buffer initialized to 0
-        let mut out_count_buf = unsafe {
-            ClBuffer::<cl_int>::create(&context, CL_MEM_READ_WRITE, 1, std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create output count buffer: {:?}", e))?
-        };
-        let initial_count_data: [cl_int; 1] = [0];
-        let _write_count_event = unsafe {
-            queue.enqueue_write_buffer(&mut out_count_buf, CL_TRUE, 0, &initial_count_data, &[])
-                .map_err(|e| format!("Failed to write output count buffer: {:?}", e))?
-        };
-
-        // Execute kernel
-        let global_work_size = if in_count == 0 { 1 } else { in_count };
-        let mut exec = ExecuteKernel::new(kernel);
-        let kernel_event = unsafe {
-            exec.set_arg(&in_pos_buf)
-                .set_arg(&in_col_buf)
-                .set_arg(&(in_count as cl_int))
-                .set_arg(&out_pos_buf)
-                .set_arg(&out_col_buf)
-                .set_arg(&out_count_buf)
-                .set_arg(&(max_vertices as cl_int));
-
-            let num_args = kernel.num_args().unwrap_or(0);
-            if num_args >= 8 {
-                exec.set_arg(&param_values_buf);
-            }
-
-            exec.set_global_work_size(global_work_size)
-                .enqueue_nd_range(&queue)
-                .map_err(|e| format!("Failed to enqueue kernel: {:?}", e))?
-        };
-
-        kernel_event.wait().map_err(|e| format!("Failed to wait for kernel: {:?}", e))?;
-
-        // Read count back
-        let mut final_count_data: [cl_int; 1] = [0];
-        let _read_count_event = unsafe {
-            queue.enqueue_read_buffer(&out_count_buf, CL_TRUE, 0, &mut final_count_data, &[])
-                .map_err(|e| format!("Failed to read output count: {:?}", e))?
-        };
-        let final_count = (final_count_data[0] as usize).min(max_vertices);
-
-        // Read output positions and colors back
-        let mut out_pos_data: Vec<cl_float> = vec![0.0; final_count * 3];
-        let mut out_col_data: Vec<cl_float> = vec![0.0; final_count * 3];
-        if final_count > 0 {
-            let _read_pos_event = unsafe {
-                queue.enqueue_read_buffer(&out_pos_buf, CL_TRUE, 0, &mut out_pos_data, &[])
-                    .map_err(|e| format!("Failed to read output positions buffer: {:?}", e))?
-            };
-            let _read_col_event = unsafe {
-                queue.enqueue_read_buffer(&out_col_buf, CL_TRUE, 0, &mut out_col_data, &[])
-                    .map_err(|e| format!("Failed to read output colors buffer: {:?}", e))?
-            };
-        }
-
-        // Rebuild geometry vertices
-        geom.vertices.clear();
-        for i in 0..final_count {
-            let mut attributes = HashMap::new();
-            attributes.insert("Norm".to_string(), GAttribute::Float3([0.0, 1.0, 0.0]));
-            attributes.insert("UV".to_string(), GAttribute::Float2([0.0, 0.0]));
-            geom.vertices.push(GVertex {
-                pos: [out_pos_data[i * 3], out_pos_data[i * 3 + 1], out_pos_data[i * 3 + 2]],
-                col: [out_col_data[i * 3], out_col_data[i * 3 + 1], out_col_data[i * 3 + 2]],
-                attributes,
-            });
-        }
-    } else {
-        let mut pos_data: Vec<cl_float> = Vec::with_capacity(geom.vertices.len() * 3);
-        let mut col_data: Vec<cl_float> = Vec::with_capacity(geom.vertices.len() * 3);
-        for v in &geom.vertices {
-            pos_data.extend_from_slice(&v.pos);
-            col_data.extend_from_slice(&v.col);
-        }
-        let count = geom.vertices.len();
-        drop(cache_guard);
-        run_deformer_flat(code, &mut pos_data, &mut col_data, count, &mut [], params)?;
-        for i in 0..count {
-            geom.vertices[i].pos = [pos_data[i * 3], pos_data[i * 3 + 1], pos_data[i * 3 + 2]];
-            geom.vertices[i].col = [col_data[i * 3], col_data[i * 3 + 1], col_data[i * 3 + 2]];
-        }
-        return Ok(());
-    }
-
-    Ok(())
-}
-
-/// Run a deformer kernel over flat per-element buffers, reading everything back
-/// in place.
-///
-/// `pos` and `col` hold three floats per element; each entry of `attrs` holds
-/// one, and they bind in the order [`parse_attr_refs`] found them — which is
-/// the order [`preprocess_opencl_code`] appended them to the signature.
-///
-/// The one implementation behind both the soup path (which binds no
-/// attributes) and the point-native path.
-fn run_deformer_flat(
-    code: &str,
-    pos: &mut Vec<cl_float>,
-    col: &mut Vec<cl_float>,
-    count: usize,
-    attrs: &mut [(String, Vec<cl_float>)],
-    params: &[f32],
-) -> Result<(), String> {
-    if count == 0 {
-        return Ok(());
-    }
-    let mut cache_guard = OPENCL_CACHE
-        .get_or_init(|| std::sync::Mutex::new(init_opencl()))
-        .lock()
-        .map_err(|e| format!("Failed to lock OpenCL cache: {:?}", e))?;
-    let cache = cache_guard.as_mut().ok_or_else(|| "No OpenCL platforms/devices found".to_string())?;
-
-    if !cache.kernels.contains_key(code) {
-        let mut program = Program::create_from_source(&cache.context, code)
-            .map_err(|e| format!("Failed to create Program: {:?}", e))?;
-        if let Err(e) = program.build(&[cache.device.id()], "") {
-            let log = program
-                .get_build_log(cache.device.id())
-                .unwrap_or_else(|_| "Failed to retrieve build log".to_string());
-            return Err(format!("OpenCL JIT compilation error: {}\nLog:\n{}", e, log));
-        }
-        let kernel = Kernel::create(&program, "process")
-            .map_err(|e| format!("Failed to create kernel 'process': {:?}", e))?;
-        cache.kernels.insert(code.to_string(), kernel);
-    }
-    let kernel = cache.kernels.get(code).unwrap();
-    let (context, queue) = (&cache.context, &cache.queue);
-
-    let mut param_data = params.to_vec();
-    if param_data.is_empty() {
-        param_data.push(0.0);
-    }
-
-    let make = |data: &[cl_float]| -> Result<ClBuffer<cl_float>, String> {
-        let mut buf = unsafe {
-            ClBuffer::<cl_float>::create(context, CL_MEM_READ_WRITE, data.len().max(1), std::ptr::null_mut())
-                .map_err(|e| format!("Failed to create buffer: {:?}", e))?
-        };
-        unsafe {
-            queue
-                .enqueue_write_buffer(&mut buf, CL_TRUE, 0, data, &[])
-                .map_err(|e| format!("Failed to write buffer: {:?}", e))?
-        };
-        Ok(buf)
-    };
-
-    let mut pos_buf = make(pos)?;
-    let mut col_buf = make(col)?;
-    let param_buf = make(&param_data)?;
-    let mut attr_bufs: Vec<ClBuffer<cl_float>> = Vec::with_capacity(attrs.len());
-    for (_, data) in attrs.iter() {
-        attr_bufs.push(make(data)?);
-    }
-
-    let mut exec = ExecuteKernel::new(kernel);
-    let kernel_event = unsafe {
-        exec.set_arg(&pos_buf).set_arg(&col_buf).set_arg(&(count as cl_int));
-        // The signature carries param_values only when the kernel names a
-        // parameter, so the attribute buffers slide up one slot when it does
-        // not. Arity is the authority, exactly as it was for the old
-        // `num_args >= 4` check.
-        let num_args = kernel.num_args().unwrap_or(0) as usize;
-        if num_args > 3 + attrs.len() {
-            exec.set_arg(&param_buf);
-        }
-        for buf in &attr_bufs {
-            exec.set_arg(buf);
-        }
-        exec.set_global_work_size(count)
-            .enqueue_nd_range(queue)
-            .map_err(|e| format!("Failed to enqueue kernel: {:?}", e))?
-    };
-    kernel_event.wait().map_err(|e| format!("Failed to wait for kernel: {:?}", e))?;
-
-    unsafe {
-        queue
-            .enqueue_read_buffer(&mut pos_buf, CL_TRUE, 0, pos, &[])
-            .map_err(|e| format!("Failed to read positions buffer: {:?}", e))?;
-        queue
-            .enqueue_read_buffer(&mut col_buf, CL_TRUE, 0, col, &[])
-            .map_err(|e| format!("Failed to read colors buffer: {:?}", e))?;
-        for (buf, (_, data)) in attr_bufs.iter().zip(attrs.iter_mut()) {
-            queue
-                .enqueue_read_buffer(buf, CL_TRUE, 0, data, &[])
-                .map_err(|e| format!("Failed to read attribute buffer: {:?}", e))?;
-        }
-    }
-    Ok(())
-}
-
-/// Run a DEFORMER kernel over a [`Detail`]'s points.
-///
-/// This is the Phase 1 ABI: the work item is a POINT, not a triangle corner,
-/// and the kernel reaches named attributes through buffers of its own. Two
-/// things follow. A deformer no longer flattens and welds — topology, groups
-/// and point identities are simply untouched, because the kernel never saw
-/// them. And "the first corner wins", which the soup round trip had to invent
-/// where corners of one point disagreed, stops being a question: there is one
-/// value per point because there is one point.
-///
-/// An attribute the kernel names but the geometry lacks is created and zeroed:
-/// naming it is the declaration.
-pub fn run_kernel_on_detail(
-    code: &str,
-    names: &[String],
-    geom: &mut Detail,
-    params: &[f32],
-) -> Result<(), String> {
-    let count = geom.num_points();
-    if count == 0 {
-        return Ok(());
-    }
-    let mut pos: Vec<cl_float> = bytemuck::cast_slice(geom.positions()).to_vec();
-    let mut col: Vec<cl_float> = Vec::with_capacity(count * 3);
-    for p in 0..count {
-        col.extend_from_slice(&geom.color(p));
-    }
-
-    // The names come from the caller, NOT from `code`: by the time a kernel
-    // reaches here the preprocessor has rewritten every `attrf("mass", i)`
-    // into `attr_0[i]`, so there is nothing left to parse. Reading them off
-    // the processed source bound zero buffers and the attribute silently never
-    // appeared — which is the gap between a test that parses raw source and a
-    // test that runs a kernel with its buffers handed to it.
-    let mut attrs: Vec<(String, Vec<cl_float>)> = Vec::with_capacity(names.len());
-    for name in names {
-        let data = match geom.points().get(name) {
-            Some(AttribData::Float(v)) => v.clone(),
-            Some(other) => (0..count)
-                .map(|p| other.get(p).map(|v| v.as_f32()).unwrap_or(0.0))
-                .collect(),
-            None => vec![0.0; count],
-        };
-        attrs.push((name.clone(), data));
-    }
-
-    let run = if crate::kernel_cpu::forced() {
-        crate::kernel_cpu::run_deformer_cpu(code, &mut pos, &mut col, count, &mut attrs, params)
-    } else {
-        match run_deformer_flat(code, &mut pos, &mut col, count, &mut attrs, params) {
-            Err(e) if e.contains("No OpenCL platforms/devices found") => {
-                note_cpu_fallback_once();
-                crate::kernel_cpu::run_deformer_cpu(code, &mut pos, &mut col, count, &mut attrs, params)
-            }
-            r => r,
-        }
-    };
-    run?;
-
-    for (p, slot) in geom.positions_mut().iter_mut().enumerate() {
-        *slot = [pos[p * 3], pos[p * 3 + 1], pos[p * 3 + 2]];
-    }
-    for p in 0..count {
-        geom.set_color(p, [col[p * 3], col[p * 3 + 1], col[p * 3 + 2]]);
-    }
-    for (name, data) in attrs {
-        let _ = geom.points_mut().insert(&name, AttribData::Float(data));
-    }
-    Ok(())
-}
-
 /// The NATIVE geometry types — the ones
 /// [`generate_single_node_geometry_with_errors`] dispatches on directly.
 ///
@@ -6079,7 +5099,7 @@ pub fn network_sphere_vertices_with_errors(
             *count += 1;
             if is_visible {
                 let mut visited = Vec::new();
-                if let Some(geom) = resolve_opencl_geometry_with_errors(root, node, &mut visited, ocl_error, sim) {
+                if let Some(geom) = retired_opencl_node(root, node, &mut visited, ocl_error, sim) {
                     out.merge(&geom);
                 }
             }
@@ -6197,7 +5217,6 @@ pub fn find_sphere_index(root: &FsNode, target: &FsNode) -> Option<usize> {
             || node.node_type.eq_ignore_ascii_case("line") 
             || node.node_type.eq_ignore_ascii_case("points")
             || node.node_type.eq_ignore_ascii_case("transform")
-            || node.node_type.eq_ignore_ascii_case("opencl")
             || node.node_type.eq_ignore_ascii_case("scatter") {
             let idx = *count;
             *count += 1;
@@ -6670,106 +5689,6 @@ mod tests {
         v
     }
 
-    /// A deformer that grows a named attribute and displaces by it — the
-    /// smallest kernel that needs the Phase 1 ABI.
-    const ATTR_KERNEL: &str = r#"
-        __kernel void process(__global float* pos, __global float* col, int count) {
-            int id = get_global_id(0);
-            if (id < count) {
-                setattrf("mass", id, attrf("mass", id) + 2.0f);
-                pos[id * 3 + 1] += attrf("mass", id);
-            }
-        }
-    "#;
-
-    #[test]
-    fn test_kernel_names_its_attributes_and_they_become_buffers() {
-        assert_eq!(parse_attr_refs(ATTR_KERNEL), vec!["mass".to_string()]);
-        // `setattrf(` contains `attrf(`; the outer call must not be counted
-        // twice, or the second buffer would shift every later binding.
-        assert_eq!(parse_attr_refs(r#"setattrf("a", i, 1.0f);"#), vec!["a".to_string()]);
-
-        let out = preprocess_opencl_code(ATTR_KERNEL);
-        assert!(out.contains("__global float* attr_0"), "{out}");
-        assert!(out.contains("attr_0[id] = (attr_0[id] + 2.0f)"), "{out}");
-        assert!(!out.contains("attrf("), "every call is rewritten: {out}");
-        // No ch* parameters here, so param_values is absent and the attribute
-        // buffer takes the fourth slot. The launchers read arity to tell.
-        assert!(!out.contains("param_values"), "{out}");
-    }
-
-    #[test]
-    fn test_deformer_runs_over_points_and_keeps_the_geometry_whole() {
-        let mut d = sphere_detail(Vec3::ZERO, 1.0, 6, 8);
-        d.points_mut().create_group("keep");
-        d.points_mut().add_to_group("keep", 3);
-        let before_ids = d.ids().to_vec();
-        let before_prims = d.num_prims();
-        let before_y: Vec<f32> = d.positions().iter().map(|p| p[1]).collect();
-
-        let code = preprocess_opencl_code(ATTR_KERNEL);
-        let count = d.num_points();
-        let mut pos: Vec<f32> = bytemuck::cast_slice(d.positions()).to_vec();
-        let mut col: Vec<f32> = (0..count).flat_map(|p| d.color(p)).collect();
-        let mut attrs = vec![("mass".to_string(), vec![0.0f32; count])];
-        crate::kernel_cpu::run_deformer_cpu(&code, &mut pos, &mut col, count, &mut attrs, &[]).unwrap();
-
-        // One work item per POINT, not per triangle corner: 42 points where the
-        // soup would have handed the kernel 240 corners and then had to decide
-        // which corner's answer a shared point takes.
-        assert_eq!(count, 42);
-        assert_eq!(attrs[0].1, vec![2.0f32; 42], "the kernel created and wrote the attribute");
-        for p in 0..count {
-            assert!((pos[p * 3 + 1] - (before_y[p] + 2.0)).abs() < 1e-5, "point {p}");
-        }
-
-        // Write back through the real entry point and check the geometry is
-        // otherwise untouched — this is what the soup round trip could not do.
-        for (p, slot) in d.positions_mut().iter_mut().enumerate() {
-            *slot = [pos[p * 3], pos[p * 3 + 1], pos[p * 3 + 2]];
-        }
-        let _ = d.points_mut().insert("mass", AttribData::Float(attrs[0].1.clone()));
-        assert_eq!(d.ids(), &before_ids[..], "identities survive a deformer");
-        assert_eq!(d.num_prims(), before_prims, "topology survives a deformer");
-        assert_eq!(d.points().group_members("keep"), vec![3], "groups survive a deformer");
-        assert_eq!(d.points().value("mass", 0), Some(AttribValue::Float(2.0)));
-    }
-
-    #[test]
-    fn test_attribute_abi_matches_across_both_backends() {
-        if !crate::geometry::has_opencl_platform() {
-            println!("Skipping cross-backend attribute ABI test: no OpenCL platform");
-            return;
-        }
-        let d = sphere_detail(Vec3::ZERO, 1.0, 6, 8);
-        let code = preprocess_opencl_code(ATTR_KERNEL);
-        let count = d.num_points();
-        let base_pos: Vec<f32> = bytemuck::cast_slice(d.positions()).to_vec();
-        let base_col: Vec<f32> = (0..count).flat_map(|p| d.color(p)).collect();
-
-        let run = |gpu: bool| -> (Vec<f32>, Vec<f32>) {
-            let (mut pos, mut col) = (base_pos.clone(), base_col.clone());
-            let mut attrs = vec![("mass".to_string(), vec![0.5f32; count])];
-            if gpu {
-                run_deformer_flat(&code, &mut pos, &mut col, count, &mut attrs, &[]).unwrap();
-            } else {
-                crate::kernel_cpu::run_deformer_cpu(&code, &mut pos, &mut col, count, &mut attrs, &[])
-                    .unwrap();
-            }
-            (pos, attrs.remove(0).1)
-        };
-
-        let (gpu_pos, gpu_mass) = run(true);
-        let (cpu_pos, cpu_mass) = run(false);
-        // The interpreter is the semantic reference, so the widened ABI has to
-        // bind the same arguments to the same slots on both sides.
-        assert_eq!(gpu_mass, cpu_mass);
-        for (i, (a, b)) in gpu_pos.iter().zip(cpu_pos.iter()).enumerate() {
-            assert!((a - b).abs() < 1e-5, "component {i}: {a} vs {b}");
-        }
-        assert_eq!(gpu_mass, vec![2.5f32; count], "an existing attribute is read, not reset");
-    }
-
     #[test]
     fn test_welded_sphere_reproduces_the_soup_it_replaced() {
         let (center, radius, lat, lon) = (Vec3::new(0.1, 0.2, 0.3), 0.7, 16, 24);
@@ -6950,95 +5869,6 @@ mod tests {
         assert_eq!(ids.len(), 16);
     }
 
-    #[test]
-    fn test_opencl_deformer_mode() {
-        if !crate::geometry::has_opencl_platform() {
-            println!("Skipping OpenCL deformer test: No OpenCL platforms found");
-            return;
-        }
-
-        let code = r#"
-            __kernel void process(__global float* pos, __global float* col, int count) {
-                int id = get_global_id(0);
-                if (id < count) {
-                    pos[id * 3 + 1] += 1.0f;
-                }
-            }
-        "#;
-
-        let mut geom = Geometry {
-            vertices: vec![GVertex {
-                pos: [1.0, 2.0, 3.0],
-                col: [1.0, 0.0, 0.0],
-                attributes: std::collections::HashMap::new(),
-            }],
-        };
-
-        run_opencl_kernel(code, &mut geom).unwrap();
-
-        assert_eq!(geom.vertices.len(), 1);
-        assert_eq!(geom.vertices[0].pos, [1.0, 3.0, 3.0]);
-    }
-
-    #[test]
-    fn test_opencl_generator_mode() {
-        if !crate::geometry::has_opencl_platform() {
-            println!("Skipping OpenCL generator test: No OpenCL platforms found");
-            return;
-        }
-
-        let code = r#"
-            __kernel void process(
-                __global const float* in_pos,
-                __global const float* in_col,
-                int in_count,
-                __global float* out_pos,
-                __global float* out_col,
-                __global int* out_count,
-                int max_out_count
-            ) {
-                int id = get_global_id(0);
-                if (id < in_count) {
-                    // Copy original
-                    int idx1 = atomic_inc(out_count);
-                    if (idx1 < max_out_count) {
-                        out_pos[idx1 * 3] = in_pos[id * 3];
-                        out_pos[idx1 * 3 + 1] = in_pos[id * 3 + 1];
-                        out_pos[idx1 * 3 + 2] = in_pos[id * 3 + 2];
-                        out_col[idx1 * 3] = in_col[id * 3];
-                        out_col[idx1 * 3 + 1] = in_col[id * 3 + 1];
-                        out_col[idx1 * 3 + 2] = in_col[id * 3 + 2];
-                    }
-                    // Generate new offset
-                    int idx2 = atomic_inc(out_count);
-                    if (idx2 < max_out_count) {
-                        out_pos[idx2 * 3] = in_pos[id * 3] + 1.0f;
-                        out_pos[idx2 * 3 + 1] = in_pos[id * 3 + 1] + 2.0f;
-                        out_pos[idx2 * 3 + 2] = in_pos[id * 3 + 2] + 3.0f;
-                        out_col[idx2 * 3] = 0.5f;
-                        out_col[idx2 * 3 + 1] = 0.5f;
-                        out_col[idx2 * 3 + 2] = 0.5f;
-                    }
-                }
-            }
-        "#;
-
-        let mut geom = Geometry {
-            vertices: vec![GVertex {
-                pos: [1.0, 2.0, 3.0],
-                col: [1.0, 0.0, 0.0],
-                attributes: std::collections::HashMap::new(),
-            }],
-        };
-
-        run_opencl_kernel(code, &mut geom).unwrap();
-
-        assert_eq!(geom.vertices.len(), 2);
-        assert_eq!(geom.vertices[0].pos, [1.0, 2.0, 3.0]);
-        assert_eq!(geom.vertices[1].pos, [2.0, 4.0, 6.0]);
-        assert_eq!(geom.vertices[1].col, [0.5, 0.5, 0.5]);
-    }
-
     #[test]
     fn test_points_node_shapes() {
         let points_node = |shape: &str| FsNode {
@@ -7302,102 +6132,6 @@ mod tests {
         assert!(geom_loop.is_none());
     }
 
-    #[test]
-    fn test_opencl_local_node() {
-        if !crate::geometry::has_opencl_platform() {
-            println!("Skipping OpenCL local node test: No OpenCL platforms found");
-            return;
-        }
-
-        let sphere = FsNode {
-            id: "id Sphere 1".to_string(),
-            inputs: 1,
-            outputs: 1,
-            name: "Sphere 1".to_string(),
-            node_type: "sphere".to_string(),
-            children: vec![],
-            params: vec![
-                ParamDef {
-                    name: "Radius".to_string(),
-                    label: String::new(),
-                    param_type: "slider".to_string(),
-                    default: "0.5".to_string(),
-                    options: vec![],
-                    min: None,
-                    max: None,
-                    step: None,
-                    show_when: String::new(), expr: false,
-                }
-            ],
-            geometry_visible: true,
-            position: (0.0, 0.0),
-        };
-
-        let opencl_node = FsNode {
-            id: "id OpenCL 1".to_string(),
-            inputs: 1,
-            outputs: 1,
-            name: "OpenCL 1".to_string(),
-            node_type: "opencl".to_string(),
-            children: vec![],
-            params: vec![
-                ParamDef {
-                    name: "Input".to_string(),
-                    label: String::new(),
-                    param_type: "text".to_string(),
-                    default: "Sphere 1".to_string(),
-                    options: vec![],
-                    min: None,
-                    max: None,
-                    step: None,
-                    show_when: String::new(), expr: false,
-                },
-                ParamDef {
-                    name: "Code".to_string(),
-                    label: String::new(),
-                    param_type: "code".to_string(),
-                    default: r#"
-                        __kernel void process(__global float* pos, __global float* col, int count) {
-                            int id = get_global_id(0);
-                            if (id < count) {
-                                pos[id * 3 + 1] += 2.0f;
-                            }
-                        }
-                    "#.to_string(),
-                    options: vec![],
-                    min: None,
-                    max: None,
-                    step: None,
-                    show_when: String::new(), expr: false,
-                }
-            ],
-            geometry_visible: true,
-            position: (0.0, 0.0),
-        };
-
-        let root = FsNode {
-            id: "id root".to_string(),
-            inputs: 1,
-            outputs: 1,
-            name: "root".to_string(),
-            node_type: "node".to_string(),
-            children: vec![sphere, opencl_node.clone()],
-            params: vec![],
-            geometry_visible: true,
-            position: (0.0, 0.0),
-        };
-
-        let mut visited = Vec::new();
-        let mut err = None;
-        let geom = resolve_opencl_geometry_with_errors(&root, &opencl_node, &mut visited, &mut err, &mut crate::geometry::EvalSim::new(0, 0, &mut crate::geometry::SimCache::default())).unwrap();
-        assert!(!geom.is_empty());
-        assert!(err.is_none());
-
-        // The sphere should be translated up by 2.0 on the y axis compared to the standard sphere (which centers around y=0.55 for index 0)
-        let avg_y = geom.positions().iter().map(|p| p[1]).sum::<f32>() / geom.num_points() as f32;
-        assert!((avg_y - 2.55).abs() < 0.01);
-    }
-
     #[test]
     fn test_scatter_node() {
         let sphere = FsNode {
@@ -8338,53 +7072,6 @@ mod simnet_tests {
         );
     }
 
-    #[test]
-    fn test_an_opencl_deformer_lands_its_named_attribute_on_the_geometry() {
-        // The gap between "parse_attr_refs reads raw source" and "run the
-        // kernel with buffers handed to it": run_kernel_on_detail was reading
-        // the names off the PROCESSED source, where every attrf() call has
-        // already become attr_0[], so it bound no buffers and the attribute
-        // silently never appeared. Only an end-to-end evaluation catches that.
-        let kernel = r#"
-            __kernel void process(__global float* pos, __global float* col, int count) {
-                int id = get_global_id(0);
-                if (id < count) {
-                    setattrf("mass", id, pos[id * 3 + 1] * 2.0f);
-                }
-            }
-        "#;
-        let sphere = node("id-s", "Sphere 1", "sphere", vec![param("Radius", "0.5")], vec![]);
-        let dfm = node(
-            "id-k",
-            "Height 1",
-            "opencl",
-            vec![param("Input", "Sphere 1"), param("Code", kernel)],
-            vec![],
-        );
-        let root = node("id-root", "root", "node", vec![], vec![sphere, dfm]);
-
-        let mut visited = Vec::new();
-        let mut err = None;
-        let mut cache = SimCache::default();
-        let mut sim = EvalSim::new(0, 0, &mut cache);
-        let g = generate_single_node_geometry_with_errors(
-            &root,
-            &root.children[1],
-            &mut visited,
-            &mut err,
-            &mut sim,
-        )
-        .expect("the deformer evaluates");
-        assert!(err.is_none(), "{err:?}");
-
-        assert!(g.points().has("mass"), "the kernel's named attribute must reach the geometry");
-        for p in 0..g.num_points() {
-            let want = g.positions()[p][1] * 2.0;
-            let got = g.points().value("mass", p).unwrap().as_f32();
-            assert!((got - want).abs() < 1e-4, "point {p}: {got} vs {want}");
-        }
-    }
-
     #[test]
     fn test_visualize_reports_a_missing_attribute_and_leaves_colour_alone() {
         let before = sphere_detail(Vec3::ZERO, 0.5, 4, 6);
diff --git a/src/kernel_cpu.rs b/src/kernel_cpu.rs
deleted file mode 100644
index 80609c2..0000000
--- a/src/kernel_cpu.rs
+++ /dev/null
@@ -1,1809 +0,0 @@
-//! CPU reference backend for the node kernels.
-//!
-//! The designer's kernel language is (a subset of) OpenCL C, which made every
-//! kernel-generated node dependent on an OpenCL platform existing at runtime —
-//! and "zero platforms" is a reachable state on this machine (nvidia-open out
-//! of kernel lockstep, rusticl opt-in), reached silently: nodes just produced
-//! empty geometry. This module is the reference implementation of that same
-//! language on the CPU: a tree-walking interpreter for the C subset the
-//! kernels use, executing the exact launcher contract of
-//! `run_opencl_kernel_with_params` (same generator/deformer split, same
-//! positional argument binding, same `max_vertices`, same output rebuild).
-//!
-//! It is the SEMANTIC REFERENCE: f32 arithmetic, C int truncation, postfix
-//! increment, short-circuit logic — anything the GPU path disagrees with here
-//! is a bug on one side or the other, and the cross-validation test compares
-//! the two directly. Out-of-bounds buffer access — undefined behavior on the
-//! GPU — reads 0 and drops writes here. A step budget guards the UI thread
-//! against a kernel that never terminates (the GPU path would hang the queue
-//! on that too; here it is a clean error instead).
-//!
-//! Deliberately NOT here: vector types (float3), barriers, local memory,
-//! images — no shipped or plausible node kernel uses them. The interpreter
-//! sees the same post-preprocessed source as OpenCL (`chf()` already rewritten
-//! to `param_values[i]` by `preprocess_opencl_code`).
-
-use crate::geometry::{GAttribute, GVertex, Geometry};
-use std::collections::HashMap;
-
-pub const MAX_VERTICES: usize = 200_000;
-/// Interpreter step budget per kernel invocation (all work items together).
-/// Generously above any real kernel (the sphere is ~100k steps, a full
-/// 200k-vertex generator ~10M) while keeping a hung kernel's UI freeze in
-/// seconds, not minutes.
-const STEP_BUDGET: u64 = 50_000_000;
-
-// ---------------------------------------------------------------- lexer
-
-#[derive(Debug, Clone, PartialEq)]
-enum Tok {
-    Ident(String),
-    Int(i64),
-    Float(f32),
-    Punct(&'static str),
-}
-
-fn lex(src: &str) -> Result<Vec<Tok>, String> {
-    let b = src.as_bytes();
-    let mut out = Vec::new();
-    let mut i = 0;
-    while i < b.len() {
-        let c = b[i] as char;
-        if c.is_whitespace() {
-            i += 1;
-        } else if c == '/' && i + 1 < b.len() && b[i + 1] == b'/' {
-            while i < b.len() && b[i] != b'\n' {
-                i += 1;
-            }
-        } else if c == '/' && i + 1 < b.len() && b[i + 1] == b'*' {
-            i += 2;
-            while i + 1 < b.len() && !(b[i] == b'*' && b[i + 1] == b'/') {
-                i += 1;
-            }
-            i = (i + 2).min(b.len());
-        } else if c.is_ascii_alphabetic() || c == '_' {
-            let s = i;
-            while i < b.len() && ((b[i] as char).is_ascii_alphanumeric() || b[i] == b'_') {
-                i += 1;
-            }
-            out.push(Tok::Ident(src[s..i].to_string()));
-        } else if c.is_ascii_digit() || (c == '.' && i + 1 < b.len() && (b[i + 1] as char).is_ascii_digit()) {
-            let s = i;
-            let mut is_float = false;
-            while i < b.len() && (b[i] as char).is_ascii_digit() {
-                i += 1;
-            }
-            if i < b.len() && b[i] == b'.' {
-                is_float = true;
-                i += 1;
-                while i < b.len() && (b[i] as char).is_ascii_digit() {
-                    i += 1;
-                }
-            }
-            if i < b.len() && (b[i] == b'e' || b[i] == b'E') {
-                is_float = true;
-                i += 1;
-                if i < b.len() && (b[i] == b'+' || b[i] == b'-') {
-                    i += 1;
-                }
-                while i < b.len() && (b[i] as char).is_ascii_digit() {
-                    i += 1;
-                }
-            }
-            let text = &src[s..i];
-            if i < b.len() && (b[i] == b'f' || b[i] == b'F') {
-                is_float = true;
-                i += 1;
-            }
-            if is_float {
-                out.push(Tok::Float(text.parse::<f32>().map_err(|e| format!("bad float literal {text}: {e}"))?));
-            } else {
-                out.push(Tok::Int(text.parse::<i64>().map_err(|e| format!("bad int literal {text}: {e}"))?));
-            }
-        } else {
-            // Longest-match puncts.
-            const P3: [&str; 0] = [];
-            const P2: [&str; 14] = ["==", "!=", "<=", ">=", "&&", "||", "+=", "-=", "*=", "/=", "%=", "++", "--", "->"];
-            let _ = P3;
-            let rest = &src[i..];
-            let mut matched = None;
-            for p in P2 {
-                if rest.starts_with(p) {
-                    matched = Some(p);
-                    break;
-                }
-            }
-            if let Some(p) = matched {
-                out.push(Tok::Punct(p));
-                i += p.len();
-            } else {
-                const P1: &str = "+-*/%<>=!&|?:;,(){}[].";
-                if P1.contains(c) {
-                    let idx = P1.find(c).unwrap();
-                    out.push(Tok::Punct(&P1[idx..idx + 1]));
-                    i += 1;
-                } else {
-                    return Err(format!("kernel_cpu: unexpected character '{c}'"));
-                }
-            }
-        }
-    }
-    Ok(out)
-}
-
-// ---------------------------------------------------------------- AST
-
-#[derive(Debug, Clone, Copy, PartialEq)]
-enum Ty {
-    Float,
-    Int,
-    Void,
-}
-
-#[derive(Debug, Clone)]
-enum Expr {
-    F(f32),
-    I(i64),
-    Var(String),
-    Index(Box<Expr>, Box<Expr>),
-    Call(String, Vec<Expr>),
-    Unary(&'static str, Box<Expr>),
-    /// Postfix ++/-- (value is the OLD one).
-    Post(&'static str, Box<Expr>),
-    Bin(&'static str, Box<Expr>, Box<Expr>),
-    Assign(&'static str, Box<Expr>, Box<Expr>),
-    Ternary(Box<Expr>, Box<Expr>, Box<Expr>),
-    Cast(Ty, Box<Expr>),
-    AddrOf(Box<Expr>),
-    Deref(Box<Expr>),
-}
-
-#[derive(Debug, Clone)]
-enum Stmt {
-    Decl { ty: Ty, name: String, arr: Option<usize>, init: Vec<Expr> },
-    Expr(Expr),
-    If(Expr, Vec<Stmt>, Vec<Stmt>),
-    For(Option<Box<Stmt>>, Option<Expr>, Option<Expr>, Vec<Stmt>),
-    While(Expr, Vec<Stmt>),
-    Return(Option<Expr>),
-    Break,
-    Continue,
-    Block(Vec<Stmt>),
-    /// Comma-separated declarators (`float a = 1, b = 2;`): a sequence run in
-    /// the CURRENT scope — Block would give the declarations their own scope
-    /// and the names would vanish at the semicolon.
-    Seq(Vec<Stmt>),
-}
-
-#[derive(Debug, Clone)]
-struct Param {
-    name: String,
-    is_ptr: bool,
-}
-
-#[derive(Debug, Clone)]
-struct FnDef {
-    params: Vec<Param>,
-    body: Vec<Stmt>,
-}
-
-struct Parser {
-    toks: Vec<Tok>,
-    pos: usize,
-}
-
-impl Parser {
-    fn peek(&self) -> Option<&Tok> {
-        self.toks.get(self.pos)
-    }
-    fn next(&mut self) -> Option<Tok> {
-        let t = self.toks.get(self.pos).cloned();
-        self.pos += 1;
-        t
-    }
-    fn eat_punct(&mut self, p: &str) -> bool {
-        if matches!(self.peek(), Some(Tok::Punct(q)) if *q == p) {
-            self.pos += 1;
-            true
-        } else {
-            false
-        }
-    }
-    fn expect_punct(&mut self, p: &str) -> Result<(), String> {
-        if self.eat_punct(p) {
-            Ok(())
-        } else {
-            Err(format!("kernel_cpu: expected '{p}' at token {:?}", self.peek()))
-        }
-    }
-    fn eat_ident(&mut self, s: &str) -> bool {
-        if matches!(self.peek(), Some(Tok::Ident(q)) if q == s) {
-            self.pos += 1;
-            true
-        } else {
-            false
-        }
-    }
-    /// Skip OpenCL address-space and const qualifiers wherever they appear.
-    fn skip_quals(&mut self) {
-        loop {
-            match self.peek() {
-                Some(Tok::Ident(s))
-                    if matches!(
-                        s.as_str(),
-                        "__kernel" | "kernel" | "__global" | "global" | "__constant" | "constant"
-                            | "__local" | "local" | "__private" | "private" | "const" | "unsigned"
-                    ) =>
-                {
-                    self.pos += 1;
-                }
-                _ => break,
-            }
-        }
-    }
-    fn peek_ty(&self) -> Option<Ty> {
-        match self.peek() {
-            Some(Tok::Ident(s)) => match s.as_str() {
-                "float" => Some(Ty::Float),
-                "int" | "bool" => Some(Ty::Int),
-                "void" => Some(Ty::Void),
-                _ => None,
-            },
-            _ => None,
-        }
-    }
-
-    fn parse_program(&mut self) -> Result<HashMap<String, FnDef>, String> {
-        let mut fns = HashMap::new();
-        while self.peek().is_some() {
-            self.skip_quals();
-            let _ret = self.peek_ty().ok_or_else(|| format!("kernel_cpu: expected function return type, got {:?}", self.peek()))?;
-            self.pos += 1;
-            self.skip_quals();
-            // Pointer return types don't occur; a stray '*' here would be one.
-            let name = match self.next() {
-                Some(Tok::Ident(n)) => n,
-                t => return Err(format!("kernel_cpu: expected function name, got {t:?}")),
-            };
-            self.expect_punct("(")?;
-            let mut params = Vec::new();
-            if !self.eat_punct(")") {
-                loop {
-                    self.skip_quals();
-                    if self.peek_ty().is_none() {
-                        return Err(format!("kernel_cpu: expected parameter type, got {:?}", self.peek()));
-                    }
-                    self.pos += 1;
-                    self.skip_quals();
-                    let mut is_ptr = false;
-                    while self.eat_punct("*") {
-                        is_ptr = true;
-                    }
-                    self.skip_quals();
-                    let pname = match self.next() {
-                        Some(Tok::Ident(n)) => n,
-                        t => return Err(format!("kernel_cpu: expected parameter name, got {t:?}")),
-                    };
-                    params.push(Param { name: pname, is_ptr });
-                    if !self.eat_punct(",") {
-                        break;
-                    }
-                }
-                self.expect_punct(")")?;
-            }
-            self.expect_punct("{")?;
-            let body = self.parse_block_body()?;
-            fns.insert(name, FnDef { params, body });
-        }
-        Ok(fns)
-    }
-
-    fn parse_block_body(&mut self) -> Result<Vec<Stmt>, String> {
-        let mut stmts = Vec::new();
-        while !self.eat_punct("}") {
-            if self.peek().is_none() {
-                return Err("kernel_cpu: unexpected end of source inside a block".into());
-            }
-            stmts.push(self.parse_stmt()?);
-        }
-        Ok(stmts)
-    }
-
-    fn parse_stmt(&mut self) -> Result<Stmt, String> {
-        if self.eat_punct("{") {
-            return Ok(Stmt::Block(self.parse_block_body()?));
-        }
-        if self.eat_ident("if") {
-            self.expect_punct("(")?;
-            let cond = self.parse_expr()?;
-            self.expect_punct(")")?;
-            let then = vec![self.parse_stmt()?];
-            let els = if self.eat_ident("else") { vec![self.parse_stmt()?] } else { Vec::new() };
-            return Ok(Stmt::If(cond, then, els));
-        }
-        if self.eat_ident("for") {
-            self.expect_punct("(")?;
-            let init = if self.eat_punct(";") { None } else { Some(Box::new(self.parse_stmt()?)) };
-            let cond = if self.eat_punct(";") {
-                None
-            } else {
-                let c = self.parse_expr()?;
-                self.expect_punct(";")?;
-                Some(c)
-            };
-            let step = if self.eat_punct(")") {
-                None
-            } else {
-                let s = self.parse_expr()?;
-                self.expect_punct(")")?;
-                Some(s)
-            };
-            let body = vec![self.parse_stmt()?];
-            return Ok(Stmt::For(init, cond, step, body));
-        }
-        if self.eat_ident("while") {
-            self.expect_punct("(")?;
-            let cond = self.parse_expr()?;
-            self.expect_punct(")")?;
-            let body = vec![self.parse_stmt()?];
-            return Ok(Stmt::While(cond, body));
-        }
-        if self.eat_ident("return") {
-            if self.eat_punct(";") {
-                return Ok(Stmt::Return(None));
-            }
-            let e = self.parse_expr()?;
-            self.expect_punct(";")?;
-            return Ok(Stmt::Return(Some(e)));
-        }
-        if self.eat_ident("break") {
-            self.expect_punct(";")?;
-            return Ok(Stmt::Break);
-        }
-        if self.eat_ident("continue") {
-            self.expect_punct(";")?;
-            return Ok(Stmt::Continue);
-        }
-        self.skip_quals();
-        if let Some(ty) = self.peek_ty() {
-            // Declaration — possibly several declarators (`float a = 1, b;`).
-            self.pos += 1;
-            let mut decls = Vec::new();
-            loop {
-                self.skip_quals();
-                let name = match self.next() {
-                    Some(Tok::Ident(n)) => n,
-                    t => return Err(format!("kernel_cpu: expected variable name, got {t:?}")),
-                };
-                let mut arr = None;
-                if self.eat_punct("[") {
-                    match self.next() {
-                        Some(Tok::Int(n)) => arr = Some(n as usize),
-                        t => return Err(format!("kernel_cpu: expected array length, got {t:?}")),
-                    }
-                    self.expect_punct("]")?;
-                }
-                let mut init = Vec::new();
-                if self.eat_punct("=") {
-                    if self.eat_punct("{") {
-                        if !self.eat_punct("}") {
-                            loop {
-                                init.push(self.parse_assign()?);
-                                if !self.eat_punct(",") {
-                                    break;
-                                }
-                            }
-                            self.expect_punct("}")?;
-                        }
-                    } else {
-                        init.push(self.parse_assign()?);
-                    }
-                }
-                decls.push(Stmt::Decl { ty, name, arr, init });
-                if !self.eat_punct(",") {
-                    break;
-                }
-            }
-            self.expect_punct(";")?;
-            return Ok(if decls.len() == 1 { decls.pop().unwrap() } else { Stmt::Seq(decls) });
-        }
-        let e = self.parse_expr()?;
-        self.expect_punct(";")?;
-        Ok(Stmt::Expr(e))
-    }
-
-    // Expression grammar: comma-free C precedence.
-    fn parse_expr(&mut self) -> Result<Expr, String> {
-        self.parse_assign()
-    }
-    fn parse_assign(&mut self) -> Result<Expr, String> {
-        let lhs = self.parse_ternary()?;
-        for op in ["=", "+=", "-=", "*=", "/=", "%="] {
-            if matches!(self.peek(), Some(Tok::Punct(p)) if *p == op) {
-                self.pos += 1;
-                let rhs = self.parse_assign()?;
-                let sop: &'static str = match op {
-                    "=" => "=",
-                    "+=" => "+=",
-                    "-=" => "-=",
-                    "*=" => "*=",
-                    "/=" => "/=",
-                    _ => "%=",
-                };
-                return Ok(Expr::Assign(sop, Box::new(lhs), Box::new(rhs)));
-            }
-        }
-        Ok(lhs)
-    }
-    fn parse_ternary(&mut self) -> Result<Expr, String> {
-        let cond = self.parse_bin(0)?;
-        if self.eat_punct("?") {
-            let a = self.parse_assign()?;
-            self.expect_punct(":")?;
-            let b = self.parse_assign()?;
-            return Ok(Expr::Ternary(Box::new(cond), Box::new(a), Box::new(b)));
-        }
-        Ok(cond)
-    }
-    fn parse_bin(&mut self, min_prec: u8) -> Result<Expr, String> {
-        let mut lhs = self.parse_unary()?;
-        loop {
-            let (op, prec): (&'static str, u8) = match self.peek() {
-                Some(Tok::Punct(p)) => match *p {
-                    "||" => ("||", 1),
-                    "&&" => ("&&", 2),
-                    "==" => ("==", 3),
-                    "!=" => ("!=", 3),
-                    "<" => ("<", 4),
-                    ">" => (">", 4),
-                    "<=" => ("<=", 4),
-                    ">=" => (">=", 4),
-                    "+" => ("+", 5),
-                    "-" => ("-", 5),
-                    "*" => ("*", 6),
-                    "/" => ("/", 6),
-                    "%" => ("%", 6),
-                    _ => break,
-                },
-                _ => break,
-            };
-            if prec < min_prec {
-                break;
-            }
-            self.pos += 1;
-            let rhs = self.parse_bin(prec + 1)?;
-            lhs = Expr::Bin(op, Box::new(lhs), Box::new(rhs));
-        }
-        Ok(lhs)
-    }
-    fn parse_unary(&mut self) -> Result<Expr, String> {
-        // Cast: '(' type ')' unary — only for our two value types.
-        if matches!(self.peek(), Some(Tok::Punct("("))) {
-            if let Some(Tok::Ident(s)) = self.toks.get(self.pos + 1) {
-                let ty = match s.as_str() {
-                    "float" => Some(Ty::Float),
-                    "int" => Some(Ty::Int),
-                    _ => None,
-                };
-                if ty.is_some() && matches!(self.toks.get(self.pos + 2), Some(Tok::Punct(")"))) {
-                    self.pos += 3;
-                    let e = self.parse_unary()?;
-                    return Ok(Expr::Cast(ty.unwrap(), Box::new(e)));
-                }
-            }
-        }
-        if self.eat_punct("-") {
-            return Ok(Expr::Unary("-", Box::new(self.parse_unary()?)));
-        }
-        if self.eat_punct("!") {
-            return Ok(Expr::Unary("!", Box::new(self.parse_unary()?)));
-        }
-        if self.eat_punct("+") {
-            return self.parse_unary();
-        }
-        if self.eat_punct("&") {
-            return Ok(Expr::AddrOf(Box::new(self.parse_unary()?)));
-        }
-        if self.eat_punct("*") {
-            return Ok(Expr::Deref(Box::new(self.parse_unary()?)));
-        }
-        if self.eat_punct("++") {
-            // Prefix inc as `x += 1`.
-            let e = self.parse_unary()?;
-            return Ok(Expr::Assign("+=", Box::new(e), Box::new(Expr::I(1))));
-        }
-        if self.eat_punct("--") {
-            let e = self.parse_unary()?;
-            return Ok(Expr::Assign("-=", Box::new(e), Box::new(Expr::I(1))));
-        }
-        self.parse_postfix()
-    }
-    fn parse_postfix(&mut self) -> Result<Expr, String> {
-        let mut e = self.parse_primary()?;
-        loop {
-            if self.eat_punct("[") {
-                let idx = self.parse_expr()?;
-                self.expect_punct("]")?;
-                e = Expr::Index(Box::new(e), Box::new(idx));
-            } else if matches!(self.peek(), Some(Tok::Punct("++"))) {
-                self.pos += 1;
-                e = Expr::Post("++", Box::new(e));
-            } else if matches!(self.peek(), Some(Tok::Punct("--"))) {
-                self.pos += 1;
-                e = Expr::Post("--", Box::new(e));
-            } else {
-                break;
-            }
-        }
-        Ok(e)
-    }
-    fn parse_primary(&mut self) -> Result<Expr, String> {
-        match self.next() {
-            Some(Tok::Int(n)) => Ok(Expr::I(n)),
-            Some(Tok::Float(f)) => Ok(Expr::F(f)),
-            Some(Tok::Punct("(")) => {
-                let e = self.parse_expr()?;
-                self.expect_punct(")")?;
-                Ok(e)
-            }
-            Some(Tok::Ident(name)) => {
-                if name == "true" {
-                    return Ok(Expr::I(1));
-                }
-                if name == "false" {
-                    return Ok(Expr::I(0));
-                }
-                if self.eat_punct("(") {
-                    let mut args = Vec::new();
-                    if !self.eat_punct(")") {
-                        loop {
-                            args.push(self.parse_assign()?);
-                            if !self.eat_punct(",") {
-                                break;
-                            }
-                        }
-                        self.expect_punct(")")?;
-                    }
-                    Ok(Expr::Call(name, args))
-                } else {
-                    Ok(Expr::Var(name))
-                }
-            }
-            t => Err(format!("kernel_cpu: unexpected token {t:?} in expression")),
-        }
-    }
-}
-
-// ---------------------------------------------------------------- interpreter
-
-#[derive(Debug, Clone, Copy, PartialEq)]
-enum BufId {
-    InPos,
-    InCol,
-    OutPos,
-    OutCol,
-    OutCount,
-    Params,
-    RwPos,
-    RwCol,
-    /// One of the kernel's named attribute buffers, by binding slot — the
-    /// Phase 1 ABI's addition. See `geometry::parse_attr_refs`.
-    Attr(usize),
-}
-
-#[derive(Debug, Clone, Copy, PartialEq)]
-enum PtrV {
-    Slot(usize),
-    Buf(BufId),
-}
-
-#[derive(Debug, Clone, Copy, PartialEq)]
-enum Val {
-    F(f32),
-    I(i64),
-    P(PtrV),
-}
-
-impl Val {
-    fn as_f(self) -> f32 {
-        match self {
-            Val::F(f) => f,
-            Val::I(i) => i as f32,
-            Val::P(_) => 0.0,
-        }
-    }
-    fn as_i(self) -> i64 {
-        match self {
-            Val::F(f) => f as i64, // C float→int truncates toward zero, as `as` does
-            Val::I(i) => i,
-            Val::P(_) => 0,
-        }
-    }
-    fn truthy(self) -> bool {
-        match self {
-            Val::F(f) => f != 0.0,
-            Val::I(i) => i != 0,
-            Val::P(_) => true,
-        }
-    }
-}
-
-#[derive(Debug, Clone)]
-enum SlotData {
-    F(f32),
-    I(i64),
-    ArrF(Vec<f32>),
-    ArrI(Vec<i64>),
-    Ptr(PtrV),
-}
-
-/// An assignable location, resolved before the store.
-enum Place {
-    Slot(usize),
-    SlotElem(usize, usize),
-    BufElem(BufId, usize),
-}
-
-enum Flow {
-    Normal,
-    Break,
-    Continue,
-    Return(Option<Val>),
-}
-
-struct Bufs<'a> {
-    in_pos: &'a [f32],
-    in_col: &'a [f32],
-    out_pos: Vec<f32>,
-    out_col: Vec<f32>,
-    out_count: i64,
-    rw_pos: Vec<f32>,
-    rw_col: Vec<f32>,
-    params: &'a [f32],
-    attrs: Vec<Vec<f32>>,
-}
-
-struct Interp<'a> {
-    fns: &'a HashMap<String, FnDef>,
-    slots: Vec<SlotData>,
-    scopes: Vec<HashMap<String, usize>>,
-    bufs: Bufs<'a>,
-    global_id: usize,
-    steps: u64,
-    budget: u64,
-}
-
-impl<'a> Interp<'a> {
-    fn step(&mut self) -> Result<(), String> {
-        self.steps += 1;
-        if self.steps > self.budget {
-            Err("kernel step budget exceeded (non-terminating kernel?)".into())
-        } else {
-            Ok(())
-        }
-    }
-
-    fn push_scope(&mut self) {
-        self.scopes.push(HashMap::new());
-    }
-    fn pop_scope(&mut self) {
-        self.scopes.pop();
-    }
-    fn declare(&mut self, name: &str, data: SlotData) -> usize {
-        let idx = self.slots.len();
-        self.slots.push(data);
-        self.scopes.last_mut().unwrap().insert(name.to_string(), idx);
-        idx
-    }
-    fn lookup(&self, name: &str) -> Option<usize> {
-        for scope in self.scopes.iter().rev() {
-            if let Some(&i) = scope.get(name) {
-                return Some(i);
-            }
-        }
-        None
-    }
-
-    fn buf_read(&self, id: BufId, idx: usize) -> Val {
-        match id {
-            BufId::InPos => Val::F(self.bufs.in_pos.get(idx).copied().unwrap_or(0.0)),
-            BufId::InCol => Val::F(self.bufs.in_col.get(idx).copied().unwrap_or(0.0)),
-            BufId::OutPos => Val::F(self.bufs.out_pos.get(idx).copied().unwrap_or(0.0)),
-            BufId::OutCol => Val::F(self.bufs.out_col.get(idx).copied().unwrap_or(0.0)),
-            BufId::OutCount => Val::I(if idx == 0 { self.bufs.out_count } else { 0 }),
-            BufId::Params => Val::F(self.bufs.params.get(idx).copied().unwrap_or(0.0)),
-            BufId::RwPos => Val::F(self.bufs.rw_pos.get(idx).copied().unwrap_or(0.0)),
-            BufId::RwCol => Val::F(self.bufs.rw_col.get(idx).copied().unwrap_or(0.0)),
-            BufId::Attr(slot) => Val::F(
-                self.bufs
-                    .attrs
-                    .get(slot)
-                    .and_then(|a| a.get(idx))
-                    .copied()
-                    .unwrap_or(0.0),
-            ),
-        }
-    }
-    fn buf_write(&mut self, id: BufId, idx: usize, v: Val) {
-        // OOB writes are dropped — the GPU's UB made safe.
-        match id {
-            BufId::OutPos => {
-                if let Some(slot) = self.bufs.out_pos.get_mut(idx) {
-                    *slot = v.as_f();
-                }
-            }
-            BufId::OutCol => {
-                if let Some(slot) = self.bufs.out_col.get_mut(idx) {
-                    *slot = v.as_f();
-                }
-            }
-            BufId::OutCount => {
-                if idx == 0 {
-                    self.bufs.out_count = v.as_i();
-                }
-            }
-            BufId::RwPos => {
-                if let Some(slot) = self.bufs.rw_pos.get_mut(idx) {
-                    *slot = v.as_f();
-                }
-            }
-            BufId::RwCol => {
-                if let Some(slot) = self.bufs.rw_col.get_mut(idx) {
-                    *slot = v.as_f();
-                }
-            }
-            BufId::Attr(slot) => {
-                if let Some(slot) = self.bufs.attrs.get_mut(slot).and_then(|a| a.get_mut(idx)) {
-                    *slot = v.as_f();
-                }
-            }
-            BufId::InPos | BufId::InCol | BufId::Params => {}
-        }
-    }
-
-    fn place_read(&self, p: &Place) -> Val {
-        match p {
-            Place::Slot(i) => match &self.slots[*i] {
-                SlotData::F(f) => Val::F(*f),
-                SlotData::I(n) => Val::I(*n),
-                SlotData::Ptr(p) => Val::P(*p),
-                SlotData::ArrF(_) | SlotData::ArrI(_) => Val::I(0),
-            },
-            Place::SlotElem(i, k) => match &self.slots[*i] {
-                SlotData::ArrF(v) => Val::F(v.get(*k).copied().unwrap_or(0.0)),
-                SlotData::ArrI(v) => Val::I(v.get(*k).copied().unwrap_or(0)),
-                _ => Val::I(0),
-            },
-            Place::BufElem(id, k) => self.buf_read(*id, *k),
-        }
-    }
-    fn place_write(&mut self, p: &Place, v: Val) {
-        match p {
-            Place::Slot(i) => {
-                let slot = &mut self.slots[*i];
-                match slot {
-                    SlotData::F(f) => *f = v.as_f(),
-                    SlotData::I(n) => *n = v.as_i(),
-                    SlotData::Ptr(q) => {
-                        if let Val::P(np) = v {
-                            *q = np;
-                        }
-                    }
-                    SlotData::ArrF(_) | SlotData::ArrI(_) => {}
-                }
-            }
-            Place::SlotElem(i, k) => {
-                let slot = &mut self.slots[*i];
-                match slot {
-                    SlotData::ArrF(vec) => {
-                        if let Some(e) = vec.get_mut(*k) {
-                            *e = v.as_f();
-                        }
-                    }
-                    SlotData::ArrI(vec) => {
-                        if let Some(e) = vec.get_mut(*k) {
-                            *e = v.as_i();
-                        }
-                    }
-                    _ => {}
-                }
-            }
-            Place::BufElem(id, k) => self.buf_write(*id, *k, v),
-        }
-    }
-
-    /// Resolve an lvalue expression to a storage location.
-    fn resolve_place(&mut self, e: &Expr) -> Result<Place, String> {
-        match e {
-            Expr::Var(name) => {
-                let idx = self.lookup(name).ok_or_else(|| format!("kernel_cpu: unknown variable '{name}'"))?;
-                Ok(Place::Slot(idx))
-            }
-            Expr::Index(base, idx) => {
-                let k = self.eval(idx)?.as_i();
-                if k < 0 {
-                    return Ok(Place::BufElem(BufId::Params, usize::MAX)); // negative index: dead place
-                }
-                let k = k as usize;
-                match &**base {
-                    Expr::Var(name) => {
-                        let slot_idx = self.lookup(name).ok_or_else(|| format!("kernel_cpu: unknown variable '{name}'"))?;
-                        match &self.slots[slot_idx] {
-                            SlotData::Ptr(PtrV::Buf(id)) => Ok(Place::BufElem(*id, k)),
-                            SlotData::Ptr(PtrV::Slot(s)) => Ok(Place::SlotElem(*s, k)),
-                            SlotData::ArrF(_) | SlotData::ArrI(_) => Ok(Place::SlotElem(slot_idx, k)),
-                            _ => Err(format!("kernel_cpu: '{name}' is not indexable")),
-                        }
-                    }
-                    other => {
-                        // e.g. (*ptr)[k] — evaluate to a pointer and index it.
-                        let v = self.eval(other)?;
-                        match v {
-                            Val::P(PtrV::Buf(id)) => Ok(Place::BufElem(id, k)),
-                            Val::P(PtrV::Slot(s)) => Ok(Place::SlotElem(s, k)),
-                            _ => Err("kernel_cpu: indexing a non-pointer expression".into()),
-                        }
-                    }
-                }
-            }
-            Expr::Deref(inner) => {
-                let v = self.eval(inner)?;
-                match v {
-                    Val::P(PtrV::Buf(id)) => Ok(Place::BufElem(id, 0)),
-                    Val::P(PtrV::Slot(s)) => Ok(Place::Slot(s)),
-                    _ => Err("kernel_cpu: dereferencing a non-pointer".into()),
-                }
-            }
-            _ => Err("kernel_cpu: expression is not assignable".into()),
-        }
-    }
-
-    fn eval(&mut self, e: &Expr) -> Result<Val, String> {
-        self.step()?;
-        match e {
-            Expr::F(f) => Ok(Val::F(*f)),
-            Expr::I(n) => Ok(Val::I(*n)),
-            Expr::Var(name) => {
-                let idx = self.lookup(name).ok_or_else(|| format!("kernel_cpu: unknown variable '{name}'"))?;
-                Ok(self.place_read(&Place::Slot(idx)))
-            }
-            Expr::Index(..) | Expr::Deref(..) => {
-                let p = self.resolve_place(e)?;
-                Ok(self.place_read(&p))
-            }
-            Expr::AddrOf(inner) => match self.resolve_place(inner)? {
-                Place::Slot(i) => Ok(Val::P(PtrV::Slot(i))),
-                Place::BufElem(id, 0) => Ok(Val::P(PtrV::Buf(id))),
-                _ => Err("kernel_cpu: unsupported address-of".into()),
-            },
-            Expr::Unary(op, inner) => {
-                let v = self.eval(inner)?;
-                match (*op, v) {
-                    ("-", Val::F(f)) => Ok(Val::F(-f)),
-                    ("-", Val::I(n)) => Ok(Val::I(n.wrapping_neg())),
-                    ("!", v) => Ok(Val::I(if v.truthy() { 0 } else { 1 })),
-                    _ => Err(format!("kernel_cpu: bad unary {op}")),
-                }
-            }
-            Expr::Post(op, inner) => {
-                let p = self.resolve_place(inner)?;
-                let old = self.place_read(&p);
-                let one = Val::I(1);
-                let new = bin_arith(if *op == "++" { "+" } else { "-" }, old, one)?;
-                self.place_write(&p, new);
-                Ok(old)
-            }
-            Expr::Bin(op, a, b) => match *op {
-                "&&" => {
-                    let va = self.eval(a)?;
-                    if !va.truthy() {
-                        return Ok(Val::I(0));
-                    }
-                    Ok(Val::I(if self.eval(b)?.truthy() { 1 } else { 0 }))
-                }
-                "||" => {
-                    let va = self.eval(a)?;
-                    if va.truthy() {
-                        return Ok(Val::I(1));
-                    }
-                    Ok(Val::I(if self.eval(b)?.truthy() { 1 } else { 0 }))
-                }
-                _ => {
-                    let va = self.eval(a)?;
-                    let vb = self.eval(b)?;
-                    bin_arith(op, va, vb)
-                }
-            },
-            Expr::Assign(op, lhs, rhs) => {
-                let rv = self.eval(rhs)?;
-                let p = self.resolve_place(lhs)?;
-                let out = if *op == "=" {
-                    rv
-                } else {
-                    let cur = self.place_read(&p);
-                    let bop = &op[..1]; // "+=" -> "+"
-                    bin_arith(bop, cur, rv)?
-                };
-                self.place_write(&p, out);
-                Ok(self.place_read(&p))
-            }
-            Expr::Ternary(c, a, b) => {
-                if self.eval(c)?.truthy() {
-                    self.eval(a)
-                } else {
-                    self.eval(b)
-                }
-            }
-            Expr::Cast(ty, inner) => {
-                let v = self.eval(inner)?;
-                Ok(match ty {
-                    Ty::Float => Val::F(v.as_f()),
-                    Ty::Int => Val::I(v.as_i()),
-                    Ty::Void => Val::I(0),
-                })
-            }
-            Expr::Call(name, args) => self.call(name, args),
-        }
-    }
-
-    fn call(&mut self, name: &str, args: &[Expr]) -> Result<Val, String> {
-        // Builtins first.
-        match name {
-            "get_global_id" => return Ok(Val::I(self.global_id as i64)),
-            "get_global_size" => return Ok(Val::I(1)),
-            _ => {}
-        }
-        if let Some(v) = self.try_math_builtin(name, args)? {
-            return Ok(v);
-        }
-        let def = self
-            .fns
-            .get(name)
-            .ok_or_else(|| format!("kernel_cpu: unknown function '{name}'"))?
-            .clone();
-        if def.params.len() != args.len() {
-            return Err(format!("kernel_cpu: {name} expects {} args, got {}", def.params.len(), args.len()));
-        }
-        let mut bound = Vec::with_capacity(args.len());
-        for (param, arg) in def.params.iter().zip(args) {
-            let v = if param.is_ptr {
-                // Pointer parameter: pass a pointer value; a bare buffer/array
-                // name decays to a pointer to it.
-                match arg {
-                    Expr::Var(n) => {
-                        let idx = self.lookup(n).ok_or_else(|| format!("kernel_cpu: unknown variable '{n}'"))?;
-                        match &self.slots[idx] {
-                            SlotData::Ptr(p) => Val::P(*p),
-                            SlotData::ArrF(_) | SlotData::ArrI(_) => Val::P(PtrV::Slot(idx)),
-                            _ => Val::P(PtrV::Slot(idx)),
-                        }
-                    }
-                    _ => self.eval(arg)?,
-                }
-            } else {
-                self.eval(arg)?
-            };
-            bound.push((param.name.clone(), param.is_ptr, v));
-        }
-        self.push_scope();
-        for (pname, is_ptr, v) in bound {
-            let data = if is_ptr {
-                match v {
-                    Val::P(p) => SlotData::Ptr(p),
-                    _ => return Err(format!("kernel_cpu: pointer argument '{pname}' is not a pointer")),
-                }
-            } else {
-                match v {
-                    Val::F(f) => SlotData::F(f),
-                    Val::I(n) => SlotData::I(n),
-                    Val::P(p) => SlotData::Ptr(p),
-                }
-            };
-            self.declare(&pname, data);
-        }
-        let flow = self.run_block(&def.body)?;
-        self.pop_scope();
-        match flow {
-            Flow::Return(Some(v)) => Ok(v),
-            _ => Ok(Val::I(0)),
-        }
-    }
-
-    fn try_math_builtin(&mut self, name: &str, args: &[Expr]) -> Result<Option<Val>, String> {
-        let f1 = |i: &mut Self, args: &[Expr]| -> Result<f32, String> { Ok(i.eval(&args[0])?.as_f()) };
-        let v = match (name, args.len()) {
-            ("sqrt", 1) => Val::F(f1(self, args)?.sqrt()),
-            ("sin", 1) => Val::F(f1(self, args)?.sin()),
-            ("cos", 1) => Val::F(f1(self, args)?.cos()),
-            ("tan", 1) => Val::F(f1(self, args)?.tan()),
-            ("fabs", 1) => Val::F(f1(self, args)?.abs()),
-            ("floor", 1) => Val::F(f1(self, args)?.floor()),
-            ("ceil", 1) => Val::F(f1(self, args)?.ceil()),
-            ("exp", 1) => Val::F(f1(self, args)?.exp()),
-            ("log", 1) => Val::F(f1(self, args)?.ln()),
-            ("abs", 1) => Val::I(self.eval(&args[0])?.as_i().wrapping_abs()),
-            ("fmod", 2) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                Val::F(a % b)
-            }
-            ("pow", 2) | ("powr", 2) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                Val::F(a.powf(b))
-            }
-            ("atan2", 2) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                Val::F(a.atan2(b))
-            }
-            ("fmin", 2) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                Val::F(a.min(b))
-            }
-            ("fmax", 2) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                Val::F(a.max(b))
-            }
-            ("min", 2) => {
-                let a = self.eval(&args[0])?;
-                let b = self.eval(&args[1])?;
-                match (a, b) {
-                    (Val::I(x), Val::I(y)) => Val::I(x.min(y)),
-                    _ => Val::F(a.as_f().min(b.as_f())),
-                }
-            }
-            ("max", 2) => {
-                let a = self.eval(&args[0])?;
-                let b = self.eval(&args[1])?;
-                match (a, b) {
-                    (Val::I(x), Val::I(y)) => Val::I(x.max(y)),
-                    _ => Val::F(a.as_f().max(b.as_f())),
-                }
-            }
-            ("clamp", 3) => {
-                let x = self.eval(&args[0])?.as_f();
-                let lo = self.eval(&args[1])?.as_f();
-                let hi = self.eval(&args[2])?.as_f();
-                Val::F(x.clamp(lo, hi))
-            }
-            ("mix", 3) => {
-                let a = self.eval(&args[0])?.as_f();
-                let b = self.eval(&args[1])?.as_f();
-                let t = self.eval(&args[2])?.as_f();
-                Val::F(a + (b - a) * t)
-            }
-            _ => return Ok(None),
-        };
-        Ok(Some(v))
-    }
-
-    fn run_block(&mut self, stmts: &[Stmt]) -> Result<Flow, String> {
-        self.push_scope();
-        let mut flow = Flow::Normal;
-        for s in stmts {
-            flow = self.run_stmt(s)?;
-            if !matches!(flow, Flow::Normal) {
-                break;
-            }
-        }
-        self.pop_scope();
-        Ok(flow)
-    }
-
-    fn run_stmt(&mut self, s: &Stmt) -> Result<Flow, String> {
-        self.step()?;
-        match s {
-            Stmt::Block(body) => self.run_block(body),
-            Stmt::Seq(stmts) => {
-                for st in stmts {
-                    match self.run_stmt(st)? {
-                        Flow::Normal => {}
-                        f => return Ok(f),
-                    }
-                }
-                Ok(Flow::Normal)
-            }
-            Stmt::Expr(e) => {
-                self.eval(e)?;
-                Ok(Flow::Normal)
-            }
-            Stmt::Return(e) => {
-                let v = match e {
-                    Some(e) => Some(self.eval(e)?),
-                    None => None,
-                };
-                Ok(Flow::Return(v))
-            }
-            Stmt::Break => Ok(Flow::Break),
-            Stmt::Continue => Ok(Flow::Continue),
-            Stmt::If(c, t, f) => {
-                if self.eval(c)?.truthy() {
-                    self.run_block(t)
-                } else {
-                    self.run_block(f)
-                }
-            }
-            Stmt::While(c, body) => {
-                loop {
-                    if !self.eval(c)?.truthy() {
-                        break;
-                    }
-                    match self.run_block(body)? {
-                        Flow::Break => break,
-                        Flow::Return(v) => return Ok(Flow::Return(v)),
-                        _ => {}
-                    }
-                }
-                Ok(Flow::Normal)
-            }
-            Stmt::For(init, cond, step, body) => {
-                self.push_scope();
-                if let Some(init) = init {
-                    self.run_stmt(init)?;
-                }
-                loop {
-                    if let Some(c) = cond {
-                        if !self.eval(c)?.truthy() {
-                            break;
-                        }
-                    }
-                    match self.run_block(body)? {
-                        Flow::Break => break,
-                        Flow::Return(v) => {
-                            self.pop_scope();
-                            return Ok(Flow::Return(v));
-                        }
-                        _ => {}
-                    }
-                    if let Some(st) = step {
-                        self.eval(st)?;
-                    }
-                }
-                self.pop_scope();
-                Ok(Flow::Normal)
-            }
-            Stmt::Decl { ty, name, arr, init } => {
-                let data = match (arr, ty) {
-                    (Some(n), Ty::Float) => {
-                        let mut v = vec![0.0f32; *n];
-                        for (i, e) in init.iter().enumerate().take(*n) {
-                            v[i] = self.eval(e)?.as_f();
-                        }
-                        SlotData::ArrF(v)
-                    }
-                    (Some(n), _) => {
-                        let mut v = vec![0i64; *n];
-                        for (i, e) in init.iter().enumerate().take(*n) {
-                            v[i] = self.eval(e)?.as_i();
-                        }
-                        SlotData::ArrI(v)
-                    }
-                    (None, Ty::Float) => {
-                        let v = init.first().map(|e| self.eval(e)).transpose()?.map(|v| v.as_f()).unwrap_or(0.0);
-                        SlotData::F(v)
-                    }
-                    (None, _) => match init.first().map(|e| self.eval(e)).transpose()? {
-                        Some(Val::P(p)) => SlotData::Ptr(p),
-                        Some(v) => SlotData::I(v.as_i()),
-                        None => SlotData::I(0),
-                    },
-                };
-                self.declare(name, data);
-                Ok(Flow::Normal)
-            }
-        }
-    }
-}
-
-/// C arithmetic: int op int stays int (truncating division, wrapping i32-ish),
-/// anything touching a float promotes to f32; comparisons yield int 0/1.
-fn bin_arith(op: &str, a: Val, b: Val) -> Result<Val, String> {
-    let both_int = matches!(a, Val::I(_)) && matches!(b, Val::I(_));
-    Ok(match op {
-        "+" | "-" | "*" | "/" | "%" => {
-            if both_int {
-                let (x, y) = (a.as_i(), b.as_i());
-                let r = match op {
-                    "+" => x.wrapping_add(y),
-                    "-" => x.wrapping_sub(y),
-                    "*" => x.wrapping_mul(y),
-                    "/" => {
-                        if y == 0 {
-                            0
-                        } else {
-                            x.wrapping_div(y)
-                        }
-                    }
-                    _ => {
-                        if y == 0 {
-                            0
-                        } else {
-                            x.wrapping_rem(y)
-                        }
-                    }
-                };
-                Val::I(r)
-            } else {
-                let (x, y) = (a.as_f(), b.as_f());
-                let r = match op {
-                    "+" => x + y,
-                    "-" => x - y,
-                    "*" => x * y,
-                    "/" => x / y,
-                    _ => x % y,
-                };
-                Val::F(r)
-            }
-        }
-        "<" | ">" | "<=" | ">=" | "==" | "!=" => {
-            let t = if both_int {
-                let (x, y) = (a.as_i(), b.as_i());
-                match op {
-                    "<" => x < y,
-                    ">" => x > y,
-                    "<=" => x <= y,
-                    ">=" => x >= y,
-                    "==" => x == y,
-                    _ => x != y,
-                }
-            } else {
-                let (x, y) = (a.as_f(), b.as_f());
-                match op {
-                    "<" => x < y,
-                    ">" => x > y,
-                    "<=" => x <= y,
-                    ">=" => x >= y,
-                    "==" => x == y,
-                    _ => x != y,
-                }
-            };
-            Val::I(if t { 1 } else { 0 })
-        }
-        _ => return Err(format!("kernel_cpu: unknown operator {op}")),
-    })
-}
-
-// ---------------------------------------------------------------- launcher
-
-/// CPU twin of `run_opencl_kernel_with_params`: same generator/deformer split
-/// (a kernel that mentions `out_count` is a generator), same positional
-/// argument binding, same `max_vertices`, same output rebuild with the default
-/// Norm/UV attributes. The `process` entry point runs once per work item with
-/// `get_global_id(0)` = the item index, exactly as the ND-range launch does.
-/// CPU twin of `geometry::run_deformer_flat`: same flat buffers, same binding
-/// order, same read-back in place.
-///
-/// The two backends share the kernel language, so they have to share the ABI
-/// too — `cpu_matches_opencl_on_every_shipped_kernel` is only meaningful while
-/// they bind the same arguments to the same slots.
-pub fn run_deformer_cpu(
-    code: &str,
-    pos: &mut [f32],
-    col: &mut [f32],
-    count: usize,
-    attrs: &mut [(String, Vec<f32>)],
-    params: &[f32],
-) -> Result<(), String> {
-    if count == 0 {
-        return Ok(());
-    }
-    let toks = lex(code)?;
-    let mut parser = Parser { toks, pos: 0 };
-    let fns = parser.parse_program()?;
-    let process = fns.get("process").ok_or("kernel_cpu: no 'process' kernel")?;
-    let arity = process.params.len();
-    let pnames: Vec<(String, bool)> = process
-        .params
-        .iter()
-        .map(|p| (p.name.clone(), p.is_ptr))
-        .collect();
-
-    let mut param_data = params.to_vec();
-    if param_data.is_empty() {
-        param_data.push(0.0);
-    }
-
-    // Binding order mirrors the OpenCL launcher exactly: the three fixed
-    // arguments, `param_values` only when the arity says the kernel declared
-    // it, then one buffer per named attribute.
-    let mut binds: Vec<Val> = vec![
-        Val::P(PtrV::Buf(BufId::RwPos)),
-        Val::P(PtrV::Buf(BufId::RwCol)),
-        Val::I(count as i64),
-    ];
-    if arity > 3 + attrs.len() {
-        binds.push(Val::P(PtrV::Buf(BufId::Params)));
-    }
-    for slot in 0..attrs.len() {
-        binds.push(Val::P(PtrV::Buf(BufId::Attr(slot))));
-    }
-    let n_bind = binds.len().min(arity);
-
-    let mut bufs = Bufs {
-        in_pos: &[],
-        in_col: &[],
-        out_pos: Vec::new(),
-        out_col: Vec::new(),
-        out_count: 0,
-        rw_pos: pos.to_vec(),
-        rw_col: col.to_vec(),
-        params: &param_data,
-        attrs: attrs.iter().map(|(_, v)| v.clone()).collect(),
-    };
-
-    let mut steps = 0u64;
-    for id in 0..count {
-        let mut interp = Interp {
-            fns: &fns,
-            slots: Vec::new(),
-            scopes: vec![HashMap::new()],
-            bufs: std::mem::replace(
-                &mut bufs,
-                Bufs { in_pos: &[], in_col: &[], out_pos: Vec::new(), out_col: Vec::new(), out_count: 0, rw_pos: Vec::new(), rw_col: Vec::new(), params: &[], attrs: Vec::new() },
-            ),
-            global_id: id,
-            steps,
-            budget: STEP_BUDGET,
-        };
-        for (i, (name, is_ptr)) in pnames.iter().enumerate().take(n_bind) {
-            let data = match binds[i] {
-                Val::P(p) if *is_ptr => SlotData::Ptr(p),
-                Val::I(n) => SlotData::I(n),
-                Val::F(f) => SlotData::F(f),
-                Val::P(p) => SlotData::Ptr(p),
-            };
-            interp.declare(name, data);
-        }
-        interp.run_block(&process.body)?;
-        steps = interp.steps;
-        bufs = interp.bufs;
-    }
-
-    pos.copy_from_slice(&bufs.rw_pos);
-    col.copy_from_slice(&bufs.rw_col);
-    for ((_, dst), src) in attrs.iter_mut().zip(bufs.attrs.into_iter()) {
-        *dst = src;
-    }
-    Ok(())
-}
-
-pub fn run_kernel_cpu(code: &str, geom: &mut Geometry, params: &[f32]) -> Result<(), String> {
-    run_kernel_cpu_with_budget(code, geom, params, STEP_BUDGET)
-}
-
-fn run_kernel_cpu_with_budget(code: &str, geom: &mut Geometry, params: &[f32], budget: u64) -> Result<(), String> {
-    let is_generator = code.contains("out_count");
-    if geom.vertices.is_empty() && !is_generator {
-        return Ok(());
-    }
-
-    let toks = lex(code)?;
-    let mut parser = Parser { toks, pos: 0 };
-    let fns = parser.parse_program()?;
-    let process = fns.get("process").ok_or("kernel_cpu: no 'process' kernel")?;
-    let arity = process.params.len();
-    let pnames: Vec<(String, bool)> = process.params.iter().map(|p| (p.name.clone(), p.is_ptr)).collect();
-
-    let mut param_data = params.to_vec();
-    if param_data.is_empty() {
-        param_data.push(0.0);
-    }
-
-    let mut in_pos = Vec::new();
-    let mut in_col = Vec::new();
-    for v in &geom.vertices {
-        in_pos.extend_from_slice(&v.pos);
-        in_col.extend_from_slice(&v.col);
-    }
-    let count = geom.vertices.len();
-
-    if is_generator {
-        // Positional binding mirrors the OpenCL set_arg order; param_values is
-        // bound only when the kernel declares the extra argument, exactly like
-        // the launcher's num_args >= 8 check.
-        let binds: Vec<Val> = vec![
-            Val::P(PtrV::Buf(BufId::InPos)),
-            Val::P(PtrV::Buf(BufId::InCol)),
-            Val::I(count as i64),
-            Val::P(PtrV::Buf(BufId::OutPos)),
-            Val::P(PtrV::Buf(BufId::OutCol)),
-            Val::P(PtrV::Buf(BufId::OutCount)),
-            Val::I(MAX_VERTICES as i64),
-            Val::P(PtrV::Buf(BufId::Params)),
-        ];
-        let n_bind = if arity >= 8 { 8 } else { 7.min(arity) };
-
-        let mut bufs = Bufs {
-            in_pos: &in_pos,
-            in_col: &in_col,
-            out_pos: vec![0.0; MAX_VERTICES * 3],
-            out_col: vec![0.0; MAX_VERTICES * 3],
-            out_count: 0,
-            rw_pos: Vec::new(),
-            rw_col: Vec::new(),
-            params: &param_data,
-            attrs: Vec::new(),
-        };
-
-        let global = count.max(1);
-        let mut steps = 0u64;
-        for id in 0..global {
-            let mut interp = Interp {
-                fns: &fns,
-                slots: Vec::new(),
-                scopes: vec![HashMap::new()],
-                bufs: std::mem::replace(
-                    &mut bufs,
-                    Bufs { in_pos: &[], in_col: &[], out_pos: Vec::new(), out_col: Vec::new(), out_count: 0, rw_pos: Vec::new(), rw_col: Vec::new(), params: &[], attrs: Vec::new() },
-                ),
-                global_id: id,
-                steps,
-                budget,
-            };
-            for (i, (name, is_ptr)) in pnames.iter().enumerate().take(n_bind) {
-                let data = match binds[i] {
-                    Val::P(p) if *is_ptr => SlotData::Ptr(p),
-                    Val::I(n) => SlotData::I(n),
-                    Val::F(f) => SlotData::F(f),
-                    Val::P(p) => SlotData::Ptr(p),
-                };
-                interp.declare(name, data);
-            }
-            let flow = interp.run_block(&process.body)?;
-            let _ = flow;
-            steps = interp.steps;
-            bufs = interp.bufs;
-        }
-
-        let final_count = (bufs.out_count.max(0) as usize).min(MAX_VERTICES);
-        geom.vertices.clear();
-        for i in 0..final_count {
-            let mut attributes = HashMap::new();
-            attributes.insert("Norm".to_string(), GAttribute::Float3([0.0, 1.0, 0.0]));
-            attributes.insert("UV".to_string(), GAttribute::Float2([0.0, 0.0]));
-            geom.vertices.push(GVertex {
-                pos: [bufs.out_pos[i * 3], bufs.out_pos[i * 3 + 1], bufs.out_pos[i * 3 + 2]],
-                col: [bufs.out_col[i * 3], bufs.out_col[i * 3 + 1], bufs.out_col[i * 3 + 2]],
-                attributes,
-            });
-        }
-    } else {
-        let binds: Vec<Val> = vec![
-            Val::P(PtrV::Buf(BufId::RwPos)),
-            Val::P(PtrV::Buf(BufId::RwCol)),
-            Val::I(count as i64),
-            Val::P(PtrV::Buf(BufId::Params)),
-        ];
-        let n_bind = if arity >= 4 { 4 } else { 3.min(arity) };
-
-        let mut bufs = Bufs {
-            in_pos: &[],
-            in_col: &[],
-            out_pos: Vec::new(),
-            out_col: Vec::new(),
-            out_count: 0,
-            rw_pos: in_pos.clone(),
-            rw_col: in_col.clone(),
-            params: &param_data,
-            attrs: Vec::new(),
-        };
-
-        let mut steps = 0u64;
-        for id in 0..count {
-            let mut interp = Interp {
-                fns: &fns,
-                slots: Vec::new(),
-                scopes: vec![HashMap::new()],
-                bufs: std::mem::replace(
-                    &mut bufs,
-                    Bufs { in_pos: &[], in_col: &[], out_pos: Vec::new(), out_col: Vec::new(), out_count: 0, rw_pos: Vec::new(), rw_col: Vec::new(), params: &[], attrs: Vec::new() },
-                ),
-                global_id: id,
-                steps,
-                budget,
-            };
-            for (i, (name, is_ptr)) in pnames.iter().enumerate().take(n_bind) {
-                let data = match binds[i] {
-                    Val::P(p) if *is_ptr => SlotData::Ptr(p),
-                    Val::I(n) => SlotData::I(n),
-                    Val::F(f) => SlotData::F(f),
-                    Val::P(p) => SlotData::Ptr(p),
-                };
-                interp.declare(name, data);
-            }
-            interp.run_block(&process.body)?;
-            steps = interp.steps;
-            bufs = interp.bufs;
-        }
-
-        for (i, v) in geom.vertices.iter_mut().enumerate() {
-            v.pos = [bufs.rw_pos[i * 3], bufs.rw_pos[i * 3 + 1], bufs.rw_pos[i * 3 + 2]];
-            v.col = [bufs.rw_col[i * 3], bufs.rw_col[i * 3 + 1], bufs.rw_col[i * 3 + 2]];
-        }
-    }
-    Ok(())
-}
-
-/// Force the CPU backend regardless of OpenCL availability (`CCE_KERNEL_CPU=1`).
-pub fn forced() -> bool {
-    std::env::var("CCE_KERNEL_CPU").is_ok_and(|v| v != "0" && !v.is_empty())
-}
-
-#[cfg(test)]
-mod tests {
-    use super::*;
-
-    fn gen(code: &str, params: &[f32]) -> Geometry {
-        let mut g = Geometry::new();
-        run_kernel_cpu(code, &mut g, params).expect("kernel runs");
-        g
-    }
-
-    /// The C corners a naive evaluator gets wrong, in one kernel: postfix ++
-    /// yielding the OLD value, int division truncating, casts, ternary,
-    /// short-circuit &&, and %.
-    #[test]
-    fn c_semantics() {
-        let code = r#"
-            __kernel void process(__global const float* in_pos, __global const float* in_col, int in_count,
-                                  __global float* out_pos, __global float* out_col, __global int* out_count, int max_vertices) {
-                int id = get_global_id(0);
-                if (id == 0) {
-                    int count = 0;
-                    int idx = count++;                 // idx = 0, count = 1
-                    int div = 7 / 2;                   // 3, not 3.5
-                    float fdiv = (float)7 / 2.0f;      // 3.5
-                    int rem = 7 % 4;                   // 3
-                    int t = div == 3 ? 10 : 20;        // 10
-                    int guard = 0;
-                    if (guard && (1 / guard) > 0) { t = 999; }   // must not divide
-                    out_pos[0] = (float)idx; out_pos[1] = (float)count; out_pos[2] = (float)div;
-                    out_col[0] = fdiv; out_col[1] = (float)rem; out_col[2] = (float)t;
-                    *out_count = 1;
-                }
-            }"#;
-        let g = gen(code, &[]);
-        assert_eq!(g.vertices.len(), 1);
-        assert_eq!(g.vertices[0].pos, [0.0, 1.0, 3.0]);
-        assert_eq!(g.vertices[0].col, [3.5, 3.0, 10.0]);
-    }
-
-    /// User-defined function with buffer pointers and &address-of an int —
-    /// the box template's add_box shape.
-    #[test]
-    fn user_function_with_pointers() {
-        let code = r#"
-            void emit(float x, __global float* out_pos, __global float* out_col, int* count, int max_vertices) {
-                int start = *count;
-                if (start < max_vertices) {
-                    out_pos[start * 3 + 0] = x;
-                    out_pos[start * 3 + 1] = x * 2.0f;
-                    out_pos[start * 3 + 2] = 0.0f;
-                    out_col[start * 3 + 0] = 1.0f;
-                }
-                *count = start + 1;
-            }
-            __kernel void process(__global const float* in_pos, __global const float* in_col, int in_count,
-                                  __global float* out_pos, __global float* out_col, __global int* out_count, int max_vertices) {
-                if (get_global_id(0) == 0) {
-                    int count = 0;
-                    for (int i = 0; i < 3; i++) {
-                        emit((float)i + 1.0f, out_pos, out_col, &count, max_vertices);
-                    }
-                    *out_count = count;
-                }
-            }"#;
-        let g = gen(code, &[]);
-        assert_eq!(g.vertices.len(), 3);
-        assert_eq!(g.vertices[1].pos, [2.0, 4.0, 0.0]);
-        assert_eq!(g.vertices[2].pos, [3.0, 6.0, 0.0]);
-    }
-
-    /// Local fixed arrays with initializer lists — the sphere/box pattern.
-    #[test]
-    fn array_initializers() {
-        let code = r#"
-            __kernel void process(__global const float* in_pos, __global const float* in_col, int in_count,
-                                  __global float* out_pos, __global float* out_col, __global int* out_count, int max_vertices) {
-                if (get_global_id(0) == 0) {
-                    float px[3] = {1.5f, 2.5f, 3.5f};
-                    int order[3] = {2, 0, 1};
-                    int n = 0;
-                    for (int i = 0; i < 3; i++) {
-                        int idx = n++;
-                        out_pos[idx * 3] = px[order[i]];
-                    }
-                    *out_count = n;
-                }
-            }"#;
-        let g = gen(code, &[]);
-        assert_eq!(g.vertices.len(), 3);
-        assert_eq!(g.vertices[0].pos[0], 3.5);
-        assert_eq!(g.vertices[1].pos[0], 1.5);
-        assert_eq!(g.vertices[2].pos[0], 2.5);
-    }
-
-    /// Deformer mode: per-vertex ids, in-place pos/col, param_values binding.
-    #[test]
-    fn deformer_mode_and_params() {
-        let code = r#"
-            __kernel void process(__global float* pos, __global float* col, int count, __global const float* param_values) {
-                int id = get_global_id(0);
-                if (id < count) {
-                    pos[id * 3 + 1] = pos[id * 3 + 1] + param_values[0];
-                }
-            }"#;
-        let mut g = Geometry::new();
-        for i in 0..4 {
-            g.vertices.push(GVertex {
-                pos: [i as f32, 1.0, 0.0],
-                col: [0.0; 3],
-                attributes: HashMap::new(),
-            });
-        }
-        run_kernel_cpu(code, &mut g, &[2.5]).unwrap();
-        for v in &g.vertices {
-            assert_eq!(v.pos[1], 3.5);
-        }
-    }
-
-    /// A kernel that never terminates errors out instead of hanging the app.
-    #[test]
-    fn step_budget_stops_runaway_kernels() {
-        let code = r#"
-            __kernel void process(__global const float* in_pos, __global const float* in_col, int in_count,
-                                  __global float* out_pos, __global float* out_col, __global int* out_count, int max_vertices) {
-                while (1) { int x = 0; }
-            }"#;
-        let mut g = Geometry::new();
-        let err = run_kernel_cpu_with_budget(code, &mut g, &[], 100_000).unwrap_err();
-        assert!(err.contains("step budget"), "unexpected error: {err}");
-    }
-}
-
-/// Template-kernel validation: the CPU backend run against the REAL shipped
-/// kernels, with absolute geometric asserts (green with or without an OpenCL
-/// runtime — the coverage the OpenCL-side tests could never give headless),
-/// and a cross-validation pass comparing CPU output to OpenCL vertex-by-vertex
-/// whenever a platform exists. The reference is only a reference if the two
-/// backends agree.
-#[cfg(test)]
-mod template_tests {
-    use super::*;
-    use crate::geometry::{parse_dynamic_params, preprocess_opencl_code, run_opencl_kernel_with_params};
-
-    /// A small generator in the kernel language: `Count` triangles of
-    /// `Size`, one per unit along x. The four template kernels this suite
-    /// used to run are native nodes now (2026-09-24), so the launcher
-    /// contract is exercised on this and on the OpenCL node's own default.
-    const GENERATOR: &str = r#"
-__kernel void process(__global const float* in_pos, __global const float* in_col, int in_count, __global float* out_pos, __global float* out_col, __global int* out_count, int max_vertices) {
-    int id = get_global_id(0);
-    if (id == 0) {
-        float s = chf("Size", 0.5f);
-        int n = chi("Count", 3);
-        int count = 0;
-        for (int t = 0; t < n; t++) {
-            float px[3] = {0.0f, s, 0.0f};
-            float py[3] = {0.0f, 0.0f, s};
-            for (int v = 0; v < 3; v++) {
-                int idx = count++;
-                if (idx < max_vertices) {
-                    out_pos[idx * 3 + 0] = px[v] + (float)t;
-                    out_pos[idx * 3 + 1] = py[v];
-                    out_pos[idx * 3 + 2] = 0.0f;
-                    out_col[idx * 3 + 0] = 0.5f + 0.5f * px[v];
-                    out_col[idx * 3 + 1] = 0.2f;
-                    out_col[idx * 3 + 2] = (float)t / (float)n;
-                }
-            }
-        }
-        *out_count = count;
-    }
-}"#;
-
-    /// The OpenCL node's shipped default: the one kernel still in a template.
-    fn opencl_node_default() -> String {
-        let templates = crate::app::load_fs_tree();
-        let node = templates.children.iter().find(|t| t.node_type == "opencl").expect("the OpenCL node template");
-        node.params.iter().find(|p| p.name == "Code").expect("Code param").default.clone()
-    }
-
-    /// A kernel prepared as the launcher prepares it: preprocessed, with its
-    /// parameters flattened in declaration order at their defaults.
-    fn prepared(code: &str) -> (String, Vec<f32>) {
-        let mut flat = Vec::new();
-        for p in parse_dynamic_params(code) {
-            if p.param_type == "float3" {
-                let parts: Vec<f32> = p.default.split(':').filter_map(|s| s.parse().ok()).collect();
-                flat.extend_from_slice(&[
-                    parts.first().copied().unwrap_or(0.0),
-                    parts.get(1).copied().unwrap_or(0.0),
-                    parts.get(2).copied().unwrap_or(0.0),
-                ]);
-            } else if p.default.eq_ignore_ascii_case("true") {
-                flat.push(1.0);
-            } else if p.default.eq_ignore_ascii_case("false") {
-                flat.push(0.0);
-            } else {
-                flat.push(p.default.parse().unwrap_or(0.0));
-            }
-        }
-        (preprocess_opencl_code(code), flat)
-    }
-
-    fn triangle() -> Geometry {
-        let mut g = Geometry::new();
-        for p in [[0.0f32, 0.0, 0.0], [1.0, 0.0, 0.0], [0.0, 1.0, 0.0]] {
-            g.vertices.push(GVertex { pos: p, col: [0.5; 3], attributes: HashMap::new() });
-        }
-        g
-    }
-
-    #[test]
-    fn cpu_runs_a_generator_kernel() {
-        let (code, params) = prepared(GENERATOR);
-        assert_eq!(params, vec![0.5, 3.0], "Size then Count, in declaration order");
-        let mut g = Geometry::new();
-        run_kernel_cpu(&code, &mut g, &params).expect("generator kernel");
-        assert_eq!(g.vertices.len(), 9, "three triangles of three corners");
-        assert_eq!(g.vertices[3].pos, [1.0, 0.0, 0.0], "the second triangle starts one unit along");
-        assert!((g.vertices[4].col[0] - 0.75).abs() < 1e-6);
-    }
-
-    #[test]
-    fn cpu_runs_the_opencl_nodes_default_deformer() {
-        let (code, params) = prepared(&opencl_node_default());
-        let mut g = triangle();
-        run_kernel_cpu(&code, &mut g, &params).expect("default deformer");
-        assert_eq!(g.vertices.len(), 3, "a deformer keeps its vertex count");
-        // y += sin(x * 4) * 0.15 at x = 1: the second corner rises.
-        assert!((g.vertices[1].pos[1] - (4.0f32).sin() * 0.15).abs() < 1e-5, "{:?}", g.vertices[1].pos);
-        assert_eq!(g.vertices[0].pos[1], 0.0, "and x = 0 does not move");
-    }
-
-    /// The reference test proper: byte-level agreement with OpenCL on every
-    /// shipped kernel. Skips silently where no platform exists — the absolute
-    /// tests above still cover the CPU side there.
-    #[test]
-    fn cpu_matches_opencl_on_every_shipped_kernel() {
-        let default = opencl_node_default();
-        for (name, source, input) in [
-            ("generator", GENERATOR.to_string(), Geometry::new()),
-            ("opencl default", default, triangle()),
-        ] {
-            let (code, params) = prepared(&source);
-
-            let mut gpu = input.clone();
-            match run_opencl_kernel_with_params(&code, &mut gpu, &params) {
-                Err(e) if e.contains("No OpenCL platforms") => return, // headless: nothing to compare against
-                Err(e) => panic!("{name} on OpenCL: {e}"),
-                Ok(()) => {}
-            }
-
-            let mut cpu = input;
-            run_kernel_cpu(&code, &mut cpu, &params).unwrap_or_else(|e| panic!("{name} on CPU: {e}"));
-
-            assert_eq!(cpu.vertices.len(), gpu.vertices.len(), "{name}: vertex count diverges");
-            for (i, (a, b)) in cpu.vertices.iter().zip(&gpu.vertices).enumerate() {
-                for k in 0..3 {
-                    assert!(
-                        (a.pos[k] - b.pos[k]).abs() < 1e-4,
-                        "{name}: pos[{k}] diverges at vertex {i}: cpu {} vs gpu {}",
-                        a.pos[k],
-                        b.pos[k]
-                    );
-                    assert!(
-                        (a.col[k] - b.col[k]).abs() < 1e-4,
-                        "{name}: col[{k}] diverges at vertex {i}: cpu {} vs gpu {}",
-                        a.col[k],
-                        b.col[k]
-                    );
-                }
-            }
-        }
-    }
-}
diff --git a/src/main.rs b/src/main.rs
index 7dc3b5d..4035cfb 100644
--- a/src/main.rs
+++ b/src/main.rs
@@ -22,7 +22,6 @@ pub mod viewport_3d;
 pub mod api;
 pub mod window;
 pub mod geometry;
-pub mod kernel_cpu;
 pub mod project;
 pub mod render;
 pub mod shortcut;
@@ -237,7 +236,7 @@ mod tests {
         let child = |name: &str| FsNode {
             id: name.to_string(),
             name: name.to_string(),
-            node_type: "opencl".to_string(),
+            node_type: "grid".to_string(),
             children: vec![],
             params: vec![],
             geometry_visible: true,
@@ -2315,7 +2314,7 @@ mod tests {
             &mut ocl_err,
             &mut crate::geometry::EvalSim::new(0, 0, &mut crate::geometry::SimCache::default()),
         ).expect("Group geometry generation failed");
-        assert!(ocl_err.is_none(), "OpenCL compilation error: {:?}", ocl_err);
+        assert!(ocl_err.is_none(), "node error: {:?}", ocl_err);
 
         let members = crate::geometry::group_member_positions(&geom, "group1");
         assert!(!members.is_empty(), "the box should tag the upper hemisphere");
@@ -2820,16 +2819,25 @@ mod tests {
         let templates_root = crate::app::load_fs_tree();
         let templates = crate::app::flatten_node_templates(&templates_root);
         let group_t = templates_root.children.iter().find(|t| t.name == "Group").unwrap();
-        let opencl_t = templates_root.children.iter().find(|t| t.node_type == "opencl").unwrap();
         let output_t = templates_root.children.iter().find(|t| t.node_type == "output").unwrap();
 
         // An "old save": a Sphere instance from before the construction
         // controls AND from before the port — a subnet of opencl1 → output1
         // holding only Radius, with a user value, and a stale kernel.
-        let mut opencl1 = opencl_t.clone();
-        opencl1.id = "s_opencl1".to_string();
-        opencl1.name = "opencl1".to_string();
-        opencl1.params.iter_mut().find(|p| p.name == "Code").unwrap().default = "OLD KERNEL".to_string();
+        let opencl1 = FsNode {
+            id: "s_opencl1".to_string(),
+            name: "opencl1".to_string(),
+            node_type: "opencl".to_string(),
+            children: vec![],
+            params: vec![crate::app::ParamDef {
+                name: "Code".into(), label: String::new(), param_type: "code".into(), default: "OLD KERNEL".into(),
+                options: vec![], min: None, max: None, step: None, show_when: String::new(), expr: false,
+            }],
+            geometry_visible: true,
+            position: (4.0, 2.0),
+            inputs: 1,
+            outputs: 1,
+        };
         let mut output1 = output_t.clone();
         output1.id = "s_output1".to_string();
         output1.name = "output1".to_string();
@@ -3215,7 +3223,7 @@ mod tests {
                 &mut ocl_err,
                 &mut crate::geometry::EvalSim::new(0, 0, &mut crate::geometry::SimCache::default()),
             ).expect("Geometry generation failed");
-            assert!(ocl_err.is_none(), "OpenCL compilation error: {:?}", ocl_err);
+            assert!(ocl_err.is_none(), "node error: {:?}", ocl_err);
             geom
         };
 
@@ -3245,54 +3253,6 @@ mod tests {
         assert_eq!(geom_3.num_points(), 5 * 9, "a 4x8 cell grid is 5x9 points");
     }
 
-    #[test]
-    fn test_dynamic_parameters_parsing_and_preprocessing() {
-        let code = r#"
-            float freq = chf("freq", 4.0f);
-            int count = chi("count", 15);
-            float3 col = chv("col", 0.8f, 0.2f, 0.2f);
-            float scale = chf("scale");
-        "#;
-        
-        let parsed = crate::geometry::parse_dynamic_params(code);
-        assert_eq!(parsed.len(), 4);
-        
-        assert_eq!(parsed[0].name, "freq");
-        assert_eq!(parsed[0].param_type, "slider");
-        assert_eq!(parsed[0].default, "4.0");
-        
-        assert_eq!(parsed[1].name, "scale");
-        assert_eq!(parsed[1].param_type, "slider");
-        assert_eq!(parsed[1].default, "0.5");
-        
-        assert_eq!(parsed[2].name, "count");
-        assert_eq!(parsed[2].param_type, "spinbox");
-        assert_eq!(parsed[2].default, "15");
-        
-        assert_eq!(parsed[3].name, "col");
-        assert_eq!(parsed[3].param_type, "float3");
-        assert_eq!(parsed[3].default, "0.80:0.20:0.20");
-        
-        let mut target = FsNode {
-            id: "node1".to_string(),
-            name: "OpenCL Node".to_string(),
-            node_type: "opencl".to_string(),
-            children: Vec::new(),
-            params: parsed,
-            geometry_visible: true,
-            position: (0.0, 0.0),
-            inputs: 1,
-            outputs: 1,
-        };
-        
-        target.params[1].default = "1.25".to_string(); // "scale"
-        let preprocessed = crate::geometry::preprocess_opencl_code(code);
-        assert!(preprocessed.contains("param_values[0]"));
-        assert!(preprocessed.contains("((int)param_values[2])"));
-        assert!(preprocessed.contains("(float3)(param_values[3], param_values[4], param_values[5])"));
-        assert!(preprocessed.contains("param_values[1]"));
-    }
-
     #[test]
     fn test_project_serialization_roundtrip() {
         let root = FsNode {
@@ -6241,7 +6201,7 @@ mod tests {
     #[test]
     fn a_code_parameter_is_never_an_expression() {
         use crate::app::{infer_template_exprs, McpAction};
-        let mut node = ref_node("w", "w1", "opencl", vec![("Code", "code", "ch(\"../a/Radius\")"), ("Radius", "slider", "ch(\"../a/Radius\")")], vec![]);
+        let mut node = ref_node("w", "w1", "wrangle", vec![("Code", "code", "ch(\"../a/Radius\")"), ("Radius", "slider", "ch(\"../a/Radius\")")], vec![]);
         for p in &mut node.params {
             p.expr = false;
         }
@@ -6251,7 +6211,7 @@ mod tests {
 
         let mut state = State::new(false);
         let mut redraw = false;
-        state.apply_action(McpAction::AddNode { template_name: "OpenCL".into(), name: Some("k".into()), x: 3.0, y: 9.0 }, &mut redraw).unwrap();
+        state.apply_action(McpAction::AddNode { template_name: "Wrangle".into(), name: Some("k".into()), x: 3.0, y: 9.0 }, &mut redraw).unwrap();
         let k = state.current_dir().children.iter().position(|c| c.name == "k").unwrap();
         state.apply_action(McpAction::SetParam { slot: k, name: "Code".into(), value: "chf(\"../sphere1/Radius\")".into() }, &mut redraw).unwrap();
         let code = state.current_dir().children[k].params.iter().find(|p| p.name == "Code").unwrap();
@@ -10616,4 +10576,27 @@ mod tests {
         let moved = (0..g.num_points()).filter(|&p| (g.pos(p).y - before.pos(p).y).abs() > 1e-6).count();
         assert!(moved > 0, "the default script deforms the input");
     }
+
+    /// An `opencl` node in a save older than the retirement is not dropped:
+    /// it passes its input through and says what it is, on the status line
+    /// and in the CLI's warning, so the fix is one rewrite as a wrangle.
+    #[test]
+    fn a_retired_opencl_node_passes_its_input_through_and_says_so() {
+        let src = ref_node("s", "src", "sphere", vec![("Radius", "slider", "1.0")], vec![]);
+        let k = ref_node("k", "opencl1", "opencl", vec![("Input", "text", "src"), ("Code", "code", "__kernel void process() {}")], vec![]);
+        let root = ref_node("root", "root", "node", vec![], vec![src, k]);
+        let before = eval(&root, &root.children[0]).0.unwrap();
+        let (g, err) = eval(&root, &root.children[1]);
+        let err = err.expect("the retired node reports itself");
+        assert!(err.starts_with("opencl1: OpenCL nodes are retired"), "{err}");
+        assert!(err.contains("wrangle"), "and points at the replacement: {err}");
+        let g = g.expect("the input passes through");
+        assert_eq!(g.num_points(), before.num_points());
+        assert_eq!(g.pos(5), before.pos(5));
+        assert!(crate::geometry::is_geometry_node_type("opencl"), "the type still resolves, to that arm");
+
+        // And no template offers it any more.
+        let templates = crate::app::load_fs_tree();
+        assert!(!templates.children.iter().any(|t| t.node_type == "opencl"), "nodes/opencl.json is gone");
+    }
 }
diff --git a/src/render.rs b/src/render.rs
index 6aee181..688e29d 100644
--- a/src/render.rs
+++ b/src/render.rs
@@ -2,7 +2,7 @@
 use cce_ui::colors;
 use cce_ui::widget::WidgetHost;
 
-use crate::app::{State, FsNode};
+use crate::app::State;
 use crate::slots::{
     WIDGET_COUNT,
     CONTENT_IDX, VIEWPORT_IDX, PARAM_IDX,
@@ -1192,23 +1192,8 @@ impl State {
         };
         self.sim_cache = sim_cache;
 
-        fn has_visible_opencl(node: &FsNode) -> bool {
-            if node.node_type.eq_ignore_ascii_case("opencl") && node.geometry_visible {
-                return true;
-            }
-            for child in &node.children {
-                if has_visible_opencl(child) {
-                    return true;
-                }
-            }
-            false
-        }
-
-        let displayed_opencl = has_visible_opencl(self.viewport_editor_dir());
         if let Some(e) = ocl_error {
-            self.update_status_text(&format!("OpenCL Error: {}", e));
-        } else if displayed_opencl {
-            self.update_status_text("OpenCL kernel executed successfully.");
+            self.update_status_text(&format!("Node error: {}", e));
         } else {
             self.update_status_text("Geometry updated successfully.");
         }
@@ -1216,7 +1201,7 @@ impl State {
         let verts = crate::geometry::detail_vertices(&geom);
         self.vertex_count_spheres = verts.len() as u32;
         // Cache for the path tracer, so RT mode never re-runs the node
-        // graph / OpenCL kernels; the version bump invalidates its scene.
+        // graph; the version bump invalidates its scene.
         // The raster mesh uploads from this same cache on the next
         // `stage_renderer` flush.
         self.rt_sphere_verts = verts;
diff --git a/src/thumbnail.rs b/src/thumbnail.rs
index fbddaae..30f4d38 100644
--- a/src/thumbnail.rs
+++ b/src/thumbnail.rs
@@ -3,7 +3,7 @@
 //!
 //! Loads a project's `state.json`, regenerates its geometry from the node
 //! graph (the same `network_sphere_vertices_with_errors` the viewport uses,
-//! OpenCL nodes included), auto-frames a camera on the scene bounds, and
+//! wrangles included), auto-frames a camera on the scene bounds, and
 //! renders through `cce_ui::vk::RtOffscreen` — no window, no compositor, any
 //! graphics-capable Vulkan device. cce-files shells out to this for its
 //! preview cache.
@@ -47,8 +47,8 @@ pub fn run(project: &Path, out: &Path, size: u32, samples: Option<u32>, frame: O
     // network level the project was saved at: root as both eval and walk root.
     let geom = network_sphere_vertices_with_errors(&proj.root, &proj.root, &mut ocl_error, &mut sim);
     if let Some(e) = ocl_error {
-        // Non-fatal: OpenCL nodes just contribute nothing, like the viewport.
-        eprintln!("thumbnail: OpenCL error (geometry partially skipped): {e}");
+        // Non-fatal: a failing node just contributes nothing, like the viewport.
+        eprintln!("thumbnail: node error (geometry partially skipped): {e}");
     }
     let verts = crate::geometry::detail_vertices(&geom);
     let (tris, mats) = rt_scene_from_verts(&verts);
diff --git a/src/wrangle.rs b/src/wrangle.rs
index 7d1e5d5..e2e0457 100644
--- a/src/wrangle.rs
+++ b/src/wrangle.rs
@@ -4,9 +4,10 @@
 //! Houdini's attribwrangle, on an engine someone else maintains. The script
 //! language is Rhai, chosen in `shapeshifter.md` Phase 7 for being pure Rust
 //! with no C toolchain, sandboxed behind an operation budget — the step budget
-//! `kernel_cpu` reimplements by hand — and compiled once to an AST that is
-//! cached by source, as `OPENCL_CACHE` keys kernels. What this module adds is
-//! the BINDING: how a script reaches the geometry, and nothing else.
+//! the retired `kernel_cpu` interpreter used to count by hand — and compiled
+//! once to an AST that is cached by source, as the retired launcher cached
+//! kernels. What this module adds is the BINDING: how a script reaches the
+//! geometry, and nothing else.
 //!
 //! The vocabulary, deliberately VEX-shaped so it reads as it does there:
 //!