#1578 STRATA_ROUTE_RESIDENT faults with illegal memory access on a two-GPU layer split (works with --pipeline-windows 0)

Description

## Environment

- 2× RTX 3080 Ti 12 GB (sm_86, no P2P between them), i5-13600K, 62 GB RAM, Linux
- CUDA 12.4 toolkit, engine 0.1.41 built locally (`BUILD.json`: `source: local`, `archs: [86]`)
- Model: Qwen3.8-Flash-Next IQ3_S (GSQ-RCO), layer split 32 (layers 0–31 / 32–47), `--pipeline-windows 2`, `--kv q4_0` + `--kv-resident 32768`, MTP draft on

## Repro

1. Start the server with `STRATA_ROUTE_RESIDENT=0.5` (also tried `1.0`).
2. Send any generation request (even `"Say hi."`, `max_tokens: 8`).
3. The request fails; engine log:
```
strata serve: pipelined decode: verify: layer 31 never rang (unspecified launch failure)
strata serve: verify: layer 0 never rang (an illegal memory access was encountered)
```

## What still works

Same config with `--pipeline-windows 0` serves cleanly, including a full LC32 (33,561 prompt tokens / 256 gen): **92.5 tok/s, 80.2% drafts accepted** — on par with our tuned baseline (91.7). So the flag itself is viable on a split; only the decode-pipeline combination faults.

## Analysis

Two independent halves; I fixed the first, the second needs your eye.

**Half 1 (fixed locally, patch below): counter placement.** `RouteResidentCfg::stats()` (`src/core/verify.cpp`) does one process-wide `cudaMalloc` with no device scoping. Verifier inits run per stage under `OnDevice`, but the 8×uint64 block lands on whichever device is current at the first init. On a split without P2P the other stage's kernels fault on it. The residency table `d_res` right next to it is correctly replicated per device in `generate.cpp` — this counter was missed. You never saw it because the flag was only measured on single cards (5070, 3060).

**Half 2 (open): the decode pipeline.** With per-device counters in place, `pipeline-windows 2` still faults while `0` is clean. Suspects, in order:
1. Stats selection at launch time (the `stats() + (n <= 8 ? 4 : 0)` call) picks the block by *host current device*, not by the kernel's stream device. The pipeline's second verifiers (`ver_b` on CUDA0, `gs.ver_b` on the later card) share one stream per card — a stage-1 kernel can easily be handed the stage-0 block.
2. The in-place `ids_` rewrite (`id[r] = best` in `route_resident_k`) interacting with the even/odd parity verifiers that share buffers.

The async fault surfacing at "layer 31, then layer 0" is consistent with either — I stopped here rather than restructure the launch path blind.

## Patch (half 1)

```diff
 struct RouteResidentCfg {
     float margin = 0.0f;
     int lo = 6, hi = 9;
-    unsigned long long* d_stats = nullptr;
+    // one counter block per GPU: a layer split inits a verifier per stage and
+    // each stage's kernels must read memory on their own device (no P2P here)
+    std::map<int, unsigned long long*> dev_stats;
     unsigned long long* stats() {
-        if (d_stats == nullptr && margin > 0.0f) {
-            cudaMalloc((void**) &d_stats, 8 * sizeof(unsigned long long));
-            cudaMemset(d_stats, 0, 8 * sizeof(unsigned long long));
+        if (margin <= 0.0f) return nullptr;
+        int dev = 0;
+        cudaGetDevice(&dev);
+        auto it = dev_stats.find(dev);
+        if (it != dev_stats.end()) return it->second;
+        unsigned long long* p = nullptr;
+        cudaMalloc((void**) &p, 8 * sizeof(unsigned long long));
+        cudaMemset(p, 0, 8 * sizeof(unsigned long long));
+        dev_stats[dev] = p;
+        if (dev_stats.size() == 1) {
             std::atexit([] {
-                unsigned long long h[8] = {};
                 RouteResidentCfg& c = route_resident_cfg();
-                cudaDeviceSynchronize();
-                cudaMemcpy(h, c.d_stats, sizeof(h), cudaMemcpyDeviceToHost);
-                std::fprintf(stderr, "route-resident: margin=%g ranks=%d-%d windows T>8 (prompt reads): tail_nonres=%llu swaps=%llu (%.1f%%) nonres_entries before=%llu after=%llu\n",
-                             c.margin, c.lo, c.hi, h[0], h[1], h[0] ? 100.0 * h[1] / h[0] : 0.0, h[2], h[3]);
-                std::fprintf(stderr, "route-resident: margin=%g ranks=%d-%d windows T<=8 (decode): tail_nonres=%llu swaps=%llu (%.1f%%) nonres_entries before=%llu after=%llu\n",
-                             c.margin, c.lo, c.hi, h[4], h[5], h[4] ? 100.0 * h[5] / h[4] : 0.0, h[6], h[7]);
+                for (auto& kv : c.dev_stats) {
+                    unsigned long long h[8] = {};
+                    cudaSetDevice(kv.first);
+                    cudaDeviceSynchronize();
+                    cudaMemcpy(h, kv.second, sizeof(h), cudaMemcpyDeviceToHost);
+                    std::fprintf(stderr, "route-resident: dev=%d margin=%g ranks=%d-%d windows T>8 (prompt reads): tail_nonres=%llu swaps=%llu (%.1f%%) nonres_entries before=%llu after=%llu\n",
+                                 kv.first, c.margin, c.lo, c.hi, h[0], h[1], h[0] ? 100.0 * h[1] / h[0] : 0.0, h[2], h[3]);
+                    std::fprintf(stderr, "route-resident: dev=%d margin=%g ranks=%d-%d windows T<=8 (decode): tail_nonres=%llu swaps=%llu (%.1f%%) nonres_entries before=%llu after=%llu\n",
+                                 kv.first, c.margin, c.lo, c.hi, h[4], h[5], h[4] ? 100.0 * h[5] / h[4] : 0.0, h[6], h[7]);
+                }
             });
         }
-        return d_stats;
+        return p;
     }
 };
```

`<map>` is already included in `verify.cpp`. Bitwise behaviour on single GPU is unchanged (same bytes, one device in the map).

## Measurements (LC32: 33,561 prompt tokens, 256 gen, greedy)

| Config | TG | acc | Note |
|---|---|---|---|
| tuned baseline (see config below), pipewin 2, no ROUTE_RESIDENT | 91.7 | 84.9% | our prod |
| + ROUTE_RESIDENT=0.5, pipewin 2 | fault | — | this issue |
| + ROUTE_RESIDENT=0.5, pipewin 0 | 92.5 | 80.2% | clean, on par |

**Baseline config ("tuned baseline" above)** — engine flags: `--layer-split 32 --pool-workers 15 --spec 4 --spec-min-p 0.60 --suffix-draft 6 --mtp-max-t 2 --pipeline-windows 2 --pcie-frac 0.00 --max-context 131072 --kv q4_0 --kv-resident 32768 --vram-reserve-mib 400`; env: `STRATA_STAGE_TRIM=1 STRATA_TSUM=1 STRATA_QFUSE=1 STRATA_LFUSE=1`. All three rows use this exact base; only the flagged row/param differs.

Note: ROUTE_RESIDENT changes answers by design (I did not evaluate quality beyond acc). Found with agentic assistance; all runs are real measurements on the hardware above, engine built from source with the patch.

<sub>issue was generated by AI</sub>