From 90a74b7d1338733cc2a6b6f03623bba080c0028f Mon Sep 17 00:00:00 2001
From: Scott Woolsey Biggart <297695+swbiggart@users.noreply.github.com>
Date: Mon, 19 Oct 2026 18:28:06 -0700
Subject: [PATCH 1/5] perf(dsv41): opt in to adjacent-edge TP4 transport

---
 README.md                                   |   9 +-
 docs/receipts/dsv41-neighbor-host/README.md |  41 +++
 docs/recipes/deepseek-v41-neighbor.md       |  74 +++++
 zig/build/cuda.zig                          |  45 +--
 zig/src/families/dsv41/draft.zig            |   8 +
 zig/src/families/dsv41/dsv41.zig            |   2 +
 zig/src/families/dsv41/model.zig            |  38 ++-
 zig/src/families/dsv41/neighbor.zig         | 304 ++++++++++++++++++++
 zig/src/families/dsv41/prompt2d.zig         |  46 ++-
 zig/src/families/dsv41/rdma.zig             |  27 ++
 10 files changed, 538 insertions(+), 56 deletions(-)
 create mode 100644 docs/receipts/dsv41-neighbor-host/README.md
 create mode 100644 docs/recipes/deepseek-v41-neighbor.md
 create mode 100644 zig/src/families/dsv41/neighbor.zig

diff --git a/README.md b/README.md
index 0a1fa048ec..4a2766f303 100644
--- a/README.md
+++ b/README.md
@@ -3,10 +3,11 @@
 > (MIT), our EXL3 quant of [deepseek-ai/DeepSeek-V4.1-Flash-1017](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash-1017)
 > (MIT). Releases up to v0.1.1 served [Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw](https://huggingface.co/Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw)
 > (MIT), an EXL3 quant of [deepseek-ai/DeepSeek-V4.1-Flash](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash).
-> The branch adds a `dsv41` family to TensorFold 1.0.12's native engine: a four-node split (on the 2.9 bpw weights its
-> replies equal the two-node family's token for token in our gates), image input, up to 1,048,576 tokens a request in
-> a shared 2,097,152-position window with kept-prompt reuse, and exact DSpark speculative decoding. Launch, measured
-> numbers, gates and the kernel kit are in the recipe, [bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10](https://github.com/bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10).
+> The branch adds a `dsv41` family to TensorFold 1.0.2's native engine: a four-node split whose replies equal the
+> two-node family's token for token in our gates, image input, up to 1,048,576 tokens a request in a shared
+> 2,097,152-position window with kept-prompt reuse, and exact DSpark speculative decoding. Launch, measured numbers,
+> gates and the kernel kit are in the recipe, [bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10](https://github.com/bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10).
+> An optional [neighbor transport](docs/recipes/deepseek-v41-neighbor.md) supports adjacent-edge RDMA on a physical ring.
 > `tools/dsv41-zig/inputs/` makes each node's weight file and token map from the checkpoint; `tools/dsv41-zig/kit/`
 > makes and checks the kit's generated files. See NOTICE and ATTRIBUTION.md. This fork is not affiliated with or
 > endorsed by DeepSeek, NVIDIA, the TensorFold authors or Mia-AiLab.
diff --git a/docs/receipts/dsv41-neighbor-host/README.md b/docs/receipts/dsv41-neighbor-host/README.md
new file mode 100644
index 0000000000..f4b5882fe9
--- /dev/null
+++ b/docs/receipts/dsv41-neighbor-host/README.md
@@ -0,0 +1,41 @@
+# TP4 neighbor transport: host checks
+
+Local checks on 2026-10-20 UTC, macOS arm64, Zig 0.17.0. This change is based on
+`bertholomus/TensorFold` branch `deepseek-v41-zig-tp4` at
+`59402086381a42c74d53053443224eec2790b4c5`.
+
+The transport is opt-in through `TF_DS_NEIGHBOR_RDMA=1`; see the
+[configuration and dispatch contract](../../recipes/deepseek-v41-neighbor.md).
+No hardware or model qualification was performed for this branch.
+
+| Check | Result |
+| --- | --- |
+| `zig build test-dsv41-neighbor --summary all` | PASS: 14 tests, 0 failures, 0 skips; no GPU/verbs execution |
+| `zig build test test-golden -Dcpu=apple_m1 --summary all` | Inherited compile failure at `dsv41/link.zig:23`: macOS `std.c.SO` has no `BUSY_POLL`; reproduced on the pinned base and candidate |
+| Frozen server corpus, from the same command | PASS on base and candidate: 312/319 equal, 7 known differences, 0 unexpected differences |
+| `zig build native -Dtarget=aarch64-linux-gnu.2.28 -Doptimize=fast --summary all` | PASS: Linux arm64 cross-build, without CUDA fatbins; not a runnable model qualification |
+| Linux arm64 `dsv41` host-test module | PASS: semantic compilation with `zig test -fno-emit-bin -Odebug -target aarch64-linux-gnu.2.28` and the build's CUDA/core/lanes/fabric modules; not executed |
+| `python3 tools/zig/lean_check.py` | Base: 848 files, 1,511 problems. Candidate: 849 files, 1,511 problems. Zero added findings; checker remains nonzero |
+| `zig fmt --check` on changed Zig files; `git diff --check` | PASS |
+
+The base host run reported 420/422 tests passed and two skips before the existing
+compile failure. The candidate reused unchanged successful test steps from that
+run and reported five newly run passes and one skip, with the same compile
+failure. The golden runner also logs an inherited allocator leak on both runs;
+its build step exits successfully. Neither complete host command is reported as
+a pass.
+
+The focused tests exercise host-side preconditions, startup-mode comparison,
+actual NCCL forwarding call sequencing with a recording API, and abort-word
+propagation. They do not validate device output, RDMA queue pairs, rail selection,
+CUDA graph execution, startup collectives on four machines, throughput, or model
+tokens. GPU/model gates in CONTRIBUTING.md remain pending, including byte/layout
+and graph checks, plain/drafted and concurrent/solo equality, prompt reuse,
+long-context behavior and matched performance. No fault, restart, or reboot gate
+was run.
+
+The earlier experimental implementation used source
+`b65c8b92d2403dfdefccf15a3b23ba609a4553e2`, a descendant of this pinned base, plus
+separate patches. Its device and serving results support the transport design;
+they are not measurements of this consolidated branch or its startup agreement.
+No unrelated model, server, forced-tool, kernel or deployment changes were ported.
diff --git a/docs/recipes/deepseek-v41-neighbor.md b/docs/recipes/deepseek-v41-neighbor.md
new file mode 100644
index 0000000000..43ca22bf66
--- /dev/null
+++ b/docs/recipes/deepseek-v41-neighbor.md
@@ -0,0 +1,74 @@
+# DeepSeek-V4.1 TP4 neighbor transport
+
+The native server and lane runner can opt into adjacent-edge RDMA on a four-node
+physical ring numbered `0 → 1 → 2 → 3 → 0`. It reuses the existing RDMA kernels
+and NCCL; it changes byte movement, not reductions or drafting policy.
+
+Set `TF_DS_NEIGHBOR_RDMA=1` on **every rank**. Unset or `0` retains the existing
+NCCL or `TF_RDMA_DEVICES` transport. There are no separate optimization switches.
+Other values fail startup. The neighbor mode requires TP4 and cannot be combined
+with `TF_RDMA_DEVICES`.
+
+For each rank, set both comma-separated device lists:
+
+| Setting | Physical neighbor | Devices |
+| --- | --- | --- |
+| `TF_DS_NEIGHBOR_FIRST` | rank xor 1 | Two rails to that neighbor |
+| `TF_DS_NEIGHBOR_SECOND` | rank xor 3 | Two rails to that neighbor |
+
+Each list must contain two distinct names; the lists cannot overlap. List rails
+in the same physical order at both ends of each edge (for example, by PCI domain).
+The mode uses RoCE GID index 3 and the existing transport's MTU and queue settings.
+It does not detect or configure networking. NCCL must also reach adjacent ranks
+and support its ordinary collectives on this fabric.
+
+A one-byte NCCL all-gather at TP4 model startup checks that every rank selected
+the same mode. This check also runs when the mode is off, so a mixed selection
+fails before neighbor setup. It creates no RDMA resources when disabled. Use the
+same binary on every rank; mixing versions with different startup collectives is
+unsupported. Device selection and mode stay fixed for the lifetime of the model
+and its captured graphs.
+
+## Operations
+
+Aligned, positive shards through 512 KiB use the following paths:
+
+- Full gathers, including Markov best scores: xor1, then xor3, in global rank order.
+- Pair gathers: xor1 only.
+- Equal-width quarters with 16-byte-aligned rows: collect directly into the final layout.
+- Parity diagonal exchange: remote-only xor1, then xor3, yielding rank xor2.
+
+Larger or unaligned shards retain NCCL. Unequal-width quarters retain padded
+gather and layout copies. In neighbor mode, diagonal NCCL traffic is forwarded
+over two physical edges; it never requires a direct diagonal link. This includes
+the Engram exchange, whose existing arithmetic and order are retained.
+
+Each rank owns two existing RDMA ring instances, two proxy threads, and 1 MiB
+of device scratch in addition to the rings' registered host buffers. Operations
+use the model stream in order; scratch is not shared across streams. A detected
+local stage failure aborts both stages' device waits, and model synchronization
+checks the failure before returning results. This adds no watchdog, retry,
+recovery policy, or distributed failure guarantee.
+
+## Validation
+
+Use the pinned Zig 0.17.0 and the existing host checks:
+
+```sh
+zig build test-dsv41-neighbor
+zig build test test-golden -Dcpu=apple_m1
+python3 tools/zig/lean_check.py
+```
+
+The focused host target checks mode disagreement, bounded dispatch, device-list
+validation, relay preconditions, the production NCCL forwarding calls at every
+rank with unequal counts, and local abort propagation. It also compiles the
+GPU/verbs declarations. It does not execute GPU kernels, RDMA traffic, captured
+graphs, or a model.
+
+Before hardware acceptance, compare all output bytes with an independent NCCL
+reference at all four ranks, including pair/full/quarter/diagonal operations,
+changing-input graph replay, cutoff and unequal-width fallbacks, and mixed
+operation sequences. Then apply CONTRIBUTING.md's plain/drafted,
+concurrent/solo and resumed/fresh token gates and matched serving benchmarks.
+Prototype measurements on a different source revision do not qualify this branch.
diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig
index e0429d8c46..c50fa9bdde 100644
--- a/zig/build/cuda.zig
+++ b/zig/build/cuda.zig
@@ -308,45 +308,12 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ .root_source_file = b.path("zig/build/check_fatbins.zig"), .target = host }) })).step);
     const cuda = runtime(b, host, .debug, &.{});
     const mods = family(b, host, .debug, cuda, draft_ids);
-    const native_modules = engines(b, host, .debug, cuda, mods.core, mods.lanes, mods.nemotron, mods.heat);
-    const native = native_modules.engines;
-    // the HTTP server's unit tests (zig/src/server/root.zig), which the macOS build runs, on Linux hosts too
-    const template = b.createModule(.{ .root_source_file = b.path("zig/src/core/template/template.zig"), .target = host, .optimize = .debug, .link_libc = true });
-    const server_tests = b.createModule(.{
-        .root_source_file = b.path("zig/src/server/root.zig"),
-        .target = host,
-        .optimize = .debug,
-        .link_libc = true,
-        .imports = &.{ .{ .name = "engine_api", .module = native_modules.api }, .{ .name = "tokenizer", .module = mods.tokenizer }, .{ .name = "template", .module = template } },
-    });
-    const server_json = @import("server.zig").addTo(b, server_tests, host, .debug);
-    const server_test = b.addRunArtifact(b.addTest(.{ .root_module = server_tests }));
-    const json_test = b.addRunArtifact(b.addTest(.{ .root_module = server_json }));
-    const server_step = b.step(
-        "test-server-cpu",
-        "The HTTP server's unit tests (routes, templates, tool parsing) without a GPU",
-    );
-    for ([_]*std.Build.Step.Run{ server_test, json_test }) |run| {
-        step.dependOn(&run.step);
-        server_step.dependOn(&run.step);
-    }
-    // the server goldens and Cortheon's fuzz step here when the target is not macOS, whose Metal build defines them
-    if (server_steps) {
-        golden_build.step(b, host, server_tests, native_modules.api, mods.tokenizer, template);
-        @import("cortheon_fuzz.zig").add(b, host, native_modules.api, mods.tokenizer, template);
-    }
-    // the tokenizer's own tests (memo, steps, regex, unicode) run here: no other test binary reaches them
-    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, lfm2Module(b, host, .debug, cuda, mods.core), mods.radio, mods.heat, mods.tokenizer, native_modules.api, native, native_modules.dsv41, stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
-    // the GLM-5.3-Flash CUDA family's host side
-    const glm = glmCuda(b, host, .debug, cuda, mods.core, mods.lanes, mods.heat);
-    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = glm })).step);
-    // the shares a GLM-5.3-Flash rank takes (cluster/canon.zig against the Python family's)
-    const shares = b.createModule(.{ .root_source_file = b.path("zig/src/cluster/canon_glm_test.zig"), .target = host, .optimize = .debug });
-    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = shares })).step);
-    // the tool-call parser's tests, run beside the server's on Macs: the GLM calls the CUDA port serves
-    const parse_tests = b.createModule(.{ .root_source_file = b.path("zig/src/server/tool_parse_test.zig"), .target = host, .optimize = .debug, .link_libc = true, .imports = &.{ .{ .name = "engine_api", .module = native_modules.api }, .{ .name = "tokenizer", .module = mods.tokenizer }, .{ .name = "template", .module = template } } });
-    _ = @import("server.zig").addTo(b, parse_tests, host, .debug);
-    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = parse_tests })).step);
+    const eng = engines(b, host, .debug, cuda, mods.lanes, mods.nemotron, mods.core);
+    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, eng.api, eng.engines, eng.dsv41, stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
+    const neighbor = b.createModule(.{ .root_source_file = b.path("zig/src/families/dsv41/neighbor.zig"), .target = host, .optimize = .debug, .link_libc = true });
+    neighbor.addImport("cuda", cuda);
+    neighbor.addImport("fabric", eng.dsv41.import_table.get("fabric").?);
+    b.step("test-dsv41-neighbor", "Host checks of TP4 neighbor transport (no GPU or verbs)").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = neighbor })).step);
     const cli = b.createModule(.{ .root_source_file = b.path("zig/src/cli/cuda_main.zig"), .target = host, .optimize = .debug, .link_libc = true });
     cli.addImport("cuda", cuda);
     cli.addImport("core", mods.core);
diff --git a/zig/src/families/dsv41/draft.zig b/zig/src/families/dsv41/draft.zig
index 5ac250b9df..9f7b1a67fb 100644
--- a/zig/src/families/dsv41/draft.zig
+++ b/zig/src/families/dsv41/draft.zig
@@ -105,6 +105,14 @@ const MarkovGather = struct {
 
     fn run(ctx: *anyopaque, send: u64, rows: usize) anyerror!u64 {
         const g: *MarkovGather = @ptrCast(@alignCast(ctx));
+        if (g.e.two) |tw| if (tw.neighbor_gather) |ng| {
+            const neighbor = @import("neighbor.zig");
+            const bytes = try std.math.mul(usize, rows, 16);
+            if (neighbor.eligible(bytes)) {
+                try ng.gather(g.e.s, send, g.dst, bytes);
+                return g.dst;
+            }
+        };
         try g.e.gatherF32(send, g.dst, rows * 4);
         return g.dst;
     }
diff --git a/zig/src/families/dsv41/dsv41.zig b/zig/src/families/dsv41/dsv41.zig
index dabfcec2a1..38ba0e44d8 100644
--- a/zig/src/families/dsv41/dsv41.zig
+++ b/zig/src/families/dsv41/dsv41.zig
@@ -40,6 +40,7 @@ pub const exl3_experts2d = @import("exl3_experts2d.zig");
 pub const exl3_experts2d_rec = @import("exl3_experts2d_rec.zig");
 pub const prompt2d = @import("prompt2d.zig");
 pub const ring2d = @import("ring2d.zig");
+pub const neighbor = @import("neighbor.zig");
 pub const text = @import("text.zig");
 pub const round = @import("round.zig");
 pub const round2d = @import("round2d.zig");
@@ -87,6 +88,7 @@ test {
     refAll(exl3_experts2d_rec);
     refAll(prompt2d);
     refAll(ring2d);
+    refAll(neighbor);
     refAll(text);
     refAll(round);
     refAll(round2d);
diff --git a/zig/src/families/dsv41/model.zig b/zig/src/families/dsv41/model.zig
index 83abb34ee1..210309f8b0 100644
--- a/zig/src/families/dsv41/model.zig
+++ b/zig/src/families/dsv41/model.zig
@@ -28,6 +28,7 @@ const comm_mod = @import("comm.zig");
 const link = @import("link.zig");
 const ring2d = @import("ring2d.zig");
 const rdma = @import("rdma.zig");
+const neighbor = @import("neighbor.zig");
 
 /// The served prompt chunk (TF_DS_PREFILL_CHUNK): a prompt runs in chunks that start at its multiples.
 pub const chunk_rows = 2048;
@@ -157,6 +158,7 @@ pub const Model = struct {
     // world 4: the 2D split's exchanges (prompt2d.zig, TP4's), with the RDMA rings when asked
     two: ?prompt2d.Two = null,
     rdma_rings: ?ring2d.Rings = null,
+    neighbor_gather: ?neighbor.Gather = null,
 
     /// Loads rank o.rank's weights from the lane's rank cache, links the ranks (rank 0 listens on o.port), opens NCCL
     /// and sets up every buffer. Needs `ctx` current on this thread.
@@ -209,6 +211,7 @@ pub const Model = struct {
         m.pgraphs = null;
         m.two = null;
         m.rdma_rings = null;
+        m.neighbor_gather = null;
         const d = ctx.d;
         m.cfg = try Config.read(a, io, o.model_dir);
         const c = &m.cfg;
@@ -244,6 +247,9 @@ pub const Model = struct {
             .pool = o.pool,
         });
         errdefer m.comm.deinit();
+        const neighbor_rdma = try neighbor.configured();
+        try neighbor.agree(&m.comm, m.stream, neighbor_rdma);
+        if (neighbor_rdma and o.rdma_devices != null) return error.NeighborDisallowsStockRdma;
         // TP2: the decode gathers' RDMA ring (both ranks connect before either returns)
         if (sp.pair == null and o.rdma_devices != null) {
             const list = o.rdma_devices.?;
@@ -297,9 +303,15 @@ pub const Model = struct {
         errdefer m.arena.deinit();
         // world 4 (layer_main's setup): the 2D exchanges' buffers, and with rdma_devices their rings (whose infos go
         // round over NCCL world 4)
+        errdefer if (m.neighbor_gather) |*ng| ng.close(gpa);
         errdefer if (m.rdma_rings) |*x| x.close(gpa);
         if (sp.pair != null) {
             m.two = try prompt2d.Two.init(c, &m.w, &m.arena, sp, chunk_rows);
+            if (neighbor_rdma) {
+                m.rdma_mod = try cuda.Module.load(d, cuda.kernels.dsv41_rdma);
+                m.neighbor_gather = try neighbor.Gather.open(gpa, d, try rdma.Kernels.load(m.rdma_mod.?), o.rank, &m.comm, m.stream);
+                m.two.?.neighbor_gather = &m.neighbor_gather.?;
+            }
             if (o.rdma_devices) |devs| {
                 const image: []const u8 = if (o.rdma_kernels) |path|
                     try std.Io.Dir.cwd().readFileAllocOptions(io, path, a, .limited(1 << 26), .@"16", null)
@@ -340,7 +352,7 @@ pub const Model = struct {
                 m.dr.?.graphs = &m.pgraphs.?;
             }
         }
-        try m.stream.synchronize();
+        try m.synchronize();
         try d.check(d.api.cuCtxSynchronize(), "cuCtxSynchronize"); // the setup's legacy-stream copies too
 
         // Engram on the host: the compressed token map, the multipliers, the tables and 64 readers
@@ -402,6 +414,7 @@ pub const Model = struct {
         if (m.paio) |x| x.deinit();
         if (m.epool) |p| p.deinit(m.gpa);
         if (m.tables) |*t| t.close();
+        if (m.neighbor_gather) |*ng| ng.close(m.gpa);
         if (m.rdma_rings) |*r| r.close(m.gpa);
         m.blas.close();
         m.arena.deinit();
@@ -429,6 +442,11 @@ pub const Model = struct {
         m.gpa.destroy(m);
     }
 
+    fn synchronize(m: *Model) !void {
+        try m.stream.synchronize();
+        if (m.neighbor_gather) |*ng| try ng.check();
+    }
+
     /// TP2's RDMA ring for the decode gathers: created on both ranks, the queue-pair infos exchanged over NCCL,
     /// connected, and started once both are (an all-gather as the barrier).
     pub fn openRing(gpa: std.mem.Allocator, d: *const cuda.Driver, k: rdma.Kernels, devices: []const []const u8, rank: u32, comm: *const comm_mod.Comm, stream: cuda.Stream) !*rdma.Ring {
@@ -570,7 +588,7 @@ pub const Model = struct {
             }
             const d = m.ctx.d;
             try d.check(d.api.cuMemcpyHtoDAsync_v2(m.img_buf.?.ptr, rows.ptr, rows.len, m.stream.handle), "cuMemcpyHtoDAsync");
-            try m.stream.synchronize(); // (the caller's rows may go once this returns)
+            try m.synchronize(); // (the caller's rows may go once this returns)
         }
         m.spans.clearRetainingCapacity();
         try m.spans.appendSlice(m.gpa, spans);
@@ -634,7 +652,7 @@ pub const Model = struct {
                     // encoder-only: no synchronize either (as at the chunk's end below), so the next chunk's host steps
                     // run while this one's last layers do; the profile alone waits
                     if (pt) |t| {
-                        try m.stream.synchronize();
+                        try m.synchronize();
                         try t.finish();
                     }
                     return false;
@@ -694,7 +712,7 @@ pub const Model = struct {
         // at 2,048 rows) run while this chunk's kernels do; readers of the results (the prompt's logits) synchronize.
         // The profile (--profile 2) times a chunk to its end, so it alone waits.
         if (pt) |t| {
-            try m.stream.synchronize();
+            try m.synchronize();
             try t.finish();
         }
         return true;
@@ -711,12 +729,12 @@ pub const Model = struct {
     /// overwrites them). absorb[i]: window i's stream drafts.
     pub fn verify(m: *Model, rows: round.Rows, absorb: []const bool) !void {
         const e = &m.eng;
-        if (m.prof) try m.stream.synchronize(); // (the GPU idle, so the issue time shows whether the host keeps it fed)
+        if (m.prof) try m.synchronize(); // (the GPU idle, so the issue time shows whether the host keeps it fed)
         const t0 = m.now();
         try round.forward(e, &m.rd, &m.ch, &m.caches, m.rings, if (m.eh) |*x| x else null, &.{}, rows, m.pool_cap, null, if (m.graphs) |*g| g else null);
         if (m.prof) {
             m.t_enqueue = m.now() - t0;
-            try m.stream.synchronize();
+            try m.synchronize();
             m.t_forward = m.now() - t0;
             if (m.timer) |t| try t.finish();
             if (m.rd.clk != 0) {
@@ -744,7 +762,7 @@ pub const Model = struct {
             }
             if (n > 0) try dr.absorbMany(e, &m.ch, &m.dpool.?, items[0..n]);
         }
-        try m.stream.synchronize();
+        try m.synchronize();
         if (m.prof) m.t_absorb = m.now() - t1;
     }
 
@@ -782,14 +800,14 @@ pub const Model = struct {
     pub fn roundArgmax(m: *Model, R: usize) ![]const u32 {
         const d = m.ctx.d;
         try m.ops.argmaxRows(m.stream, m.rd.logits, m.vocab, m.vocab, R, m.amax);
-        try m.stream.synchronize();
+        try m.synchronize();
         try d.check(d.api.cuMemcpyDtoH_v2(&m.amax_host, m.amax, R * 4), "cuMemcpyDtoH");
         return m.amax_host[0..R];
     }
 
     pub fn roundLogits(m: *Model, R: usize) ![]const f32 {
         const d = m.ctx.d;
-        try m.stream.synchronize();
+        try m.synchronize();
         try d.check(d.api.cuMemcpyDtoH_v2(m.logits.ptr, m.rd.logits, R * m.vocab * 4), "cuMemcpyDtoH");
         return m.logits[0 .. R * m.vocab];
     }
@@ -797,7 +815,7 @@ pub const Model = struct {
     /// The prompt's logits on the host (after a fill's last chunk).
     pub fn promptLogitsHost(m: *Model) ![]const f32 {
         const d = m.ctx.d;
-        try m.stream.synchronize();
+        try m.synchronize();
         try d.check(d.api.cuMemcpyDtoH_v2(m.logits.ptr, m.ch.head_g, m.vocab * 4), "cuMemcpyDtoH");
         return m.logits[0..m.vocab];
     }
diff --git a/zig/src/families/dsv41/neighbor.zig b/zig/src/families/dsv41/neighbor.zig
new file mode 100644
index 0000000000..33d5bc92b2
--- /dev/null
+++ b/zig/src/families/dsv41/neighbor.zig
@@ -0,0 +1,304 @@
+//! TP4 on a physical ring: adjacent xor1 then xor3 exchanges, using the existing RDMA kernels.
+const std = @import("std");
+const cuda = @import("cuda");
+const rdma = @import("rdma.zig");
+const Comm = @import("comm.zig").Comm;
+pub const cutoff: usize = 512 * 1024;
+pub fn configured() !bool {
+    const raw = std.c.getenv("TF_DS_NEIGHBOR_RDMA") orelse return false;
+    const value = std.mem.span(raw);
+    if (std.mem.eql(u8, value, "1")) return true;
+    if (std.mem.eql(u8, value, "0")) return false;
+    return error.BadNeighborMode;
+}
+/// Every TP4 rank agrees before any neighbor queue pair or captured graph is created.
+pub fn agree(comm: *const Comm, stream: cuda.Stream, enabled: bool) !void {
+    if (comm.world != 4) {
+        if (enabled) return error.NotFourNodes;
+        return;
+    }
+    var dev = try cuda.DeviceBuffer.alloc(stream.d, 5);
+    defer dev.free();
+    try dev.upload(0, &.{@intFromBool(enabled)});
+    try comm.allGather(dev.ptr, dev.ptr + 1, 1, .u8, stream);
+    try stream.synchronize();
+    var modes: [4]u8 = undefined;
+    try dev.download(1, &modes);
+    try sameMode(modes);
+}
+
+fn sameMode(modes: [4]u8) !void {
+    for (modes) |mode| if (mode != modes[0]) return error.NeighborModeMismatch;
+}
+
+pub fn eligible(bytes: usize) bool {
+    return bytes > 0 and bytes <= cutoff and bytes % 16 == 0;
+}
+pub const Gather = struct {
+    first: *rdma.Ring,
+    second: *rdma.Ring,
+    scratch: cuda.DeviceBuffer,
+
+    /// Each stage uses two rails in matching order at both ends of its physical edge.
+    pub fn open(gpa: std.mem.Allocator, d: *const cuda.Driver, k: rdma.Kernels, g: u32, comm: *const Comm, stream: cuda.Stream) !Gather {
+        if (g >= 4 or comm.world != 4) return error.NotFourNodes;
+        const first_raw = std.c.getenv("TF_DS_NEIGHBOR_FIRST") orelse return error.NoFirstDevices;
+        const second_raw = std.c.getenv("TF_DS_NEIGHBOR_SECOND") orelse return error.NoSecondDevices;
+        const first_devices = try devices(std.mem.span(first_raw));
+        const second_devices = try devices(std.mem.span(second_raw));
+        for (first_devices) |x| for (second_devices) |y| {
+            if (std.mem.eql(u8, x, y)) return error.NeighborDevicesOverlap;
+        };
+        const s: rdma.Settings = .{ .max_bytes = 2 * cutoff, .gid_index = 3 };
+        const first = try rdma.Ring.create(gpa, d, k, &first_devices, g % 2, 2, s);
+        errdefer first.destroy(gpa);
+        const second = try rdma.Ring.create(gpa, d, k, &second_devices, g / 2, 2, s);
+        errdefer second.destroy(gpa);
+        var scratch = try cuda.DeviceBuffer.alloc(d, 2 * cutoff);
+        errdefer scratch.free();
+        const per = @sizeOf(rdma.Info);
+        const mine = [4]rdma.Info{ first.info(0), first.info(1), second.info(0), second.info(1) };
+        var dev = try cuda.DeviceBuffer.alloc(d, 5 * 4 * per);
+        defer dev.free();
+        try dev.upload(0, std.mem.sliceAsBytes(&mine));
+        try comm.allGather(dev.ptr, dev.ptr + 4 * per, 4 * per, .u8, stream);
+        try stream.synchronize();
+        var all: [4][4]rdma.Info = undefined;
+        try dev.download(4 * per, std.mem.sliceAsBytes(&all));
+        var a: [2]rdma.Info = undefined;
+        var b: [2]rdma.Info = undefined;
+        a[g % 2] = mine[g % 2];
+        a[1 - g % 2] = all[g ^ 1][g % 2];
+        b[g / 2] = mine[2 + g / 2];
+        b[1 - g / 2] = all[g ^ 3][2 + g / 2];
+        try first.connect(&a);
+        try second.connect(&b);
+        try comm.allGather(dev.ptr, dev.ptr + 4 * per, 1, .u8, stream);
+        try stream.synchronize();
+        first.failure_peer = second;
+        second.failure_peer = first;
+        try first.start();
+        errdefer first.stop();
+        try second.start();
+        return .{ .first = first, .second = second, .scratch = scratch };
+    }
+    pub fn close(n: *Gather, gpa: std.mem.Allocator) void {
+        n.first.stop();
+        n.second.stop();
+        n.first.failure_peer = null;
+        n.second.failure_peer = null;
+        n.first.destroy(gpa);
+        n.second.destroy(gpa);
+        n.scratch.free();
+    }
+    pub fn check(n: *const Gather) !void {
+        if (n.first.failure() != null or n.second.failure() != null) return error.NeighborTransportFailed;
+    }
+    /// Opposite rank (xor2) via adjacent xor1 then xor3; only remote slices are copied.
+    pub fn relay(n: *const Gather, stream: cuda.Stream, src: u64, dst: u64, bytes: usize) !void {
+        if (!eligible(bytes)) return error.NeighborSize;
+        try ranges(src, bytes, dst, bytes, n.scratch.ptr, 2 * cutoff, 16);
+        const n4: u32 = @intCast(bytes / 16);
+        var first_dsts: [2]rdma.Ring.Dst = .{ .{}, .{} };
+        var second_dsts: [2]rdma.Ring.Dst = .{ .{}, .{} };
+        first_dsts[1 - n.first.rank] = .{ .ptr = n.scratch.ptr, .row4 = n4, .stride4 = n4 };
+        second_dsts[1 - n.second.rank] = .{ .ptr = dst, .row4 = n4, .stride4 = n4 };
+        try n.check();
+        try n.first.gatherInto(stream, src, 1, n4, &first_dsts);
+        try n.second.gatherInto(stream, n.scratch.ptr, 1, n4, &second_dsts);
+    }
+    pub fn pair(n: *const Gather, stream: cuda.Stream, src: u64, dst: u64, bytes: usize) !void {
+        if (!eligible(bytes)) return error.NeighborSize;
+        try n.check();
+        try n.first.gather(stream, src, @intCast(bytes / 16), dst);
+    }
+    /// Equal-width quarters, final [2, rows, 2*width] byte layout.
+    pub fn quarters(n: *const Gather, stream: cuda.Stream, src: u64, dst: u64, rows: usize, width_bytes: usize) !void {
+        const bytes = try std.math.mul(usize, rows, width_bytes);
+        if (!eligible(bytes) or width_bytes % 16 != 0) return error.NeighborSize;
+        // All validation precedes either exchange; no partial fallback.
+        const rows2: u32 = @intCast(try std.math.mul(usize, rows, 2));
+        const row4: u32 = @intCast(width_bytes / 16);
+        const stride4 = 2 * @as(u64, row4);
+        const dsts = [_]rdma.Ring.Dst{
+            .{ .ptr = dst, .row4 = row4, .stride4 = stride4 },
+            .{ .ptr = dst + width_bytes, .row4 = row4, .stride4 = stride4 },
+        };
+        try n.check();
+        try n.first.gather(stream, src, @intCast(bytes / 16), n.scratch.ptr);
+        try n.second.gatherInto(stream, n.scratch.ptr, rows2, row4, &dsts);
+    }
+    pub fn gather(n: *const Gather, stream: cuda.Stream, src: u64, dst: u64, bytes: usize) !void {
+        if (!eligible(bytes)) return error.NeighborSize;
+        try n.check();
+        try n.first.gather(stream, src, @intCast(bytes / 16), n.scratch.ptr);
+        try n.second.gather(stream, n.scratch.ptr, @intCast(bytes / 8), dst);
+    }
+};
+/// Larger or unequal diagonal parts travel two physical edges by NCCL, never across a missing diagonal.
+pub fn forward(comm: *const Comm, src: u64, dst: u64, counts: [4]usize, dtype: cuda.nccl.DataType, scratch: u64, capacity: usize, stream: cuda.Stream) !void {
+    if (comm.world != 4 or comm.rank >= 4) return error.NotFourNodes;
+    if (dtype != .f16 and dtype != .bf16) return error.DiagonalType;
+    for (counts) |n| {
+        if (n == 0 or n > capacity / 2) return error.DiagonalScratchTooSmall;
+    }
+    const g = comm.rank;
+    try ranges(src, counts[g] * 2, dst, counts[g ^ 2] * 2, scratch, capacity, 2);
+    const next = (g + 1) % 4;
+    const prev = (g + 3) % 4;
+    try hop(comm, src, counts[g], next, scratch, counts[prev], prev, dtype, stream);
+    try hop(comm, scratch, counts[prev], next, dst, counts[g ^ 2], prev, dtype, stream);
+}
+
+fn ranges(src: u64, src_bytes: usize, dst: u64, dst_bytes: usize, scratch: u64, scratch_bytes: usize, alignment: u64) !void {
+    const src_end = try std.math.add(u64, src, src_bytes);
+    const dst_end = try std.math.add(u64, dst, dst_bytes);
+    const scratch_end = try std.math.add(u64, scratch, scratch_bytes);
+    if (src == 0 or dst == 0 or scratch == 0 or src % alignment != 0 or dst % alignment != 0 or
+        scratch % alignment != 0 or (src < dst_end and dst < src_end) or
+        (src < scratch_end and scratch < src_end) or (dst < scratch_end and scratch < dst_end))
+        return error.NeighborAlias;
+}
+
+fn hop(comm: *const Comm, src: u64, count: usize, next: u32, dst: u64, rcount: usize, prev: u32, dtype: cuda.nccl.DataType, stream: cuda.Stream) !void {
+    const lib = &comm.lib;
+    try lib.check(lib.api.ncclGroupStart(), "neighbor group start");
+    const s = lib.api.ncclSend(src, count, dtype, @intCast(next), comm.comm, stream.handle);
+    const r = lib.api.ncclRecv(dst, rcount, dtype, @intCast(prev), comm.comm, stream.handle);
+    const end = lib.api.ncclGroupEnd();
+    try lib.check(s, "neighbor send");
+    try lib.check(r, "neighbor receive");
+    try lib.check(end, "neighbor group end");
+}
+
+fn devices(raw: []const u8) ![2][]const u8 {
+    var it = std.mem.splitScalar(u8, raw, ',');
+    const a = it.next() orelse return error.BadNeighborDevices;
+    const b = it.next() orelse return error.BadNeighborDevices;
+    if (a.len == 0 or b.len == 0 or it.next() != null or std.mem.eql(u8, a, b)) return error.BadNeighborDevices;
+    return .{ a, b };
+}
+test "neighbor stages produce global rank order and bounded dispatch" {
+    for (0..4) |g| {
+        var output: [4]usize = undefined;
+        for (0..2) |p| {
+            const owner = if (p == g / 2) g else g ^ 3;
+            output[2 * p] = owner & ~@as(usize, 1);
+            output[2 * p + 1] = owner | 1;
+        }
+        try std.testing.expectEqual([4]usize{ 0, 1, 2, 3 }, output);
+    }
+    try std.testing.expect(eligible(cutoff));
+    try std.testing.expect(!eligible(cutoff + 16));
+    try std.testing.expect(!eligible(15));
+    try std.testing.expect(!eligible(0));
+}
+
+test "neighbor modes agree before rings open" {
+    try sameMode(.{ 0, 0, 0, 0 });
+    try sameMode(.{ 1, 1, 1, 1 });
+    for (0..4) |g| {
+        var modes: [4]u8 = @splat(0);
+        modes[g] = 1;
+        try std.testing.expectError(error.NeighborModeMismatch, sameMode(modes));
+    }
+}
+
+test "a stage requires two distinct device names" {
+    const names = try devices("mlx5_0,mlx5_1");
+    try std.testing.expectEqualStrings("mlx5_0", names[0]);
+    try std.testing.expectEqualStrings("mlx5_1", names[1]);
+    for ([_][]const u8{ "", "mlx5_0", "mlx5_0,", ",mlx5_1", "mlx5_0,mlx5_0", "mlx5_0,mlx5_1,mlx5_2" }) |bad| {
+        try std.testing.expectError(error.BadNeighborDevices, devices(bad));
+    }
+}
+
+test "the GPU and verbs declarations compile on the host" {
+    std.testing.refAllDecls(@This());
+    std.testing.refAllDecls(Gather);
+}
+
+test "invalid relay shapes and buffers fail before either stage is enqueued" {
+    var n: Gather = undefined;
+    n.scratch.ptr = 0x100000;
+    const stream: cuda.Stream = undefined;
+    try std.testing.expectError(error.NeighborSize, n.relay(stream, 16, 32, 0));
+    try std.testing.expectError(error.NeighborSize, n.relay(stream, 16, 32, cutoff + 16));
+    try std.testing.expectError(error.NeighborAlias, n.relay(stream, 16, 32, 32));
+    try std.testing.expectError(error.NeighborAlias, n.relay(stream, 17, 64, 16));
+    try std.testing.expectError(error.NeighborAlias, n.relay(stream, n.scratch.ptr, 64, 16));
+    try std.testing.expectError(error.Overflow, n.relay(stream, std.math.maxInt(u64) - 15, 32, 16));
+}
+
+test "diagonal fallback validates every count before its first NCCL group" {
+    var comm: Comm = undefined;
+    comm.rank = 0;
+    comm.world = 4;
+    const stream: cuda.Stream = undefined;
+    try std.testing.expectError(error.DiagonalType, forward(&comm, 16, 64, .{ 1, 1, 1, 1 }, .f32, 128, 32, stream));
+    try std.testing.expectError(error.DiagonalScratchTooSmall, forward(&comm, 16, 64, .{ 1, 1, 1, 17 }, .f16, 128, 32, stream));
+    try std.testing.expectError(error.NeighborAlias, forward(&comm, 16, 64, .{ 1, 1, 1, 1 }, .bf16, 16, 32, stream));
+    comm.world = 2;
+    try std.testing.expectError(error.NotFourNodes, forward(&comm, 16, 64, .{ 1, 1, 1, 1 }, .f16, 128, 32, stream));
+}
+
+const ForwardTest = struct {
+    const Call = struct { send: bool, ptr: u64, count: usize, peer: c_int };
+    var calls: [4]Call = undefined;
+    var used: usize = 0;
+    var groups: usize = 0;
+    var depth: usize = 0;
+
+    fn start() callconv(.c) c_int {
+        std.debug.assert(depth == 0);
+        depth = 1;
+        groups += 1;
+        return 0;
+    }
+    fn end() callconv(.c) c_int {
+        std.debug.assert(depth == 1);
+        depth = 0;
+        return 0;
+    }
+    fn send(ptr: u64, count: usize, dtype: cuda.nccl.DataType, peer: c_int, _: cuda.nccl.Comm, _: cuda.abi.Stream) callconv(.c) c_int {
+        std.debug.assert(depth == 1 and dtype == .bf16);
+        calls[used] = .{ .send = true, .ptr = ptr, .count = count, .peer = peer };
+        used += 1;
+        return 0;
+    }
+    fn recv(ptr: u64, count: usize, dtype: cuda.nccl.DataType, peer: c_int, _: cuda.nccl.Comm, _: cuda.abi.Stream) callconv(.c) c_int {
+        _ = send(ptr, count, dtype, peer, null, null);
+        calls[used - 1].send = false;
+        return 0;
+    }
+};
+
+test "production diagonal fallback forwards unequal parts on physical edges at every rank" {
+    const F = ForwardTest;
+    var comm: Comm = undefined;
+    comm.world = 4;
+    comm.comm = null;
+    comm.lib.api.ncclGroupStart = F.start;
+    comm.lib.api.ncclGroupEnd = F.end;
+    comm.lib.api.ncclSend = F.send;
+    comm.lib.api.ncclRecv = F.recv;
+    const stream: cuda.Stream = .{ .d = undefined, .handle = null };
+    const counts = [4]usize{ 3, 3, 5, 5 };
+    for (0..4) |g| {
+        comm.rank = @intCast(g);
+        F.used = 0;
+        F.groups = 0;
+        try forward(&comm, 16, 64, counts, .bf16, 128, 32, stream);
+        const next: c_int = @intCast((g + 1) % 4);
+        const prev: c_int = @intCast((g + 3) % 4);
+        const expected = [_]F.Call{
+            .{ .send = true, .ptr = 16, .count = counts[g], .peer = next },
+            .{ .send = false, .ptr = 128, .count = counts[@intCast(prev)], .peer = prev },
+            .{ .send = true, .ptr = 128, .count = counts[@intCast(prev)], .peer = next },
+            .{ .send = false, .ptr = 64, .count = counts[g ^ 2], .peer = prev },
+        };
+        try std.testing.expectEqual(@as(usize, 2), F.groups);
+        try std.testing.expectEqual(@as(usize, 0), F.depth);
+        try std.testing.expectEqualSlices(F.Call, &expected, F.calls[0..F.used]);
+    }
+}
diff --git a/zig/src/families/dsv41/prompt2d.zig b/zig/src/families/dsv41/prompt2d.zig
index 6c1106a041..32ba6474d0 100644
--- a/zig/src/families/dsv41/prompt2d.zig
+++ b/zig/src/families/dsv41/prompt2d.zig
@@ -12,6 +12,7 @@
 //! prompt.zig takes it through thin hooks behind Engine.two; a TP2 engine has none and runs as before. The decode
 //! rounds' part is round2d.zig's, on this node state.
 const std = @import("std");
+const neighbor = @import("neighbor.zig");
 const cuda = @import("cuda");
 const Config = @import("config.zig").Config;
 const plan = @import("plan.zig");
@@ -183,6 +184,7 @@ fn len(r: ?[2]usize) usize {
 /// A 2D node's part of the prompt path and the decode rounds: the pairs' widths of each split output and the exchange
 /// buffers.
 pub const Two = struct {
+    neighbor_gather: ?*const neighbor.Gather,
     r: usize,
     p: usize,
     heads: usize, // the heads here (its pair's half of the rank's)
@@ -215,6 +217,7 @@ pub const Two = struct {
         const x0 = s0.expertParts(c.*);
         const x1 = s1.expertParts(c.*);
         var t: Two = undefined;
+        t.neighbor_gather = null;
         t.r = s.rank;
         t.p = pp;
         t.heads = s.headSpan(c.*)[1];
@@ -286,8 +289,19 @@ pub const Two = struct {
                 else => return err,
             }
         }
+        const count = try std.math.mul(usize, rows, width);
+        if (t.neighbor_gather) |ng| {
+            const bytes = try std.math.mul(usize, count, 2);
+            if (neighbor.eligible(bytes)) {
+                if (bytes > t.part_bytes) return error.PartTooBig;
+                try ng.relay(e.s, mine, t.peer, bytes);
+            } else {
+                try neighbor.forward(e.comm, mine, t.peer, @splat(count), .f16, t.pad, t.part_bytes, e.s);
+            }
+            return;
+        }
         const g: u32 = @intCast(t.r + 2 * t.p);
-        try e.comm.exchange(mine, rows * width, t.peer, rows * width, .f16, g ^ 2, e.s);
+        try e.comm.exchange(mine, count, t.peer, count, .f16, g ^ 2, e.s);
     }
 
     /// moe_2d's parity experts after the main gate / up and the rest's: the packs exchanged with the column partner,
@@ -311,7 +325,14 @@ pub const Two = struct {
             try e.ops.copyRows(e.s, src, w[t.p] * esize, t.pad, wmax * esize, w[t.p] * esize, n);
             send = t.pad;
         }
-        try e.comm.allGather(send, t.parts, n * wmax * esize, .u8, e.s);
+        const bytes = n * wmax * esize;
+        if (t.neighbor_gather) |ng| {
+            if (neighbor.eligible(bytes)) {
+                try ng.gather(e.s, send, t.parts, bytes);
+                return wmax;
+            }
+        }
+        try e.comm.allGather(send, t.parts, bytes, .u8, e.s);
         return wmax;
     }
 
@@ -328,6 +349,14 @@ pub const Two = struct {
     /// [2, n, w0 + w1]: rank r's partial, pair 0's columns then pair 1's, for every r.
     pub fn quarters(t: *const Two, e: *const prompt.Engine, src: u64, dst: u64, n: usize, w: [2]usize, esize: usize) !void {
         if (t.rings) |rs| if (n < 64 and try ringed(rs.quarters(e.s, src, dst, n, w, esize))) return;
+        if (t.neighbor_gather) |ng| {
+            const width = try std.math.mul(usize, w[0], esize);
+            const bytes = try std.math.mul(usize, n, width);
+            if (w[0] == w[1] and width % 16 == 0 and neighbor.eligible(bytes)) {
+                try ng.quarters(e.s, src, dst, n, width);
+                return;
+            }
+        }
         const wmax = try t.allParts(e, src, n, w, esize);
         const row = (w[0] + w[1]) * esize;
         for (0..4) |g| {
@@ -354,6 +383,13 @@ pub const Two = struct {
     /// (Comm2D.gather4(x)[2p:2p + 2]).
     pub fn pairGather(t: *const Two, e: *const prompt.Engine, src: u64, dst: u64, n: usize, width: usize, esize: usize) !void {
         if (t.rings) |rs| if (try ringed(rs.pairGather(e.s, src, dst, n, width, esize))) return;
+        if (t.neighbor_gather) |ng| {
+            const bytes = try std.math.mul(usize, n, try std.math.mul(usize, width, esize));
+            if (neighbor.eligible(bytes)) {
+                try ng.pair(e.s, src, dst, bytes);
+                return;
+            }
+        }
         const wmax = try t.allParts(e, src, n, .{ width, width }, esize);
         try e.ops.copyRows(e.s, t.parts + 2 * t.p * n * wmax * esize, wmax * esize, dst, width * esize, width * esize, 2 * n);
     }
@@ -396,7 +432,11 @@ pub const Two = struct {
         const g: u32 = @intCast(t.r + 2 * t.p);
         try e.comm.exchange(src, mine, other, mine, .f32, g ^ 1, e.s);
         try e.ops.add2Bf16(e.s, src, other, half, mine);
-        try e.comm.exchange(half, mine, their, n * w[1 - t.p], .bf16, g ^ 2, e.s);
+        if (t.neighbor_gather != null) {
+            try neighbor.forward(e.comm, half, their, .{ n * w[0], n * w[0], n * w[1], n * w[1] }, .bf16, t.pad, t.part_bytes, e.s);
+        } else {
+            try e.comm.exchange(half, mine, their, n * w[1 - t.p], .bf16, g ^ 2, e.s);
+        }
         const row = (w[0] + w[1]) * 2;
         for (0..2) |pg| {
             const from = if (pg == t.p) half else their;
diff --git a/zig/src/families/dsv41/rdma.zig b/zig/src/families/dsv41/rdma.zig
index dfc06a9c6a..73ff3ee63b 100644
--- a/zig/src/families/dsv41/rdma.zig
+++ b/zig/src/families/dsv41/rdma.zig
@@ -145,6 +145,7 @@ pub const Ring = struct {
     nd: u32 = 0,
     peers: [max_ranks]Peer = @splat(.{}),
     proxy: ?std.Thread = null,
+    failure_peer: ?*Ring = null,
     running: std.atomic.Value(bool) = .init(false),
     failed: std.atomic.Value(bool) = .init(false),
     why: [128]u8 = @splat(0),
@@ -310,6 +311,11 @@ pub const Ring = struct {
     }
 
     fn fail(r: *Ring, why: []const u8) void {
+        r.failLocal(why);
+        if (r.failure_peer) |peer| peer.failLocal(why);
+    }
+
+    fn failLocal(r: *Ring, why: []const u8) void {
         if (!r.failed.swap(true, .acq_rel)) {
             const n = @min(why.len, r.why.len - 1);
             @memcpy(r.why[0..n], why[0..n]);
@@ -519,3 +525,24 @@ test "the info a peer reads is a plain extern struct (it crosses the link as byt
     try std.testing.expectEqual(@as(usize, 128), @sizeOf(Info));
     try std.testing.expectEqual(@as(usize, 120), @offsetOf(Info, "recv"));
 }
+
+test "a failed neighbor stage releases both local stages' device waits" {
+    var first: Ring = undefined;
+    var second: Ring = undefined;
+    var first_meta: [32]u64 align(std.heap.page_size_min) = @splat(0);
+    var second_meta: [32]u64 align(std.heap.page_size_min) = @splat(0);
+    for ([_]*Ring{ &first, &second }, [_][]u64{ &first_meta, &second_meta }) |r, meta| {
+        r.s.slots = 4;
+        r.world = 2;
+        r.meta.mem = @alignCast(std.mem.sliceAsBytes(meta));
+        r.failed = .init(false);
+        r.why = @splat(0);
+    }
+    first.failure_peer = &second;
+    second.failure_peer = &first;
+    first.abort("peer failed");
+    try std.testing.expectEqualStrings("peer failed", first.failure().?);
+    try std.testing.expectEqualStrings("peer failed", second.failure().?);
+    try std.testing.expectEqual(@as(u64, 1), first.abortWord().*);
+    try std.testing.expectEqual(@as(u64, 1), second.abortWord().*);
+}

From d876b6c653bb359a62d7fd0679bd83621e015328 Mon Sep 17 00:00:00 2001
From: Scott Woolsey Biggart <297695+swbiggart@users.noreply.github.com>
Date: Mon, 19 Oct 2026 19:37:46 -0700
Subject: [PATCH 2/5] docs(dsv41): record consolidated neighbor hardware
 qualification

---
 .../dsv41-neighbor-hardware/README.md         | 106 +++
 .../dsv41-neighbor-hardware/summary.json      | 132 +++
 .../dsv41-neighbor-hardware/transport.json    | 590 ++++++++++++
 .../dsv41-neighbor-hardware/workloads.json    | 864 ++++++++++++++++++
 docs/receipts/dsv41-neighbor-host/README.md   |  10 +-
 docs/recipes/deepseek-v41-neighbor.md         |   5 +
 6 files changed, 1702 insertions(+), 5 deletions(-)
 create mode 100644 docs/receipts/dsv41-neighbor-hardware/README.md
 create mode 100644 docs/receipts/dsv41-neighbor-hardware/summary.json
 create mode 100644 docs/receipts/dsv41-neighbor-hardware/transport.json
 create mode 100644 docs/receipts/dsv41-neighbor-hardware/workloads.json

diff --git a/docs/receipts/dsv41-neighbor-hardware/README.md b/docs/receipts/dsv41-neighbor-hardware/README.md
new file mode 100644
index 0000000000..669f129e6e
--- /dev/null
+++ b/docs/receipts/dsv41-neighbor-hardware/README.md
@@ -0,0 +1,106 @@
+# TP4 neighbor transport: four-host hardware qualification
+
+Tested 2026-10-20 UTC at code commit
+`a9225d72ec63b19867c72eba1cca48a49648805c`, based on
+`deepseek-v41-zig-tp4` at `59402086381a42c74d53053443224eec2790b4c5`.
+Subsequent receipt-only changes do not change the tested engine code.
+
+We do not have a network switch. Four GB10 hosts form a physical ring, with two
+PCI-domain RoCE rails per adjacent edge and GID 3. This was the fastest ring
+implementation selected in our experiments. These results are **not a direct
+performance comparison with the target branch or with a switched/full-mesh
+fabric**. The attempted target-branch comparison was discontinued; its partial
+measurements are not included here.
+
+## Build and runtime
+
+Both native executables were built with Zig 0.17.0, Linux/aarch64 glibc 2.28,
+fast optimization, SM121 and the existing CUDA fatbins. The build used the pinned
+NGC image `nvcr.io/nvidia/pytorch@sha256:7531d90bcbe0e43e1f7363029c7e145ce90eebeb494a7b4695fdba0329d7c3c3`:
+
+```sh
+zig build native neighbor-lanes neighbor-probe \
+  -Dtarget=aarch64-linux-gnu.2.28 -Doptimize=fast -Dsm=121 -j2 -Dfatbins=/fatbins
+```
+
+`neighbor-lanes` and `neighbor-probe` were external qualification build targets,
+not release targets added to this branch. The latter exposes a C ABI around the
+production code for a weight-free Python/CUDA driver. Only build target additions
+and the new adapter were applied to the source archive; all 423 tracked files
+under `zig/src` were verified byte-identical to the tested commit. The native
+executables have no test-path substitutions. Binary digests and per-rank OS,
+driver and kernel versions are in [summary.json](summary.json).
+
+Model: DeepSeek-V4.1-Flash-1017 EXL3 revision
+`a2f91e2794923d0beebebe0d267fb3468eaf0b65`, original Engram revision
+`1bae0e1088993b9b82de08f59de5be5a34e42170`, DSpark.
+Effective server arguments: `--context 1048576 --parallel 16 --max-tokens 32768`,
+shared window 2097152, round rows 48, graphs enabled, served-depth policy `solo`.
+`TF_DS_NEIGHBOR_RDMA=1` on all four ranks; no per-optimization switches or stock
+`TF_RDMA_DEVICES`. NCCL 2.30.7 with the existing switchless-routing patches,
+IB network, Ring algorithm, four channels, and LL128 excluded. No forced-tool,
+model arithmetic, or server patches were added to the tested commit.
+
+## Passed checks
+
+- **56 rank/shape cases:** independent NCCL all-gather byte oracle, output guards,
+  changing-input CUDA graph replay, and mixed operation graphs on all four ranks.
+  The adapter calls production pair/quarter/diagonal operations and the full
+  gather helper used by Markov. Sizes include the 512 KiB cutoff and alignment
+  fallbacks. Four-rank startup mode agreement also executed successfully.
+- **9 short model replies:** three prompts through serial, drafted and concurrent
+  paths, 64 generated tokens each; exact token IDs and UTF-8 agree.
+- **24 benchmark replies:** four requests × six timed bursts, 256 tokens each;
+  UTF-8 agrees across repetitions for each request.
+- **11 longer-prompt replies:** four approximately 8K-token prompts (requested
+  lengths 8192), plus a seeded-sampling case, burst/solo replay and one serial
+  reference. All 5/5 burst/solo cases match exact token IDs; drafting matches
+  serial. All four recall markers were found in both greedy paths. Solo replays
+  reported 6144 cached tokens each. This exercises prompt chunking and retained
+  prefixes, not the full configured context.
+- **Cancellation:** disconnect after an active content chunk, then exact equality
+  of the next serial and drafted 64-token replies.
+- **Sustained serving:** 90-second four-stream run, 62 requests, no worker errors.
+
+Transport timings below are complete-operation candidate latencies: slowest
+rank's mean of two measurement rounds. They are not a target-branch comparison.
+
+| Diagonal payload | Latency |
+|---|---:|
+| 8.75 KiB | 17.81 µs |
+| 52.5 KiB | 27.77 µs |
+| 140 KiB | 42.89 µs |
+| 420 KiB | 87.10 µs |
+
+[transport.json](transport.json) contains every candidate rank/shape result.
+
+## Serving measurements
+
+The recorded workload is in [workloads.json](workloads.json). Greedy decoding,
+four streams and 256-token maximum outputs were used. Set-a is the same public
+prose/code/structured-output workload used during transport selection. One warmup
+preceded six measured bursts; steady admission used set-a for 30 seconds. The
+separate 90-second sustained run used set-b. Different workloads are not directly
+comparable, and these short measurements are not confidence estimates.
+
+| Workload | Aggregate tokens/s | Detail |
+|---|---:|---|
+| Set-a four-stream bursts | **132.63** | Mean of six; range 131.94–133.31 |
+| Set-a steady admission | **148.83** | 30 seconds, 19 requests |
+| Set-b sustained | **169.77** | 90 seconds, 62 requests; TTFT p50 0.159 s, p95 0.283 s |
+
+## Host checks and limits
+
+A fresh `zig build test-dsv41-neighbor --cache-dir /tmp/neighbor-host-cache --summary all`
+passed **14/14**. `zig build test test-golden -Dcpu=apple_m1 --summary all` again
+reported the inherited macOS `std.c.SO.BUSY_POLL` compile blocker; the frozen
+corpus remained 312/319 equal, seven known differences, zero unexpected. The full
+host command is not reported as a pass. See the [host receipt](../dsv41-neighbor-host/README.md)
+for the pinned-base comparison and inherited allocator/lean findings.
+
+Full-1M quality, 160K cold prefill, 16-way sustained load, fault injection,
+restart/reboot and long-duration reliability were not tested on this consolidated
+branch. Unequal-width diagonal forwarding has host call-sequence coverage; this
+GPU probe does not claim exhaustive model shapes or distributed failure handling.
+No target-branch speedup is claimed. The previously deployed build was resumed
+after successful bounded qualification; the PR was not deployed as a replacement.
diff --git a/docs/receipts/dsv41-neighbor-hardware/summary.json b/docs/receipts/dsv41-neighbor-hardware/summary.json
new file mode 100644
index 0000000000..1c5a1bda40
--- /dev/null
+++ b/docs/receipts/dsv41-neighbor-hardware/summary.json
@@ -0,0 +1,132 @@
+{
+  "tested_commit": "a9225d72ec63b19867c72eba1cca48a49648805c",
+  "target_commit": "59402086381a42c74d53053443224eec2790b4c5",
+  "direct_target_comparison": false,
+  "reason": "No network switch is available; these are measurements of the fastest implementation selected in our ring experiments, not a direct target-branch or switched-fabric comparison.",
+  "hardware": {
+    "gpu": "NVIDIA GB10",
+    "architecture": "aarch64",
+    "ranks": [
+      {
+        "rank": 0,
+        "kernel": "6.17.0-1029-nvidia",
+        "ubuntu": "24.04.4",
+        "driver": "580.173.02",
+        "management_mbps": 10000
+      },
+      {
+        "rank": 1,
+        "kernel": "6.17.0-1032-nvidia",
+        "ubuntu": "24.04.3",
+        "driver": "580.173.02",
+        "management_mbps": 10000
+      },
+      {
+        "rank": 2,
+        "kernel": "6.17.0-1022-nvidia",
+        "ubuntu": "24.04.4",
+        "driver": "580.159.03",
+        "management_mbps": 2500
+      },
+      {
+        "rank": 3,
+        "kernel": "6.17.0-1022-nvidia",
+        "ubuntu": "24.04.4",
+        "driver": "580.159.03",
+        "management_mbps": 2500
+      }
+    ],
+    "nccl": "2.30.7, patched",
+    "nccl_sha256": "1cdf5f3e2d617d5c07a87903eb8849d5bd7758f6c900455d613a2e0eb289d63e"
+  },
+  "binaries": {
+    "tensorfold-native": "ff0a33c176e621115a4f98d662f56deed7ba54c1e005ac5b8742cf17a6279cae",
+    "libtf-neighbor.so": "38867278afe5e9560d64e81b1f1ce6bf1a953713a5243057bb0e4ff14ff8e00a",
+    "probe.py": "8d1900817f5e8b240c4be327fef72ca2f5bd187701c109f64007cc34b9d4dba2",
+    "tf-dsv41-lanes": "fb85c5476b7ef946bf91de6fd3ef9380ab0f04d2e809abeaf050c7915dfccc69"
+  },
+  "transport": {
+    "rank_shape_cases": 56,
+    "pass_bytes_and_graphs": true,
+    "latencies": [
+      {
+        "bytes": 8960,
+        "slowest_rank_mean_us": 17.814800143241882
+      },
+      {
+        "bytes": 53760,
+        "slowest_rank_mean_us": 27.772000432014465
+      },
+      {
+        "bytes": 143360,
+        "slowest_rank_mean_us": 42.89499968290329
+      },
+      {
+        "bytes": 430080,
+        "slowest_rank_mean_us": 87.09779977798462
+      }
+    ]
+  },
+  "exact_tokens": {
+    "prompts": 3,
+    "paths": 3,
+    "tokens_per_reply": 64,
+    "replies": 9,
+    "pass_test": true
+  },
+  "long_context": {
+    "burst_equal_to_solo": 5,
+    "of": 5,
+    "drafted_equals_serial": true,
+    "needles_found_burst": 4,
+    "needles_found_solo": 4,
+    "prompts": 4,
+    "burst_s": 17.4,
+    "solo_s": 7.7,
+    "serial_s": 1.8,
+    "cached_tokens_solo": [
+      6144,
+      6144,
+      6144,
+      6144,
+      6144
+    ]
+  },
+  "cancellation_pass": true,
+  "burst": {
+    "streams": 4,
+    "tokens": 256,
+    "warmups": 1,
+    "repetitions": 6,
+    "mean_aggregate_tps": 132.6345814307147,
+    "min_aggregate_tps": 131.9383995892958,
+    "max_aggregate_tps": 133.30848622707344,
+    "exact_utf8_replies": 24
+  },
+  "steady": {
+    "streams": 4,
+    "seconds": 30,
+    "max_tokens": 256,
+    "set": "a",
+    "requests": 19,
+    "aggregate_tps": 148.83,
+    "per_stream_tps_p50": 47.475,
+    "per_stream_tps_p95": 62.115,
+    "per_stream_tps_p05": 22.925,
+    "ttft_s_p50": 0.189,
+    "ttft_s_p95": 0.49
+  },
+  "sustained": {
+    "streams": 4,
+    "seconds": 90,
+    "max_tokens": 256,
+    "set": "b",
+    "requests": 62,
+    "aggregate_tps": 169.77,
+    "per_stream_tps_p50": 51.476,
+    "per_stream_tps_p95": 82.968,
+    "per_stream_tps_p05": 25.957,
+    "ttft_s_p50": 0.159,
+    "ttft_s_p95": 0.283
+  }
+}
diff --git a/docs/receipts/dsv41-neighbor-hardware/transport.json b/docs/receipts/dsv41-neighbor-hardware/transport.json
new file mode 100644
index 0000000000..4c8bad74fa
--- /dev/null
+++ b/docs/receipts/dsv41-neighbor-hardware/transport.json
@@ -0,0 +1,590 @@
+[
+  {
+    "rank": 0,
+    "pass_gate": true,
+    "mixed_graph": true,
+    "records": [
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 409600,
+        "exact": true,
+        "candidate_us": [
+          30.493998527526855,
+          30.445200204849243
+        ]
+      },
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.60559749603271,
+          202.56800651550293
+        ]
+      },
+      {
+        "kind": 1,
+        "rows": 4,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.526399970054626,
+          14.788000285625458
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          18.897199630737305,
+          18.829600512981415
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 48,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          128.52879762649536,
+          122.29399681091309
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          226.56400203704834,
+          198.59280586242676
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 7,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          17.79559999704361,
+          17.834000289440155
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 42,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          27.924799919128418,
+          27.32599973678589
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 112,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          46.43999934196472,
+          39.35000002384186
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 336,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          92.94919967651367,
+          80.39720058441162
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.654800295829773,
+          14.564000070095062
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524288,
+        "exact": true,
+        "candidate_us": [
+          94.62159872055054,
+          94.88400220870972
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.54680347442627,
+          216.4423942565918
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 18,
+        "exact": true,
+        "candidate_us": [
+          109.96840000152588,
+          109.81719493865967
+        ]
+      }
+    ]
+  },
+  {
+    "rank": 1,
+    "pass_gate": true,
+    "mixed_graph": true,
+    "records": [
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 409600,
+        "exact": true,
+        "candidate_us": [
+          30.48959970474243,
+          30.421200394630432
+        ]
+      },
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.48959493637085,
+          202.37560272216797
+        ]
+      },
+      {
+        "kind": 1,
+        "rows": 4,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.542800188064575,
+          14.771999418735504
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          18.901599943637848,
+          18.756400048732758
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 48,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          128.51239442825317,
+          121.68760299682617
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          226.52881145477295,
+          198.57800006866455
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 7,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          17.684000730514526,
+          17.757999897003174
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 42,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          28.045600652694702,
+          27.455198764801025
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 112,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          46.21039927005768,
+          39.081600308418274
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 336,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          94.0060019493103,
+          80.18959760665894
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.591200649738312,
+          14.52919989824295
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524288,
+        "exact": true,
+        "candidate_us": [
+          94.61640119552612,
+          95.49999833106995
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.21360158920288,
+          217.42920875549316
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 18,
+        "exact": true,
+        "candidate_us": [
+          110.02440452575684,
+          109.80960130691528
+        ]
+      }
+    ]
+  },
+  {
+    "rank": 2,
+    "pass_gate": true,
+    "mixed_graph": true,
+    "records": [
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 409600,
+        "exact": true,
+        "candidate_us": [
+          30.64880073070526,
+          30.66360056400299
+        ]
+      },
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.56400537490845,
+          202.4616003036499
+        ]
+      },
+      {
+        "kind": 1,
+        "rows": 4,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.658400416374207,
+          14.724799990653992
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          18.907199800014496,
+          18.756400048732758
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 48,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          128.5156011581421,
+          121.51639461517334
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          226.47359371185303,
+          198.57200384140015
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 7,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          17.76439994573593,
+          17.847999930381775
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 42,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          28.082001209259033,
+          27.461999654769897
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 112,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          46.42600119113922,
+          39.333200454711914
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 336,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          94.12440061569214,
+          79.66960072517395
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.567600190639496,
+          14.519999921321869
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524288,
+        "exact": true,
+        "candidate_us": [
+          95.23000121116638,
+          95.53840160369873
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.3708028793335,
+          217.13440418243408
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 18,
+        "exact": true,
+        "candidate_us": [
+          110.25960445404053,
+          109.65759754180908
+        ]
+      }
+    ]
+  },
+  {
+    "rank": 3,
+    "pass_gate": true,
+    "mixed_graph": true,
+    "records": [
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 409600,
+        "exact": true,
+        "candidate_us": [
+          30.607199668884277,
+          30.65280020236969
+        ]
+      },
+      {
+        "kind": 0,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.57120561599731,
+          202.49919891357422
+        ]
+      },
+      {
+        "kind": 1,
+        "rows": 4,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.633199572563171,
+          14.882799983024597
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          18.878400325775146,
+          18.8060000538826
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 48,
+        "width": 10240,
+        "exact": true,
+        "candidate_us": [
+          128.46280336380005,
+          122.10240364074707
+        ]
+      },
+      {
+        "kind": 2,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          226.35560035705566,
+          198.80399703979492
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 7,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          17.788800597190857,
+          17.838799953460693
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 42,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          28.051599860191345,
+          27.460798621177673
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 112,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          46.43479883670807,
+          39.18159902095795
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 336,
+        "width": 1280,
+        "exact": true,
+        "candidate_us": [
+          93.40720176696777,
+          79.79999780654907
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 16,
+        "exact": true,
+        "candidate_us": [
+          14.601600170135498,
+          14.541999995708466
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524288,
+        "exact": true,
+        "candidate_us": [
+          95.30079960823059,
+          95.05879878997803
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 524304,
+        "exact": true,
+        "candidate_us": [
+          185.6376051902771,
+          216.42999649047852
+        ]
+      },
+      {
+        "kind": 3,
+        "rows": 1,
+        "width": 18,
+        "exact": true,
+        "candidate_us": [
+          110.12599468231201,
+          109.82120037078857
+        ]
+      }
+    ]
+  }
+]
diff --git a/docs/receipts/dsv41-neighbor-hardware/workloads.json b/docs/receipts/dsv41-neighbor-hardware/workloads.json
new file mode 100644
index 0000000000..fa49c61a50
--- /dev/null
+++ b/docs/receipts/dsv41-neighbor-hardware/workloads.json
@@ -0,0 +1,864 @@
+{
+  "set-a": {
+    "aad77d0c88401c4ae209805bc3043a877c99318a25f59e626cb47031e167bfe0": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Return a JSON array of 15 fictional employees. Each object has the fields id (int), name, department, title, salary (int), start_date (YYYY-MM-DD), manager_id (int or null) and skills (array of strings). Output only the JSON, pretty-printed with two-space indents. (variant 2)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "e3c84d75cf8e397ef7fb1325767f28956aa28dee34b78feb5f0f93caaf74b8aa": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a complete Python implementation of a thread-safe LRU cache with per-entry TTL expiry. Include type hints, docstrings, and a set of unittest test cases covering eviction and expiry. (variant 0)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "f9f80b2afdd2bde32b848b34a4e590ba3f672968bb2f192934c173268476b804": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a complete Python implementation of a thread-safe LRU cache with per-entry TTL expiry. Include type hints, docstrings, and a set of unittest test cases covering eviction and expiry. (variant 3)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "a89565c2f6f30f4f810ebf439719bfa90d2438914a43da8f71ced3c047dc380e": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a long, vivid essay about the history of lighthouses: how they were built, how their lamps and lenses evolved, and the lives of the keepers who tended them through storms. (variant 1)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    }
+  },
+  "set-b": {
+    "74d158f3297d120977d203b27efc15bcebeb9ce9d50d17403bd7235b96c73abf": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 2)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "028435d19643b27ced3885536e85042449a68642389143bb64b7d8ffc1dd3a54": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 0)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "760e69453ee1cf28d8e9f40bb4a1cf6e1e9b79abf6074c01d279d5f55cf9b777": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 3)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "de33478bcdca36eaadd4cc13f71c3ad350a765844b9c9813355ca8b4cc131cd7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 5)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "bdbdd70cc48017b4527e14d4ab31a59d661dd5d8022d0000808fe5de532569ed": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 1)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "beef6bcf1f303b6517c9cafd60acfbc884cc2d556b3472532189a724ddeb7878": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 6)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "c053f7d0c8f3fdcef1c1eb5b52897e5fc2beef8ebac60634c080b28d517821a3": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 4)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "ee688488e64bc3dcd37d3e70ffb8db08cfea1ba1bf941f949092ee9cccbdcbb6": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 8)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "2ed4b9756bf08f14734a74e6a410cc26e9ced4b0f0a65036d2537e41f957c3d5": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 9)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "6057b4a21bf73ada4b79a4039a2ba74c7446aee7e80ce7cc145291c64334ffcf": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 11)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "39254edafb61cbe2abad3bf00bc6b722b15e8675567f3299d93c6cda76641cd7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 7)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "f2556a3d3f1bb1544c1d3a0726a95bea91e7d1257c3d5c0e561e1ee4c148930b": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 12)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "c9832353d1b5d1a41c571abf029044223e632416e945f3e45c489eb6360f6d9a": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 14)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "61bb00afa3a5f4228ee8c3d6624341158bc031d9c1a801ed4481f5fe095241fc": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 10)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "b150394af0a6a0d6239ada94af62ce4fe40a2c991e29cb122512f8cd0704f388": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 15)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "0ea704a4299bf930d45b13a1d1379dd25b8714848b3a8872dcad9198d89cf8d7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 17)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "394c16d01708d75ee622a98491a01e09224e817d61c63fb65b7a57e0ce387b2c": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 13)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "8612277e4d4bd632fd50301f8ca8c3fdd415beec3905572768e3fbf8f10ab98b": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 20)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "d195904172f21d23aa3966070551f352886f1d0d3942c8d8e6425aad6af8d08e": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 18)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "02e086eef437d961a2ae0ab4956aacfd7bbee9483b79bae8ff175da8a7e1759a": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 16)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "f9ec88032a4e09da597cceccbce33f9fbf947a44771c310cbcc677f9aa649dd2": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 23)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "e543c9d9632fa172bdb5bb283932e4232cd31ca1709cb1b47d2fd9e5a20537cd": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 21)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "fb10b35f572b9ad21af79cf3ea188dc3143501e29ff3dd707fa9fa2df4d1a52c": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 19)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "7563bc1009722bdd43ccf5846ac50df5e2d1e2e752816fddec00561fc8143dc7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 26)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "ff679871d37c1c6eaf2ac3496ceb92ff26fc7e98137f596be2ce454aeb143cb1": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 24)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "4d3d92f905b87db04348e187953e0924968f2d7f64a187536ecc39bd942f385b": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 22)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "b665dea5d6dcaf15c3543bf734d0dfcc8ad438a5f1c02525d236e910f6c488f6": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 29)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "7c755be3ee3ae6e8afa8f93afcbaf8b14f82fa16a464b8785038a389595c0572": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 27)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "c19257b66f04616c8319d1aa30020aff31269db84233a43e391adde9c2bc2564": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 25)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "4aa21ea63407d194592a2b9be7b51afbd5b6adb64d6312b51e4388e9a617f6d7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 32)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "2720e7d96aa99e5ecab04864014db6d60f6029cef294cb4bc406cb81ff8b1728": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 30)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "05d8fbfbef8bcbc73ae595328ccd41e461976b0a9fee1176c3210a9e8c784833": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 28)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "d59ac2bce4aa6b88a5f62340a496a6534a372bff18ed537ea08f2e4bfd44ad85": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 35)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "17eba2f5aafe710a6311640e36ea7a43d942d10c7ae13cad12be5a18c9b6422e": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 33)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "cfa99e206932c70fac7b4a256675c76fa6d94c260fe1ad610309c13459d9bef8": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 31)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "9ff73c975a93918f9ec4438bc33de12ec94d6d20c347f5d4a25d68ae50856939": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 38)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "22aca259e5dda9bc8b2cb3ce641dbab0b4dbf5db17ce08eaf9297b570b9f8227": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 36)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "02491c407f3733b9add505531e47f6ff92583ee2c722b78c0c539e50e10b4e28": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 34)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "a4c1592b3209d76b6d42111eb5082bb83b4ef57211cdcd312c165a91c34ddfb0": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 41)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "d88699ff4479e041ebe22e19249dc2bb630298865d02b558353ad9cb0e5088f3": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 39)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "2b96801251e63f96e4186b6e22bce9f066a31d4497e8bd2f24798ee20cbd2ea8": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 37)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "998041c2cab307dde851507b6c910c7adc432d5a79bec70bab2810fc0f39d083": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 44)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "921f25ef9cf42bd31f1c97f2e246bb48ff0795243509de0dfe22256a928ed43d": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 40)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "dff54ab3305dfaab814633b805b4fcbba025d6ebdd3c1df4330b9606df87597a": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 42)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "fe3672f37f07c25956e462afb0585df9dfc8a399c6be9189f87068ca74edaa4a": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 47)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "d9f1d22a50579a4f09566c42ed4d56a06be105c8da5db10e39f2911645010a57": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 45)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "a39df1ee979126d0736c8b58e60ba184fdbb96af815064dc4162448f533a96be": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 43)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "f5c3d8680c034af74350f02e5d58481b0e1f9d007ba31ac45aaa2c7f34518784": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 50)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "88eb3af4bdee8a7e8f2654d8734b74c4b2fc4d0734bb98d5510419068327b555": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 48)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "19b7ded5fb54746b8f735c470234c41d1a038541eb026980437329b282023fbd": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 46)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "3f44ebfc6a88fadecb809a0c3a086aa544bec6a6b3e2a16ea6ecddcfb457b876": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 53)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "857b334534247f0fdd065c2e2edc1777bee96f2d7ce227c4b41b847f2394738f": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 51)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "948e2c5dfca9aa15cf228bcd8f2a873a525b6147c4430e08abeedc2103ab74e2": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 49)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "1c42d25a06e66fc5dec813da632968208dd00723d0c846f541fe93f19a4de438": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 56)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "e668ae955d6fce15ae772967d7fcea44f913f085a5d96150f6b97f50274c6f4e": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 54)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "5b22e772a086a5dc357161a39a36daf36fe76407aaae2f84d055a288dce858b3": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 52)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "552e7f415a949eb1ac0476619da6ffd47183f79152b9af214232c5b97844c00e": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Count from 1 to 200. Write each number on its own line and nothing else. (request 59)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "42bdbddd34e965406ad3200e7fc87c0e1fb7fc435e94f986f16818c5700c5cc3": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 57)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "a978f17183b2e8cb1d2d805bd104f73e73067f1c62c1aeafc3ed3da691da32c6": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 55)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "a00465c2148cb0965aae033f90cb95c8a653d937ab6be5b99ee1d1e4a6cc81f7": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a Python class LRUCache with get(key) and put(key, value), both O(1), using a dict and a doubly linked list. Then write unittest test cases for it. (request 60)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "ab1ad3a00104824ad67f1bb44c2222225d6cc18aff6797e62853ecd9793469e4": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 58)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    },
+    "6179488294e705da6ad4955798cbd6281ec774b4c13137a515f01ca63661c142": {
+      "messages": [
+        {
+          "role": "user",
+          "content": "Write a 400-word essay about the invention of the printing press and how it changed Europe. (request 61)"
+        }
+      ],
+      "max_tokens": 256,
+      "temperature": 0,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      }
+    }
+  }
+}
diff --git a/docs/receipts/dsv41-neighbor-host/README.md b/docs/receipts/dsv41-neighbor-host/README.md
index f4b5882fe9..ee5b5b4a79 100644
--- a/docs/receipts/dsv41-neighbor-host/README.md
+++ b/docs/receipts/dsv41-neighbor-host/README.md
@@ -6,7 +6,8 @@ Local checks on 2026-10-20 UTC, macOS arm64, Zig 0.17.0. This change is based on
 
 The transport is opt-in through `TF_DS_NEIGHBOR_RDMA=1`; see the
 [configuration and dispatch contract](../../recipes/deepseek-v41-neighbor.md).
-No hardware or model qualification was performed for this branch.
+These were the initial host-only checks. Subsequent exact-commit device and model
+results are recorded in the [hardware receipt](../dsv41-neighbor-hardware/README.md).
 
 | Check | Result |
 | --- | --- |
@@ -29,10 +30,9 @@ The focused tests exercise host-side preconditions, startup-mode comparison,
 actual NCCL forwarding call sequencing with a recording API, and abort-word
 propagation. They do not validate device output, RDMA queue pairs, rail selection,
 CUDA graph execution, startup collectives on four machines, throughput, or model
-tokens. GPU/model gates in CONTRIBUTING.md remain pending, including byte/layout
-and graph checks, plain/drafted and concurrent/solo equality, prompt reuse,
-long-context behavior and matched performance. No fault, restart, or reboot gate
-was run.
+tokens. At this initial host-only stage, device/model gates remained pending. The linked
+hardware receipt records the subsequently tested subset and remaining limits.
+No fault, restart, or reboot gate was run.
 
 The earlier experimental implementation used source
 `b65c8b92d2403dfdefccf15a3b23ba609a4553e2`, a descendant of this pinned base, plus
diff --git a/docs/recipes/deepseek-v41-neighbor.md b/docs/recipes/deepseek-v41-neighbor.md
index 43ca22bf66..20757a59c2 100644
--- a/docs/recipes/deepseek-v41-neighbor.md
+++ b/docs/recipes/deepseek-v41-neighbor.md
@@ -72,3 +72,8 @@ changing-input graph replay, cutoff and unequal-width fallbacks, and mixed
 operation sequences. Then apply CONTRIBUTING.md's plain/drafted,
 concurrent/solo and resumed/fresh token gates and matched serving benchmarks.
 Prototype measurements on a different source revision do not qualify this branch.
+
+The [consolidated-branch hardware receipt](../receipts/dsv41-neighbor-hardware/README.md)
+records the tested code commit, ring setup, byte/graph/token checks and serving
+measurements. It does not claim a direct comparison with a switched fabric or the
+target branch.

From 9b0112eb482b368e78a45f424db7ee0faedc503d Mon Sep 17 00:00:00 2001
From: Scott Woolsey Biggart <297695+swbiggart@users.noreply.github.com>
Date: Tue, 20 Oct 2026 09:35:58 -0700
Subject: [PATCH 3/5] fix(dsv41): reconcile two v0.2.0 rebase conflicts

---
 README.md          |  8 ++++----
 zig/build/cuda.zig | 43 ++++++++++++++++++++++++++++++++++++++++---
 2 files changed, 44 insertions(+), 7 deletions(-)

diff --git a/README.md b/README.md
index 4a2766f303..d52b5bf4c1 100644
--- a/README.md
+++ b/README.md
@@ -3,10 +3,10 @@
 > (MIT), our EXL3 quant of [deepseek-ai/DeepSeek-V4.1-Flash-1017](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash-1017)
 > (MIT). Releases up to v0.1.1 served [Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw](https://huggingface.co/Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw)
 > (MIT), an EXL3 quant of [deepseek-ai/DeepSeek-V4.1-Flash](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash).
-> The branch adds a `dsv41` family to TensorFold 1.0.2's native engine: a four-node split whose replies equal the
-> two-node family's token for token in our gates, image input, up to 1,048,576 tokens a request in a shared
-> 2,097,152-position window with kept-prompt reuse, and exact DSpark speculative decoding. Launch, measured numbers,
-> gates and the kernel kit are in the recipe, [bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10](https://github.com/bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10).
+> The branch adds a `dsv41` family to TensorFold 1.0.12's native engine: a four-node split (on the 2.9 bpw weights its
+> replies equal the two-node family's token for token in our gates), image input, up to 1,048,576 tokens a request in
+> a shared 2,097,152-position window with kept-prompt reuse, and exact DSpark speculative decoding. Launch, measured
+> numbers, gates and the kernel kit are in the recipe, [bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10](https://github.com/bertholomus/deepseek-v4.1-tensorfold-tp4-4xgb10).
 > An optional [neighbor transport](docs/recipes/deepseek-v41-neighbor.md) supports adjacent-edge RDMA on a physical ring.
 > `tools/dsv41-zig/inputs/` makes each node's weight file and token map from the checkpoint; `tools/dsv41-zig/kit/`
 > makes and checks the kit's generated files. See NOTICE and ATTRIBUTION.md. This fork is not affiliated with or
diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig
index c50fa9bdde..885ab97d45 100644
--- a/zig/build/cuda.zig
+++ b/zig/build/cuda.zig
@@ -308,11 +308,48 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ .root_source_file = b.path("zig/build/check_fatbins.zig"), .target = host }) })).step);
     const cuda = runtime(b, host, .debug, &.{});
     const mods = family(b, host, .debug, cuda, draft_ids);
-    const eng = engines(b, host, .debug, cuda, mods.lanes, mods.nemotron, mods.core);
-    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, eng.api, eng.engines, eng.dsv41, stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
+    const native_modules = engines(b, host, .debug, cuda, mods.core, mods.lanes, mods.nemotron, mods.heat);
+    const native = native_modules.engines;
+    // the HTTP server's unit tests (zig/src/server/root.zig), which the macOS build runs, on Linux hosts too
+    const template = b.createModule(.{ .root_source_file = b.path("zig/src/core/template/template.zig"), .target = host, .optimize = .debug, .link_libc = true });
+    const server_tests = b.createModule(.{
+        .root_source_file = b.path("zig/src/server/root.zig"),
+        .target = host,
+        .optimize = .debug,
+        .link_libc = true,
+        .imports = &.{ .{ .name = "engine_api", .module = native_modules.api }, .{ .name = "tokenizer", .module = mods.tokenizer }, .{ .name = "template", .module = template } },
+    });
+    const server_json = @import("server.zig").addTo(b, server_tests, host, .debug);
+    const server_test = b.addRunArtifact(b.addTest(.{ .root_module = server_tests }));
+    const json_test = b.addRunArtifact(b.addTest(.{ .root_module = server_json }));
+    const server_step = b.step(
+        "test-server-cpu",
+        "The HTTP server's unit tests (routes, templates, tool parsing) without a GPU",
+    );
+    for ([_]*std.Build.Step.Run{ server_test, json_test }) |run| {
+        step.dependOn(&run.step);
+        server_step.dependOn(&run.step);
+    }
+    // the server goldens and Cortheon's fuzz step here when the target is not macOS, whose Metal build defines them
+    if (server_steps) {
+        golden_build.step(b, host, server_tests, native_modules.api, mods.tokenizer, template);
+        @import("cortheon_fuzz.zig").add(b, host, native_modules.api, mods.tokenizer, template);
+    }
+    // the tokenizer's own tests (memo, steps, regex, unicode) run here: no other test binary reaches them
+    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, lfm2Module(b, host, .debug, cuda, mods.core), mods.radio, mods.heat, mods.tokenizer, native_modules.api, native, native_modules.dsv41, stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
+    // the GLM-5.3-Flash CUDA family's host side
+    const glm = glmCuda(b, host, .debug, cuda, mods.core, mods.lanes, mods.heat);
+    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = glm })).step);
+    // the shares a GLM-5.3-Flash rank takes (cluster/canon.zig against the Python family's)
+    const shares = b.createModule(.{ .root_source_file = b.path("zig/src/cluster/canon_glm_test.zig"), .target = host, .optimize = .debug });
+    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = shares })).step);
+    // the tool-call parser's tests, run beside the server's on Macs: the GLM calls the CUDA port serves
+    const parse_tests = b.createModule(.{ .root_source_file = b.path("zig/src/server/tool_parse_test.zig"), .target = host, .optimize = .debug, .link_libc = true, .imports = &.{ .{ .name = "engine_api", .module = native_modules.api }, .{ .name = "tokenizer", .module = mods.tokenizer }, .{ .name = "template", .module = template } } });
+    _ = @import("server.zig").addTo(b, parse_tests, host, .debug);
+    step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = parse_tests })).step);
     const neighbor = b.createModule(.{ .root_source_file = b.path("zig/src/families/dsv41/neighbor.zig"), .target = host, .optimize = .debug, .link_libc = true });
     neighbor.addImport("cuda", cuda);
-    neighbor.addImport("fabric", eng.dsv41.import_table.get("fabric").?);
+    neighbor.addImport("fabric", native_modules.dsv41.import_table.get("fabric").?);
     b.step("test-dsv41-neighbor", "Host checks of TP4 neighbor transport (no GPU or verbs)").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = neighbor })).step);
     const cli = b.createModule(.{ .root_source_file = b.path("zig/src/cli/cuda_main.zig"), .target = host, .optimize = .debug, .link_libc = true });
     cli.addImport("cuda", cuda);

From d6953038565dd8b14f2acfcbdf281756b39bfe2f Mon Sep 17 00:00:00 2001
From: Scott Woolsey Biggart <297695+swbiggart@users.noreply.github.com>
Date: Tue, 20 Oct 2026 09:35:58 -0700
Subject: [PATCH 4/5] fix(dsv41): agree neighbor mode at startup and configure
 GID

---
 docs/recipes/deepseek-v41-neighbor.md | 30 ++++++++++----
 zig/src/families/dsv41/model.zig      | 11 +++--
 zig/src/families/dsv41/neighbor.zig   | 60 ++++++++++++++-------------
 zig/src/families/dsv41/prompt2d.zig   | 27 +++++++++++-
 4 files changed, 85 insertions(+), 43 deletions(-)

diff --git a/docs/recipes/deepseek-v41-neighbor.md b/docs/recipes/deepseek-v41-neighbor.md
index 20757a59c2..9e36fbd7be 100644
--- a/docs/recipes/deepseek-v41-neighbor.md
+++ b/docs/recipes/deepseek-v41-neighbor.md
@@ -18,16 +18,20 @@ For each rank, set both comma-separated device lists:
 
 Each list must contain two distinct names; the lists cannot overlap. List rails
 in the same physical order at both ends of each edge (for example, by PCI domain).
-The mode uses RoCE GID index 3 and the existing transport's MTU and queue settings.
+Set `TF_DS_NEIGHBOR_GID_INDEX` to the local RoCE GID table index (decimal 0–255).
+It defaults to 5, matching upstream's rings; sites using GID 3 must explicitly set
+`TF_DS_NEIGHBOR_GID_INDEX=3`. This is separate from NCCL's GID setting and is only
+read when neighbor mode is enabled. The existing MTU and queue settings are retained.
 It does not detect or configure networking. NCCL must also reach adjacent ranks
 and support its ordinary collectives on this fabric.
 
-A one-byte NCCL all-gather at TP4 model startup checks that every rank selected
-the same mode. This check also runs when the mode is off, so a mixed selection
-fails before neighbor setup. It creates no RDMA resources when disabled. Use the
-same binary on every rank; mixing versions with different startup collectives is
-unsupported. Device selection and mode stay fixed for the lifetime of the model
-and its captured graphs.
+The existing startup `Settings` exchange checks each follower’s neighbor mode
+against rank 0 before that follower enters NCCL communicator creation. This is
+the upstream per-follower handshake, not a global rejection broadcast. No separate mode all-gather
+runs, and disabled mode creates no neighbor RDMA resources. Use the same binary
+on every rank: the settings frame now includes the neighbor flag. Device and GID
+selection stay local to each rank; they must describe compatible physical links.
+All settings stay fixed for the lifetime of the model and its captured graphs.
 
 ## Operations
 
@@ -60,10 +64,12 @@ zig build test test-golden -Dcpu=apple_m1
 python3 tools/zig/lean_check.py
 ```
 
-The focused host target checks mode disagreement, bounded dispatch, device-list
+The focused host target checks default-off mode and GID parsing, bounded dispatch, device-list
 validation, relay preconditions, the production NCCL forwarding calls at every
 rank with unequal counts, and local abort propagation. It also compiles the
-GPU/verbs declarations. It does not execute GPU kernels, RDMA traffic, captured
+GPU/verbs declarations. The family host suite checks startup mode disagreement,
+including both directions of mixed enablement over the existing settings link.
+The host tests do not execute GPU kernels, RDMA traffic, captured
 graphs, or a model.
 
 Before hardware acceptance, compare all output bytes with an independent NCCL
@@ -77,3 +83,9 @@ The [consolidated-branch hardware receipt](../receipts/dsv41-neighbor-hardware/R
 records the tested code commit, ring setup, byte/graph/token checks and serving
 measurements. It does not claim a direct comparison with a switched fabric or the
 target branch.
+
+The [v0.2.0 update receipt](../receipts/dsv41-neighbor-v020/README.md) records
+the rebased code checks separately from the historical hardware measurements.
+
+The [v0.2.0 hardware repeat](../receipts/dsv41-neighbor-v020-hardware/README.md)
+records fresh byte/graph/token checks and the same serving workloads on the rebased code.
diff --git a/zig/src/families/dsv41/model.zig b/zig/src/families/dsv41/model.zig
index 210309f8b0..6442e43d81 100644
--- a/zig/src/families/dsv41/model.zig
+++ b/zig/src/families/dsv41/model.zig
@@ -232,6 +232,11 @@ pub const Model = struct {
         errdefer m.w.deinit();
         if (o.drafts and m.w.dspark == null) return error.NoDrafter;
 
+        const neighbor_rdma = try neighbor.configured();
+        if (neighbor_rdma and o.world != 4) return error.NotFourNodes;
+        if (neighbor_rdma and o.rdma_devices != null) return error.NeighborDisallowsStockRdma;
+        const neighbor_gid = if (neighbor_rdma) try neighbor.gidIndex() else 0;
+
         // the ranks' link and NCCL (rank 0 listens: rank 1 retries its connect while rank 0 loads)
         m.fan = try prompt2d.Fan.open(io, o.master, o.port, o.rank, o.world);
         errdefer m.fan.close();
@@ -244,12 +249,10 @@ pub const Model = struct {
             .drafts = @intFromBool(o.drafts),
             .replay_floor = @intFromBool(prompt.replay_floor),
             .idx_split = @intCast(prompt.split_from),
+            .neighbor_rdma = @intFromBool(neighbor_rdma),
             .pool = o.pool,
         });
         errdefer m.comm.deinit();
-        const neighbor_rdma = try neighbor.configured();
-        try neighbor.agree(&m.comm, m.stream, neighbor_rdma);
-        if (neighbor_rdma and o.rdma_devices != null) return error.NeighborDisallowsStockRdma;
         // TP2: the decode gathers' RDMA ring (both ranks connect before either returns)
         if (sp.pair == null and o.rdma_devices != null) {
             const list = o.rdma_devices.?;
@@ -309,7 +312,7 @@ pub const Model = struct {
             m.two = try prompt2d.Two.init(c, &m.w, &m.arena, sp, chunk_rows);
             if (neighbor_rdma) {
                 m.rdma_mod = try cuda.Module.load(d, cuda.kernels.dsv41_rdma);
-                m.neighbor_gather = try neighbor.Gather.open(gpa, d, try rdma.Kernels.load(m.rdma_mod.?), o.rank, &m.comm, m.stream);
+                m.neighbor_gather = try neighbor.Gather.open(gpa, d, try rdma.Kernels.load(m.rdma_mod.?), o.rank, &m.comm, m.stream, neighbor_gid);
                 m.two.?.neighbor_gather = &m.neighbor_gather.?;
             }
             if (o.rdma_devices) |devs| {
diff --git a/zig/src/families/dsv41/neighbor.zig b/zig/src/families/dsv41/neighbor.zig
index 33d5bc92b2..d4747fd8be 100644
--- a/zig/src/families/dsv41/neighbor.zig
+++ b/zig/src/families/dsv41/neighbor.zig
@@ -5,30 +5,26 @@ const rdma = @import("rdma.zig");
 const Comm = @import("comm.zig").Comm;
 pub const cutoff: usize = 512 * 1024;
 pub fn configured() !bool {
-    const raw = std.c.getenv("TF_DS_NEIGHBOR_RDMA") orelse return false;
-    const value = std.mem.span(raw);
+    const raw = std.c.getenv("TF_DS_NEIGHBOR_RDMA");
+    return parseMode(if (raw) |v| std.mem.span(v) else null);
+}
+
+fn parseMode(raw: ?[]const u8) !bool {
+    const value = raw orelse return false;
     if (std.mem.eql(u8, value, "1")) return true;
     if (std.mem.eql(u8, value, "0")) return false;
     return error.BadNeighborMode;
 }
-/// Every TP4 rank agrees before any neighbor queue pair or captured graph is created.
-pub fn agree(comm: *const Comm, stream: cuda.Stream, enabled: bool) !void {
-    if (comm.world != 4) {
-        if (enabled) return error.NotFourNodes;
-        return;
-    }
-    var dev = try cuda.DeviceBuffer.alloc(stream.d, 5);
-    defer dev.free();
-    try dev.upload(0, &.{@intFromBool(enabled)});
-    try comm.allGather(dev.ptr, dev.ptr + 1, 1, .u8, stream);
-    try stream.synchronize();
-    var modes: [4]u8 = undefined;
-    try dev.download(1, &modes);
-    try sameMode(modes);
+pub fn gidIndex() !u8 {
+    const raw = std.c.getenv("TF_DS_NEIGHBOR_GID_INDEX");
+    return parseGid(if (raw) |v| std.mem.span(v) else null);
 }
 
-fn sameMode(modes: [4]u8) !void {
-    for (modes) |mode| if (mode != modes[0]) return error.NeighborModeMismatch;
+fn parseGid(value: ?[]const u8) !u8 {
+    const raw = value orelse return 5;
+    if (raw.len == 0) return error.BadNeighborGid;
+    for (raw) |c| if (c < '0' or c > '9') return error.BadNeighborGid;
+    return std.fmt.parseInt(u8, raw, 10) catch error.BadNeighborGid;
 }
 
 pub fn eligible(bytes: usize) bool {
@@ -40,7 +36,7 @@ pub const Gather = struct {
     scratch: cuda.DeviceBuffer,
 
     /// Each stage uses two rails in matching order at both ends of its physical edge.
-    pub fn open(gpa: std.mem.Allocator, d: *const cuda.Driver, k: rdma.Kernels, g: u32, comm: *const Comm, stream: cuda.Stream) !Gather {
+    pub fn open(gpa: std.mem.Allocator, d: *const cuda.Driver, k: rdma.Kernels, g: u32, comm: *const Comm, stream: cuda.Stream, gid_index: u8) !Gather {
         if (g >= 4 or comm.world != 4) return error.NotFourNodes;
         const first_raw = std.c.getenv("TF_DS_NEIGHBOR_FIRST") orelse return error.NoFirstDevices;
         const second_raw = std.c.getenv("TF_DS_NEIGHBOR_SECOND") orelse return error.NoSecondDevices;
@@ -49,7 +45,7 @@ pub const Gather = struct {
         for (first_devices) |x| for (second_devices) |y| {
             if (std.mem.eql(u8, x, y)) return error.NeighborDevicesOverlap;
         };
-        const s: rdma.Settings = .{ .max_bytes = 2 * cutoff, .gid_index = 3 };
+        const s: rdma.Settings = .{ .max_bytes = 2 * cutoff, .gid_index = gid_index };
         const first = try rdma.Ring.create(gpa, d, k, &first_devices, g % 2, 2, s);
         errdefer first.destroy(gpa);
         const second = try rdma.Ring.create(gpa, d, k, &second_devices, g / 2, 2, s);
@@ -194,14 +190,22 @@ test "neighbor stages produce global rank order and bounded dispatch" {
     try std.testing.expect(!eligible(0));
 }
 
-test "neighbor modes agree before rings open" {
-    try sameMode(.{ 0, 0, 0, 0 });
-    try sameMode(.{ 1, 1, 1, 1 });
-    for (0..4) |g| {
-        var modes: [4]u8 = @splat(0);
-        modes[g] = 1;
-        try std.testing.expectError(error.NeighborModeMismatch, sameMode(modes));
-    }
+test "neighbor mode is off unless explicitly enabled" {
+    try std.testing.expect(!try parseMode(null));
+    try std.testing.expect(!try parseMode("0"));
+    try std.testing.expect(try parseMode("1"));
+    for ([_][]const u8{ "", "2", "true", " 1" }) |bad|
+        try std.testing.expectError(error.BadNeighborMode, parseMode(bad));
+}
+
+test "neighbor GID index accepts site settings and rejects invalid values" {
+    try std.testing.expectEqual(@as(u8, 5), try parseGid(null));
+    try std.testing.expectEqual(@as(u8, 3), try parseGid("3"));
+    try std.testing.expectEqual(@as(u8, 5), try parseGid("5"));
+    try std.testing.expectEqual(@as(u8, 0), try parseGid("0"));
+    try std.testing.expectEqual(@as(u8, 255), try parseGid("255"));
+    for ([_][]const u8{ "", "-1", "+3", "256", "0x3", " 3", "three" }) |bad|
+        try std.testing.expectError(error.BadNeighborGid, parseGid(bad));
 }
 
 test "a stage requires two distinct device names" {
diff --git a/zig/src/families/dsv41/prompt2d.zig b/zig/src/families/dsv41/prompt2d.zig
index 32ba6474d0..6ecef491f3 100644
--- a/zig/src/families/dsv41/prompt2d.zig
+++ b/zig/src/families/dsv41/prompt2d.zig
@@ -144,6 +144,8 @@ pub const Settings = extern struct {
     drafts: u32, // the drafter (--drafts): the same
     replay_floor: u32, // TF_DS_REPLAY_FLOOR: the same
     idx_split: u32 = 0, // TF_DS_IDX_SPLIT (prompt.split_from): the same (the keyed indexer layers all-gather inside)
+    neighbor_rdma: u32 = 0, // TF_DS_NEIGHBOR_RDMA: the same exchange path on every rank
+    reserved: u32 = 0,
     pool: u64, // rank 0's --context (a follower's --pool): a follower's holds rank 0's
 
     pub const tag: u32 = 0x31345344; // "DS41"
@@ -168,6 +170,8 @@ pub const Settings = extern struct {
             w.print("{s}rank {d} has TF_DS_REPLAY_FLOOR={d} and rank 0 {d}: set the same on every rank", .{ sepIf(&w, sep), rank, mine.replay_floor, lead.replay_floor }) catch {};
         if (mine.idx_split != lead.idx_split)
             w.print("{s}rank {d} splits the keyed indexer rows from {d} columns and rank 0 from {d} (TF_DS_IDX_SPLIT, 0: unset): set the same on every rank", .{ sepIf(&w, sep), rank, mine.idx_split, lead.idx_split }) catch {};
+        if (mine.neighbor_rdma != lead.neighbor_rdma)
+            w.print("{s}rank {d} has TF_DS_NEIGHBOR_RDMA={d} and rank 0 {d}: set the same on every rank", .{ sepIf(&w, sep), rank, mine.neighbor_rdma, lead.neighbor_rdma }) catch {};
         return w.buffered();
     }
 
@@ -487,7 +491,7 @@ test "a follower splits the keyed indexer rows as rank 0 does, or says to set th
     lead.idx_split = 32768;
     try std.testing.expectEqualStrings("", lead.differ(lead, 2, &buf));
     try std.testing.expectEqualStrings("rank 2 splits the keyed indexer rows from 0 columns and rank 0 from 32768 (TF_DS_IDX_SPLIT, 0: unset): set the same on every rank", served.differ(lead, 2, &buf));
-    try std.testing.expectEqual(@as(usize, 40), @sizeOf(Settings)); // rank 0's frame: as before (idx_split took the pad)
+    try std.testing.expectEqual(@as(usize, 48), @sizeOf(Settings));
 }
 
 test "rank 0's settings go with NCCL's id: the follower answers, and an older rank 0's id alone is taken as before" {
@@ -495,10 +499,15 @@ test "rank 0's settings go with NCCL's id: the follower answers, and an older ra
     const served: Settings = .{ .world = 2, .streams = 8, .round_rows = 16, .rails = 0, .drafts = 1, .replay_floor = 1, .pool = 1 << 19 };
     var four = served;
     four.streams = 4;
+    var neighbor_on = served;
+    neighbor_on.neighbor_rdma = 1;
     const Case = struct { lead: ?Settings, follower: ?Settings, differs: bool };
     const cases = [_]Case{
         .{ .lead = served, .follower = served, .differs = false },
         .{ .lead = served, .follower = four, .differs = true },
+        .{ .lead = served, .follower = neighbor_on, .differs = true },
+        .{ .lead = neighbor_on, .follower = served, .differs = true },
+        .{ .lead = neighbor_on, .follower = neighbor_on, .differs = false },
         .{ .lead = null, .follower = four, .differs = false }, // an older rank 0: the id alone, nothing to check
     };
     const base: u16 = 39500 + @as(u16, @intCast(@mod(std.Io.Timestamp.now(io, .awake).nanoseconds, 400)));
@@ -539,6 +548,20 @@ test "rank 0's settings go with NCCL's id: the follower answers, and an older ra
         try std.testing.expectEqualSlices(u8, &id, &peer.id);
         try std.testing.expectEqualStrings(peer.said, try heard);
         try std.testing.expectEqual(cs.differs, peer.said.len > 0);
-        if (cs.differs) try std.testing.expect(std.mem.indexOf(u8, peer.said, "--streams 8") != null);
+        if (cs.differs) {
+            const setting = if (cs.follower.?.streams != cs.lead.?.streams) "--streams 8" else "TF_DS_NEIGHBOR_RDMA";
+            try std.testing.expect(std.mem.indexOf(u8, peer.said, setting) != null);
+        }
     }
 }
+
+test "startup settings reject mixed neighbor modes before communicator creation" {
+    const off: Settings = .{ .world = 4, .streams = 16, .round_rows = 48, .rails = 0, .drafts = 1, .replay_floor = 0, .pool = 1 << 21 };
+    var on = off;
+    on.neighbor_rdma = 1;
+    var buf: [1024]u8 = undefined;
+    try std.testing.expectEqualStrings("", off.differ(off, 1, &buf));
+    try std.testing.expectEqualStrings("", on.differ(on, 1, &buf));
+    try std.testing.expectEqualStrings("rank 1 has TF_DS_NEIGHBOR_RDMA=0 and rank 0 1: set the same on every rank", off.differ(on, 1, &buf));
+    try std.testing.expectEqualStrings("rank 1 has TF_DS_NEIGHBOR_RDMA=1 and rank 0 0: set the same on every rank", on.differ(off, 1, &buf));
+}

From d08028564559ebfad465e5e566b6573b601cf3bd Mon Sep 17 00:00:00 2001
From: Scott Woolsey Biggart <297695+swbiggart@users.noreply.github.com>
Date: Tue, 20 Oct 2026 09:35:58 -0700
Subject: [PATCH 5/5] docs(dsv41): update host and four-node qualification
 receipts

---
 .../dsv41-neighbor-hardware/README.md         |  38 +--
 .../dsv41-neighbor-hardware/summary.json      |  82 ++++---
 .../dsv41-neighbor-hardware/transport.json    | 224 +++++++++---------
 docs/receipts/dsv41-neighbor-host/README.md   |  25 ++
 docs/recipes/deepseek-v41-neighbor.md         |   8 +-
 5 files changed, 217 insertions(+), 160 deletions(-)

diff --git a/docs/receipts/dsv41-neighbor-hardware/README.md b/docs/receipts/dsv41-neighbor-hardware/README.md
index 669f129e6e..a521a9d1f5 100644
--- a/docs/receipts/dsv41-neighbor-hardware/README.md
+++ b/docs/receipts/dsv41-neighbor-hardware/README.md
@@ -1,8 +1,8 @@
 # TP4 neighbor transport: four-host hardware qualification
 
 Tested 2026-10-20 UTC at code commit
-`a9225d72ec63b19867c72eba1cca48a49648805c`, based on
-`deepseek-v41-zig-tp4` at `59402086381a42c74d53053443224eec2790b4c5`.
+`660bf9aba19b7ac42ea7ca1b4e26c50b4850bce1`, based on
+`deepseek-v41-zig-tp4` at `b65c8b92d2403dfdefccf15a3b23ba609a4553e2`.
 Subsequent receipt-only changes do not change the tested engine code.
 
 We do not have a network switch. Four GB10 hosts form a physical ring, with two
@@ -26,7 +26,7 @@ zig build native neighbor-lanes neighbor-probe \
 `neighbor-lanes` and `neighbor-probe` were external qualification build targets,
 not release targets added to this branch. The latter exposes a C ABI around the
 production code for a weight-free Python/CUDA driver. Only build target additions
-and the new adapter were applied to the source archive; all 423 tracked files
+and the new adapter were applied to the source archive; all 1,200 tracked files
 under `zig/src` were verified byte-identical to the tested commit. The native
 executables have no test-path substitutions. Binary digests and per-rank OS,
 driver and kernel versions are in [summary.json](summary.json).
@@ -36,7 +36,8 @@ Model: DeepSeek-V4.1-Flash-1017 EXL3 revision
 `1bae0e1088993b9b82de08f59de5be5a34e42170`, DSpark.
 Effective server arguments: `--context 1048576 --parallel 16 --max-tokens 32768`,
 shared window 2097152, round rows 48, graphs enabled, served-depth policy `solo`.
-`TF_DS_NEIGHBOR_RDMA=1` on all four ranks; no per-optimization switches or stock
+`TF_DS_NEIGHBOR_RDMA=1` and `TF_DS_NEIGHBOR_GID_INDEX=3` on all four ranks
+(the GID default is 5); no per-optimization switches or stock
 `TF_RDMA_DEVICES`. NCCL 2.30.7 with the existing switchless-routing patches,
 IB network, Ring algorithm, four channels, and LL128 excluded. No forced-tool,
 model arithmetic, or server patches were added to the tested commit.
@@ -49,7 +50,8 @@ model arithmetic, or server patches were added to the tested commit.
   gather helper used by Markov. Sizes include the 512 KiB cutoff and alignment
   fallbacks. Four-rank startup mode agreement also executed successfully.
 - **9 short model replies:** three prompts through serial, drafted and concurrent
-  paths, 64 generated tokens each; exact token IDs and UTF-8 agree.
+  paths, 64 generated tokens each; exact token IDs and UTF-8 agree. All nine
+  also match the original commit’s replies for identical requests.
 - **24 benchmark replies:** four requests × six timed bursts, 256 tokens each;
   UTF-8 agrees across repetitions for each request.
 - **11 longer-prompt replies:** four approximately 8K-token prompts (requested
@@ -67,10 +69,10 @@ rank's mean of two measurement rounds. They are not a target-branch comparison.
 
 | Diagonal payload | Latency |
 |---|---:|
-| 8.75 KiB | 17.81 µs |
-| 52.5 KiB | 27.77 µs |
-| 140 KiB | 42.89 µs |
-| 420 KiB | 87.10 µs |
+| 8.75 KiB | 18.34 µs |
+| 52.5 KiB | 27.91 µs |
+| 140 KiB | 44.79 µs |
+| 420 KiB | 82.12 µs |
 
 [transport.json](transport.json) contains every candidate rank/shape result.
 
@@ -85,18 +87,22 @@ comparable, and these short measurements are not confidence estimates.
 
 | Workload | Aggregate tokens/s | Detail |
 |---|---:|---|
-| Set-a four-stream bursts | **132.63** | Mean of six; range 131.94–133.31 |
-| Set-a steady admission | **148.83** | 30 seconds, 19 requests |
-| Set-b sustained | **169.77** | 90 seconds, 62 requests; TTFT p50 0.159 s, p95 0.283 s |
+| Set-a four-stream bursts | **133.11** | Mean of six; range 132.49–133.52 |
+| Set-a steady admission | **149.00** | 30 seconds, 19 requests |
+| Set-b sustained | **170.20** | 90 seconds, 62 requests; TTFT p50 0.166 s, p95 0.262 s |
+
+The suite repeats the original commit’s workloads on the same ring. Historical
+burst/steady/sustained rates were 132.63/148.83/169.77 tok/s; these were not
+alternating A/B runs, and the sub-percent differences do not establish a speedup.
 
 ## Host checks and limits
 
-A fresh `zig build test-dsv41-neighbor --cache-dir /tmp/neighbor-host-cache --summary all`
-passed **14/14**. `zig build test test-golden -Dcpu=apple_m1 --summary all` again
+A fresh `zig build test-dsv41-neighbor --summary all`
+passed **15/15**. `zig build test test-golden -Dcpu=apple_m1 --summary all` again
 reported the inherited macOS `std.c.SO.BUSY_POLL` compile blocker; the frozen
-corpus remained 312/319 equal, seven known differences, zero unexpected. The full
+corpus remained 309/319 equal, ten known differences, zero unexpected. The full
 host command is not reported as a pass. See the [host receipt](../dsv41-neighbor-host/README.md)
-for the pinned-base comparison and inherited allocator/lean findings.
+for the Linux node build, local CI checks and inherited failures.
 
 Full-1M quality, 160K cold prefill, 16-way sustained load, fault injection,
 restart/reboot and long-duration reliability were not tested on this consolidated
diff --git a/docs/receipts/dsv41-neighbor-hardware/summary.json b/docs/receipts/dsv41-neighbor-hardware/summary.json
index 1c5a1bda40..c5955bdc63 100644
--- a/docs/receipts/dsv41-neighbor-hardware/summary.json
+++ b/docs/receipts/dsv41-neighbor-hardware/summary.json
@@ -1,8 +1,9 @@
 {
-  "tested_commit": "a9225d72ec63b19867c72eba1cca48a49648805c",
-  "target_commit": "59402086381a42c74d53053443224eec2790b4c5",
+  "tested_commit": "660bf9aba19b7ac42ea7ca1b4e26c50b4850bce1",
+  "target_commit": "b65c8b92d2403dfdefccf15a3b23ba609a4553e2",
   "direct_target_comparison": false,
   "reason": "No network switch is available; these are measurements of the fastest implementation selected in our ring experiments, not a direct target-branch or switched-fabric comparison.",
+  "comparison_scope": "Repeated original candidate workloads on the same fleet; historical comparison, not alternating A/B.",
   "hardware": {
     "gpu": "NVIDIA GB10",
     "architecture": "aarch64",
@@ -40,30 +41,39 @@
     "nccl_sha256": "1cdf5f3e2d617d5c07a87903eb8849d5bd7758f6c900455d613a2e0eb289d63e"
   },
   "binaries": {
-    "tensorfold-native": "ff0a33c176e621115a4f98d662f56deed7ba54c1e005ac5b8742cf17a6279cae",
-    "libtf-neighbor.so": "38867278afe5e9560d64e81b1f1ce6bf1a953713a5243057bb0e4ff14ff8e00a",
+    "tensorfold-native": "6358c0a5c84923cebe3220ce2eb00d2210c93272921005d98b7eb82cbe5a239b",
+    "libtf-neighbor.so": "82897cfc01cf3aede02077a0ab7cc142f38828fc5ae11e4e762e9c6252d2347f",
     "probe.py": "8d1900817f5e8b240c4be327fef72ca2f5bd187701c109f64007cc34b9d4dba2",
-    "tf-dsv41-lanes": "fb85c5476b7ef946bf91de6fd3ef9380ab0f04d2e809abeaf050c7915dfccc69"
+    "tf-dsv41-lanes": "83084a74f05fa50a7680b433194f37eec999ab158dbf4446d50c18f476264727",
+    "transport_base.py": "10a908eeaca9959b501e1652c095ba0d4334d74a58861aac2ee65007d6acf1c7",
+    "direct.py": "df7e5c052010da1f3f63e54c1c4f33f6853ed282d20dea2fa8e96c69473d3a80"
   },
+  "production_source": {
+    "commit": "660bf9aba19b7ac42ea7ca1b4e26c50b4850bce1",
+    "files": 1200,
+    "all_equal": true
+  },
+  "gid_index": 3,
+  "default_gid_index": 5,
   "transport": {
     "rank_shape_cases": 56,
     "pass_bytes_and_graphs": true,
     "latencies": [
       {
         "bytes": 8960,
-        "slowest_rank_mean_us": 17.814800143241882
+        "slowest_rank_mean_us": 18.343599885702133
       },
       {
         "bytes": 53760,
-        "slowest_rank_mean_us": 27.772000432014465
+        "slowest_rank_mean_us": 27.905799448490143
       },
       {
         "bytes": 143360,
-        "slowest_rank_mean_us": 42.89499968290329
+        "slowest_rank_mean_us": 44.79059875011444
       },
       {
         "bytes": 430080,
-        "slowest_rank_mean_us": 87.09779977798462
+        "slowest_rank_mean_us": 82.11579918861389
       }
     ]
   },
@@ -72,7 +82,8 @@
     "paths": 3,
     "tokens_per_reply": 64,
     "replies": 9,
-    "pass_test": true
+    "pass_test": true,
+    "also_equal_to_original_commit": true
   },
   "long_context": {
     "burst_equal_to_solo": 5,
@@ -81,8 +92,8 @@
     "needles_found_burst": 4,
     "needles_found_solo": 4,
     "prompts": 4,
-    "burst_s": 17.4,
-    "solo_s": 7.7,
+    "burst_s": 17.5,
+    "solo_s": 7.8,
     "serial_s": 1.8,
     "cached_tokens_solo": [
       6144,
@@ -98,9 +109,9 @@
     "tokens": 256,
     "warmups": 1,
     "repetitions": 6,
-    "mean_aggregate_tps": 132.6345814307147,
-    "min_aggregate_tps": 131.9383995892958,
-    "max_aggregate_tps": 133.30848622707344,
+    "mean_aggregate_tps": 133.10816911985555,
+    "min_aggregate_tps": 132.48854455631133,
+    "max_aggregate_tps": 133.52409085662865,
     "exact_utf8_replies": 24
   },
   "steady": {
@@ -109,12 +120,12 @@
     "max_tokens": 256,
     "set": "a",
     "requests": 19,
-    "aggregate_tps": 148.83,
-    "per_stream_tps_p50": 47.475,
-    "per_stream_tps_p95": 62.115,
-    "per_stream_tps_p05": 22.925,
-    "ttft_s_p50": 0.189,
-    "ttft_s_p95": 0.49
+    "aggregate_tps": 149.0,
+    "per_stream_tps_p50": 48.066,
+    "per_stream_tps_p95": 61.506,
+    "per_stream_tps_p05": 23.011,
+    "ttft_s_p50": 0.196,
+    "ttft_s_p95": 0.429
   },
   "sustained": {
     "streams": 4,
@@ -122,11 +133,28 @@
     "max_tokens": 256,
     "set": "b",
     "requests": 62,
-    "aggregate_tps": 169.77,
-    "per_stream_tps_p50": 51.476,
-    "per_stream_tps_p95": 82.968,
-    "per_stream_tps_p05": 25.957,
-    "ttft_s_p50": 0.159,
-    "ttft_s_p95": 0.283
+    "aggregate_tps": 170.2,
+    "per_stream_tps_p50": 51.742,
+    "per_stream_tps_p95": 83.059,
+    "per_stream_tps_p05": 25.56,
+    "ttft_s_p50": 0.166,
+    "ttft_s_p95": 0.262
+  },
+  "original_commit_comparison": {
+    "burst": {
+      "original": 132.6345814307147,
+      "current": 133.10816911985555,
+      "change_percent": 0.3570619999945013
+    },
+    "steady": {
+      "original": 148.83,
+      "current": 149.0,
+      "change_percent": 0.11422428273868412
+    },
+    "sustained": {
+      "original": 169.77,
+      "current": 170.2,
+      "change_percent": 0.2532838546268401
+    }
   }
 }
diff --git a/docs/receipts/dsv41-neighbor-hardware/transport.json b/docs/receipts/dsv41-neighbor-hardware/transport.json
index 4c8bad74fa..05825ef086 100644
--- a/docs/receipts/dsv41-neighbor-hardware/transport.json
+++ b/docs/receipts/dsv41-neighbor-hardware/transport.json
@@ -10,8 +10,8 @@
         "width": 409600,
         "exact": true,
         "candidate_us": [
-          30.493998527526855,
-          30.445200204849243
+          30.510801076889038,
+          30.60159981250763
         ]
       },
       {
@@ -20,8 +20,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.60559749603271,
-          202.56800651550293
+          208.8956117630005,
+          202.35960483551025
         ]
       },
       {
@@ -30,8 +30,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.526399970054626,
-          14.788000285625458
+          14.606000483036041,
+          14.523200690746307
         ]
       },
       {
@@ -40,8 +40,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          18.897199630737305,
-          18.829600512981415
+          18.849200010299683,
+          18.719999492168427
         ]
       },
       {
@@ -50,8 +50,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          128.52879762649536,
-          122.29399681091309
+          121.89359664916992,
+          121.00640535354614
         ]
       },
       {
@@ -60,8 +60,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          226.56400203704834,
-          198.59280586242676
+          201.19280815124512,
+          206.52880668640137
         ]
       },
       {
@@ -70,8 +70,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          17.79559999704361,
-          17.834000289440155
+          18.53639930486679,
+          18.086400628089905
         ]
       },
       {
@@ -80,8 +80,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          27.924799919128418,
-          27.32599973678589
+          27.906399965286255,
+          27.90519893169403
         ]
       },
       {
@@ -90,8 +90,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          46.43999934196472,
-          39.35000002384186
+          47.65839874744415,
+          41.92279875278473
         ]
       },
       {
@@ -100,8 +100,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          92.94919967651367,
-          80.39720058441162
+          83.07600021362305,
+          80.69239854812622
         ]
       },
       {
@@ -110,8 +110,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.654800295829773,
-          14.564000070095062
+          14.602400362491608,
+          14.568400382995605
         ]
       },
       {
@@ -120,8 +120,8 @@
         "width": 524288,
         "exact": true,
         "candidate_us": [
-          94.62159872055054,
-          94.88400220870972
+          95.41760087013245,
+          97.06680178642273
         ]
       },
       {
@@ -130,8 +130,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.54680347442627,
-          216.4423942565918
+          200.18160343170166,
+          191.20919704437256
         ]
       },
       {
@@ -140,8 +140,8 @@
         "width": 18,
         "exact": true,
         "candidate_us": [
-          109.96840000152588,
-          109.81719493865967
+          111.08880043029785,
+          81.72240257263184
         ]
       }
     ]
@@ -157,8 +157,8 @@
         "width": 409600,
         "exact": true,
         "candidate_us": [
-          30.48959970474243,
-          30.421200394630432
+          30.526000261306763,
+          30.595600605010986
         ]
       },
       {
@@ -167,8 +167,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.48959493637085,
-          202.37560272216797
+          208.95318984985352,
+          202.19480991363525
         ]
       },
       {
@@ -177,8 +177,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.542800188064575,
-          14.771999418735504
+          14.607599377632141,
+          14.505599439144135
         ]
       },
       {
@@ -187,8 +187,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          18.901599943637848,
-          18.756400048732758
+          18.774400651454926,
+          18.72600018978119
         ]
       },
       {
@@ -197,8 +197,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          128.51239442825317,
-          121.68760299682617
+          122.24199771881104,
+          121.17400169372559
         ]
       },
       {
@@ -207,8 +207,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          226.52881145477295,
-          198.57800006866455
+          201.1728048324585,
+          206.53119087219238
         ]
       },
       {
@@ -217,8 +217,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          17.684000730514526,
-          17.757999897003174
+          18.56359988451004,
+          18.093200027942657
         ]
       },
       {
@@ -227,8 +227,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          28.045600652694702,
-          27.455198764801025
+          27.94800102710724,
+          27.698001265525818
         ]
       },
       {
@@ -237,8 +237,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          46.21039927005768,
-          39.081600308418274
+          47.57080078125,
+          41.88840091228485
         ]
       },
       {
@@ -247,8 +247,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          94.0060019493103,
-          80.18959760665894
+          83.43240022659302,
+          80.7099997997284
         ]
       },
       {
@@ -257,8 +257,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.591200649738312,
-          14.52919989824295
+          14.508000016212463,
+          14.536799490451813
         ]
       },
       {
@@ -267,8 +267,8 @@
         "width": 524288,
         "exact": true,
         "candidate_us": [
-          94.61640119552612,
-          95.49999833106995
+          95.46639919281006,
+          96.30399942398071
         ]
       },
       {
@@ -277,8 +277,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.21360158920288,
-          217.42920875549316
+          200.1267910003662,
+          191.1736011505127
         ]
       },
       {
@@ -287,8 +287,8 @@
         "width": 18,
         "exact": true,
         "candidate_us": [
-          110.02440452575684,
-          109.80960130691528
+          111.23640537261963,
+          81.77880048751831
         ]
       }
     ]
@@ -304,8 +304,8 @@
         "width": 409600,
         "exact": true,
         "candidate_us": [
-          30.64880073070526,
-          30.66360056400299
+          30.682000517845154,
+          30.62080144882202
         ]
       },
       {
@@ -314,8 +314,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.56400537490845,
-          202.4616003036499
+          208.96799564361572,
+          202.25400924682617
         ]
       },
       {
@@ -324,8 +324,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.658400416374207,
-          14.724799990653992
+          14.661599695682526,
+          14.50120061635971
         ]
       },
       {
@@ -334,8 +334,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          18.907199800014496,
-          18.756400048732758
+          18.75,
+          18.72040033340454
         ]
       },
       {
@@ -344,8 +344,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          128.5156011581421,
-          121.51639461517334
+          122.23520278930664,
+          121.12840414047241
         ]
       },
       {
@@ -354,8 +354,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          226.47359371185303,
-          198.57200384140015
+          201.13320350646973,
+          206.45320415496826
         ]
       },
       {
@@ -364,8 +364,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          17.76439994573593,
-          17.847999930381775
+          17.870399355888367,
+          18.094399571418762
         ]
       },
       {
@@ -374,8 +374,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          28.082001209259033,
-          27.461999654769897
+          27.934399247169495,
+          27.528798580169678
         ]
       },
       {
@@ -384,8 +384,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          46.42600119113922,
-          39.333200454711914
+          47.58960008621216,
+          41.90320074558258
         ]
       },
       {
@@ -394,8 +394,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          94.12440061569214,
-          79.66960072517395
+          83.41760039329529,
+          80.70080280303955
         ]
       },
       {
@@ -404,8 +404,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.567600190639496,
-          14.519999921321869
+          14.497199654579163,
+          14.518000185489655
         ]
       },
       {
@@ -414,8 +414,8 @@
         "width": 524288,
         "exact": true,
         "candidate_us": [
-          95.23000121116638,
-          95.53840160369873
+          95.5016016960144,
+          95.70919871330261
         ]
       },
       {
@@ -424,8 +424,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.3708028793335,
-          217.13440418243408
+          200.3783941268921,
+          191.52400493621826
         ]
       },
       {
@@ -434,8 +434,8 @@
         "width": 18,
         "exact": true,
         "candidate_us": [
-          110.25960445404053,
-          109.65759754180908
+          111.23199462890625,
+          82.0527970790863
         ]
       }
     ]
@@ -451,8 +451,8 @@
         "width": 409600,
         "exact": true,
         "candidate_us": [
-          30.607199668884277,
-          30.65280020236969
+          30.68319857120514,
+          30.62160015106201
         ]
       },
       {
@@ -461,8 +461,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.57120561599731,
-          202.49919891357422
+          208.86359214782715,
+          202.32839584350586
         ]
       },
       {
@@ -471,8 +471,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.633199572563171,
-          14.882799983024597
+          14.539200067520142,
+          14.544400572776794
         ]
       },
       {
@@ -481,8 +481,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          18.878400325775146,
-          18.8060000538826
+          18.855200707912445,
+          18.73600035905838
         ]
       },
       {
@@ -491,8 +491,8 @@
         "width": 10240,
         "exact": true,
         "candidate_us": [
-          128.46280336380005,
-          122.10240364074707
+          121.83560132980347,
+          120.99080085754395
         ]
       },
       {
@@ -501,8 +501,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          226.35560035705566,
-          198.80399703979492
+          201.1143922805786,
+          206.4228057861328
         ]
       },
       {
@@ -511,8 +511,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          17.788800597190857,
-          17.838799953460693
+          18.53799968957901,
+          18.149200081825256
         ]
       },
       {
@@ -521,8 +521,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          28.051599860191345,
-          27.460798621177673
+          27.94399857521057,
+          27.766400575637817
         ]
       },
       {
@@ -531,8 +531,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          46.43479883670807,
-          39.18159902095795
+          47.61880040168762,
+          41.60760045051575
         ]
       },
       {
@@ -541,8 +541,8 @@
         "width": 1280,
         "exact": true,
         "candidate_us": [
-          93.40720176696777,
-          79.79999780654907
+          83.53959918022156,
+          80.69199919700623
         ]
       },
       {
@@ -551,8 +551,8 @@
         "width": 16,
         "exact": true,
         "candidate_us": [
-          14.601600170135498,
-          14.541999995708466
+          14.508399367332458,
+          14.535599946975708
         ]
       },
       {
@@ -561,8 +561,8 @@
         "width": 524288,
         "exact": true,
         "candidate_us": [
-          95.30079960823059,
-          95.05879878997803
+          95.4695999622345,
+          96.30839824676514
         ]
       },
       {
@@ -571,8 +571,8 @@
         "width": 524304,
         "exact": true,
         "candidate_us": [
-          185.6376051902771,
-          216.42999649047852
+          199.35840368270874,
+          191.1903977394104
         ]
       },
       {
@@ -581,8 +581,8 @@
         "width": 18,
         "exact": true,
         "candidate_us": [
-          110.12599468231201,
-          109.82120037078857
+          111.20480298995972,
+          81.78759813308716
         ]
       }
     ]
diff --git a/docs/receipts/dsv41-neighbor-host/README.md b/docs/receipts/dsv41-neighbor-host/README.md
index ee5b5b4a79..b7ab8eec9b 100644
--- a/docs/receipts/dsv41-neighbor-host/README.md
+++ b/docs/receipts/dsv41-neighbor-host/README.md
@@ -39,3 +39,28 @@ The earlier experimental implementation used source
 separate patches. Its device and serving results support the transport design;
 they are not measurements of this consolidated branch or its startup agreement.
 No unrelated model, server, forced-tool, kernel or deployment changes were ported.
+
+## v0.2.0 repeat and local CI checks
+
+Rebased onto `b65c8b92d2403dfdefccf15a3b23ba609a4553e2`; node-tested code
+`22ac1f72e9e74d3f38f7aa2136bcccb432828cf0`. Zig 0.17.0 focused tests passed
+15/15, including configurable GID parsing; all 124 DeepSeek-family tests passed,
+including both startup neighbor-mode mismatch directions. Lean findings were
+1,615 on base and head, zero new. Linux node validation completed 100/100 steps,
+1,007 newly run tests passed, two skipped; other groups were cached. Native build
+with the existing SM121 fatbins passed (SHA-256
+`4f69f434dfd56c968ddbc70e89ba05fe79de054dee74dd32fcae2a03cf6c76d0`).
+
+Local CI commands at `45af9710f6b78afbcdb198ade9db62c868ad54fe` (same production
+code; later edits are documentation only) used Zig 0.17.0, Node 22.23.2 and
+npm 10.9.8. Ubuntu 24.04 arm64: `zig build test` passed 88/88 steps, 1,683 tests,
+seven skips; kernel-contract 18/18 and fuzz 7/7 passed. Goldens: 309/319 equal,
+ten known differences, zero unexpected. Whitespace and lean differential passed.
+macOS 27.2 arm64: desktop typecheck and mobile verification passed; desktop tests
+passed 547/548, with unchanged `sandbox-git-control.test.ts` failing because the
+sandbox blocks Xcode's `libxcrun.dylib`. macOS Zig retained the `SO.BUSY_POLL`
+compile failure (1,958 passed, eight skipped). Full-tree formatting fails on five
+unchanged base files: `zig/src/server/effort_values_test.zig` and
+`zig/src/families/dsv41/{exl3_linear,round,tri_basic,split_bench_main}.zig`.
+GitHub jobs did not start due to billing/spending limits; these local platforms
+differ from its runner matrix. These are not claims that all CI checks pass.
diff --git a/docs/recipes/deepseek-v41-neighbor.md b/docs/recipes/deepseek-v41-neighbor.md
index 9e36fbd7be..a58cb82f83 100644
--- a/docs/recipes/deepseek-v41-neighbor.md
+++ b/docs/recipes/deepseek-v41-neighbor.md
@@ -84,8 +84,6 @@ records the tested code commit, ring setup, byte/graph/token checks and serving
 measurements. It does not claim a direct comparison with a switched fabric or the
 target branch.
 
-The [v0.2.0 update receipt](../receipts/dsv41-neighbor-v020/README.md) records
-the rebased code checks separately from the historical hardware measurements.
-
-The [v0.2.0 hardware repeat](../receipts/dsv41-neighbor-v020-hardware/README.md)
-records fresh byte/graph/token checks and the same serving workloads on the rebased code.
+The [host and build receipt](../receipts/dsv41-neighbor-host/README.md) records
+the rebased code checks and local CI results. The hardware receipt above includes
+the fresh repeat and historical comparison of the same serving workloads.