fix(gpu-router): dispatch-limit 2D grid + keep-best ripup rounds

crates/vcad-ecad-pcb/Cargo.toml

@@ -25,6 +25,7 @@ log = "0.4"
 
 [dev-dependencies]
 vcad-ecad-sim = { path = "../vcad-ecad-sim" }
+vcad-ecad-export = { path = "../vcad-ecad-export" }
 vcad-ecad-symbols = { workspace = true }
 env_logger = "0.11"
 

crates/vcad-ecad-pcb/examples/board_export.rs

@@ -0,0 +1,52 @@
+//! Export a routed board (`.pcb.json` from cm5_bench) to every supported
+//! fabrication/interchange format: Gerbers, Excellon drills, KiCad board,
+//! BOM CSV, and pick-and-place CSV.
+//!
+//! ```bash
+//! cargo run --release -p vcad-ecad-pcb --example board_export -- routed.pcb.json out_dir/
+//! ```
+
+use std::fs;
+use std::path::Path;
+use vcad_ir::ecad::Pcb;
+
+fn main() {
+    let mut args = std::env::args().skip(1);
+    let input = args
+        .next()
+        .expect("usage: board_export <board.pcb.json> <out_dir>");
+    let out = args
+        .next()
+        .expect("usage: board_export <board.pcb.json> <out_dir>");
+    let out = Path::new(&out);
+    fs::create_dir_all(out).expect("create out dir");
+
+    let pcb: Pcb = serde_json::from_str(&fs::read_to_string(&input).expect("read board"))
+        .expect("parse board json");
+
+    let gerbers = vcad_ecad_export::gerber::generate_gerbers(&pcb).expect("gerbers");
+    for (name, content) in &gerbers {
+        fs::write(out.join(name), content).expect("write gerber");
+    }
+    println!("gerbers: {} files", gerbers.len());
+
+    let mut drill = Vec::new();
+    vcad_ecad_export::excellon::write_excellon(&mut drill, &pcb).expect("excellon");
+    fs::write(out.join("drill.drl"), &drill).expect("write drill");
+
+    fs::write(
+        out.join("board.kicad_pcb"),
+        vcad_ecad_symbols::write_kicad_pcb(&pcb),
+    )
+    .expect("write kicad");
+
+    let mut bom = Vec::new();
+    vcad_ecad_export::bom::write_bom(&mut bom, &pcb).expect("bom");
+    fs::write(out.join("bom.csv"), &bom).expect("write bom");
+
+    let mut pnp = Vec::new();
+    vcad_ecad_export::pick_place::write_pick_place(&mut pnp, &pcb).expect("pick place");
+    fs::write(out.join("pick_place.csv"), &pnp).expect("write pnp");
+
+    println!("exported to {}", out.display());
+}

crates/vcad-ecad-pcb/src/router/auto.rs

@@ -1086,6 +1086,12 @@ fn route_pass(
         if pending.is_empty() {
             break;
         }
+        // Keep-best: a rip-up round can end with FEWER placed connections
+        // than it started (it rips victims it then fails to restore — the
+        // CM5 logs show 504 -> 489 over a 46-minute round). The negotiation
+        // loop above snapshots and discards regressing rounds; this loop
+        // must too, or one bad round permanently costs the board copper.
+        let snapshot = (session.clone(), placed.clone(), pending.clone());
         let placed_before = placed.len();
         let sw_rip = Stopwatch::start();
         pending = ripup_pass(
@@ -1109,7 +1115,12 @@ fn route_pass(
             pending.len(),
             sw_rip.ms() / 1000.0,
         );
-        if placed.len() <= placed_before {
+        if placed.len() < placed_before {
+            log::info!("ripup round regressed — restoring pre-round snapshot");
+            (session, placed, pending) = snapshot;
+            break;
+        }
+        if placed.len() == placed_before {
             break;
         }
     }

crates/vcad-kernel-gpu/src/shaders/wavefront_batch.wgsl

@@ -81,9 +81,15 @@ fn candidate(search: u32, base: u32, x: i32, y: i32, l: u32, cost: u32) -> u32 {
 }
 
 @compute @workgroup_size(256)
-fn relax_batch(@builtin(global_invocation_id) gid: vec3<u32>) {
+fn relax_batch(
+    @builtin(global_invocation_id) gid: vec3<u32>,
+    @builtin(num_workgroups) nwg: vec3<u32>,
+) {
     let total = total_nodes();
-    let flat = gid.x;
+    // 2D dispatch grid: hosts split workgroups over (x, y) because a single
+    // dispatch dimension is capped at 65535 — a full-board batch needs ~2x
+    // that. Flatten back to the linear node index.
+    let flat = gid.y * (nwg.x * 256u) + gid.x;
     if flat >= total * params.n_searches {
         return;
     }

crates/vcad-kernel-gpu/src/wavefront_batch.rs

@@ -274,7 +274,15 @@ pub async fn distance_fields_batch_async(
     let bind_ab = bind(&dist_a, &dist_b);
     let bind_ba = bind(&dist_b, &dist_a);
 
+    // Metal/wgpu cap each dispatch dimension at 65535 workgroups; a
+    // full-board batch exceeds that in one dimension (observed 129052 —
+    // the dispatch poisoned the encoder and silently killed the GPU path
+    // for the rest of the run). Split over a 2D grid; the shader flattens
+    // (x, y) back to the linear index via num_workgroups.
     let groups = ((n * total) as u32).div_ceil(256);
+    let max_dim = ctx.device.limits().max_compute_workgroups_per_dimension;
+    let groups_x = groups.min(max_dim);
+    let groups_y = groups.div_ceil(groups_x);
     // Worst-case sweep bound: longest shortest path < total nodes.
     let max_sweeps = 4 * (nx + ny + nl);
     let mut sweeps_done = 0usize;
@@ -294,7 +302,7 @@ pub async fn distance_fields_batch_async(
             });
             pass.set_pipeline(&pipeline);
             pass.set_bind_group(0, if current_is_a { &bind_ab } else { &bind_ba }, &[]);
-            pass.dispatch_workgroups(groups, 1, 1);
+            pass.dispatch_workgroups(groups_x, groups_y, 1);
             drop(pass);
             current_is_a = !current_is_a;
             sweeps_done += 1;