diff --git a/README.md b/README.md index 0a1fa048ec..d52b5bf4c1 100644 --- a/README.md +++ b/README.md @@ -7,6 +7,7 @@ > 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-hardware/README.md b/docs/receipts/dsv41-neighbor-hardware/README.md new file mode 100644 index 0000000000..a521a9d1f5 --- /dev/null +++ b/docs/receipts/dsv41-neighbor-hardware/README.md @@ -0,0 +1,112 @@ +# TP4 neighbor transport: four-host hardware qualification + +Tested 2026-10-20 UTC at code commit +`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 +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 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). + +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` 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. + +## 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. 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 + 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 | 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. + +## 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 | **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 --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 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 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 +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..c5955bdc63 --- /dev/null +++ b/docs/receipts/dsv41-neighbor-hardware/summary.json @@ -0,0 +1,160 @@ +{ + "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", + "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": "6358c0a5c84923cebe3220ce2eb00d2210c93272921005d98b7eb82cbe5a239b", + "libtf-neighbor.so": "82897cfc01cf3aede02077a0ab7cc142f38828fc5ae11e4e762e9c6252d2347f", + "probe.py": "8d1900817f5e8b240c4be327fef72ca2f5bd187701c109f64007cc34b9d4dba2", + "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": 18.343599885702133 + }, + { + "bytes": 53760, + "slowest_rank_mean_us": 27.905799448490143 + }, + { + "bytes": 143360, + "slowest_rank_mean_us": 44.79059875011444 + }, + { + "bytes": 430080, + "slowest_rank_mean_us": 82.11579918861389 + } + ] + }, + "exact_tokens": { + "prompts": 3, + "paths": 3, + "tokens_per_reply": 64, + "replies": 9, + "pass_test": true, + "also_equal_to_original_commit": 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.5, + "solo_s": 7.8, + "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": 133.10816911985555, + "min_aggregate_tps": 132.48854455631133, + "max_aggregate_tps": 133.52409085662865, + "exact_utf8_replies": 24 + }, + "steady": { + "streams": 4, + "seconds": 30, + "max_tokens": 256, + "set": "a", + "requests": 19, + "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, + "seconds": 90, + "max_tokens": 256, + "set": "b", + "requests": 62, + "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 new file mode 100644 index 0000000000..05825ef086 --- /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.510801076889038, + 30.60159981250763 + ] + }, + { + "kind": 0, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 208.8956117630005, + 202.35960483551025 + ] + }, + { + "kind": 1, + "rows": 4, + "width": 16, + "exact": true, + "candidate_us": [ + 14.606000483036041, + 14.523200690746307 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 10240, + "exact": true, + "candidate_us": [ + 18.849200010299683, + 18.719999492168427 + ] + }, + { + "kind": 2, + "rows": 48, + "width": 10240, + "exact": true, + "candidate_us": [ + 121.89359664916992, + 121.00640535354614 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 201.19280815124512, + 206.52880668640137 + ] + }, + { + "kind": 3, + "rows": 7, + "width": 1280, + "exact": true, + "candidate_us": [ + 18.53639930486679, + 18.086400628089905 + ] + }, + { + "kind": 3, + "rows": 42, + "width": 1280, + "exact": true, + "candidate_us": [ + 27.906399965286255, + 27.90519893169403 + ] + }, + { + "kind": 3, + "rows": 112, + "width": 1280, + "exact": true, + "candidate_us": [ + 47.65839874744415, + 41.92279875278473 + ] + }, + { + "kind": 3, + "rows": 336, + "width": 1280, + "exact": true, + "candidate_us": [ + 83.07600021362305, + 80.69239854812622 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 16, + "exact": true, + "candidate_us": [ + 14.602400362491608, + 14.568400382995605 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524288, + "exact": true, + "candidate_us": [ + 95.41760087013245, + 97.06680178642273 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 200.18160343170166, + 191.20919704437256 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 18, + "exact": true, + "candidate_us": [ + 111.08880043029785, + 81.72240257263184 + ] + } + ] + }, + { + "rank": 1, + "pass_gate": true, + "mixed_graph": true, + "records": [ + { + "kind": 0, + "rows": 1, + "width": 409600, + "exact": true, + "candidate_us": [ + 30.526000261306763, + 30.595600605010986 + ] + }, + { + "kind": 0, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 208.95318984985352, + 202.19480991363525 + ] + }, + { + "kind": 1, + "rows": 4, + "width": 16, + "exact": true, + "candidate_us": [ + 14.607599377632141, + 14.505599439144135 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 10240, + "exact": true, + "candidate_us": [ + 18.774400651454926, + 18.72600018978119 + ] + }, + { + "kind": 2, + "rows": 48, + "width": 10240, + "exact": true, + "candidate_us": [ + 122.24199771881104, + 121.17400169372559 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 201.1728048324585, + 206.53119087219238 + ] + }, + { + "kind": 3, + "rows": 7, + "width": 1280, + "exact": true, + "candidate_us": [ + 18.56359988451004, + 18.093200027942657 + ] + }, + { + "kind": 3, + "rows": 42, + "width": 1280, + "exact": true, + "candidate_us": [ + 27.94800102710724, + 27.698001265525818 + ] + }, + { + "kind": 3, + "rows": 112, + "width": 1280, + "exact": true, + "candidate_us": [ + 47.57080078125, + 41.88840091228485 + ] + }, + { + "kind": 3, + "rows": 336, + "width": 1280, + "exact": true, + "candidate_us": [ + 83.43240022659302, + 80.7099997997284 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 16, + "exact": true, + "candidate_us": [ + 14.508000016212463, + 14.536799490451813 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524288, + "exact": true, + "candidate_us": [ + 95.46639919281006, + 96.30399942398071 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 200.1267910003662, + 191.1736011505127 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 18, + "exact": true, + "candidate_us": [ + 111.23640537261963, + 81.77880048751831 + ] + } + ] + }, + { + "rank": 2, + "pass_gate": true, + "mixed_graph": true, + "records": [ + { + "kind": 0, + "rows": 1, + "width": 409600, + "exact": true, + "candidate_us": [ + 30.682000517845154, + 30.62080144882202 + ] + }, + { + "kind": 0, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 208.96799564361572, + 202.25400924682617 + ] + }, + { + "kind": 1, + "rows": 4, + "width": 16, + "exact": true, + "candidate_us": [ + 14.661599695682526, + 14.50120061635971 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 10240, + "exact": true, + "candidate_us": [ + 18.75, + 18.72040033340454 + ] + }, + { + "kind": 2, + "rows": 48, + "width": 10240, + "exact": true, + "candidate_us": [ + 122.23520278930664, + 121.12840414047241 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 201.13320350646973, + 206.45320415496826 + ] + }, + { + "kind": 3, + "rows": 7, + "width": 1280, + "exact": true, + "candidate_us": [ + 17.870399355888367, + 18.094399571418762 + ] + }, + { + "kind": 3, + "rows": 42, + "width": 1280, + "exact": true, + "candidate_us": [ + 27.934399247169495, + 27.528798580169678 + ] + }, + { + "kind": 3, + "rows": 112, + "width": 1280, + "exact": true, + "candidate_us": [ + 47.58960008621216, + 41.90320074558258 + ] + }, + { + "kind": 3, + "rows": 336, + "width": 1280, + "exact": true, + "candidate_us": [ + 83.41760039329529, + 80.70080280303955 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 16, + "exact": true, + "candidate_us": [ + 14.497199654579163, + 14.518000185489655 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524288, + "exact": true, + "candidate_us": [ + 95.5016016960144, + 95.70919871330261 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 200.3783941268921, + 191.52400493621826 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 18, + "exact": true, + "candidate_us": [ + 111.23199462890625, + 82.0527970790863 + ] + } + ] + }, + { + "rank": 3, + "pass_gate": true, + "mixed_graph": true, + "records": [ + { + "kind": 0, + "rows": 1, + "width": 409600, + "exact": true, + "candidate_us": [ + 30.68319857120514, + 30.62160015106201 + ] + }, + { + "kind": 0, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 208.86359214782715, + 202.32839584350586 + ] + }, + { + "kind": 1, + "rows": 4, + "width": 16, + "exact": true, + "candidate_us": [ + 14.539200067520142, + 14.544400572776794 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 10240, + "exact": true, + "candidate_us": [ + 18.855200707912445, + 18.73600035905838 + ] + }, + { + "kind": 2, + "rows": 48, + "width": 10240, + "exact": true, + "candidate_us": [ + 121.83560132980347, + 120.99080085754395 + ] + }, + { + "kind": 2, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 201.1143922805786, + 206.4228057861328 + ] + }, + { + "kind": 3, + "rows": 7, + "width": 1280, + "exact": true, + "candidate_us": [ + 18.53799968957901, + 18.149200081825256 + ] + }, + { + "kind": 3, + "rows": 42, + "width": 1280, + "exact": true, + "candidate_us": [ + 27.94399857521057, + 27.766400575637817 + ] + }, + { + "kind": 3, + "rows": 112, + "width": 1280, + "exact": true, + "candidate_us": [ + 47.61880040168762, + 41.60760045051575 + ] + }, + { + "kind": 3, + "rows": 336, + "width": 1280, + "exact": true, + "candidate_us": [ + 83.53959918022156, + 80.69199919700623 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 16, + "exact": true, + "candidate_us": [ + 14.508399367332458, + 14.535599946975708 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524288, + "exact": true, + "candidate_us": [ + 95.4695999622345, + 96.30839824676514 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 524304, + "exact": true, + "candidate_us": [ + 199.35840368270874, + 191.1903977394104 + ] + }, + { + "kind": 3, + "rows": 1, + "width": 18, + "exact": true, + "candidate_us": [ + 111.20480298995972, + 81.78759813308716 + ] + } + ] + } +] 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 new file mode 100644 index 0000000000..b7ab8eec9b --- /dev/null +++ b/docs/receipts/dsv41-neighbor-host/README.md @@ -0,0 +1,66 @@ +# 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). +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 | +| --- | --- | +| `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. 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 +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 new file mode 100644 index 0000000000..a58cb82f83 --- /dev/null +++ b/docs/recipes/deepseek-v41-neighbor.md @@ -0,0 +1,89 @@ +# 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). +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. + +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 + +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 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. 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 +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. + +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. + +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. diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig index e0429d8c46..885ab97d45 100644 --- a/zig/build/cuda.zig +++ b/zig/build/cuda.zig @@ -347,6 +347,10 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st 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", 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); 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..6442e43d81 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; @@ -229,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(); @@ -241,6 +249,7 @@ 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(); @@ -297,9 +306,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, neighbor_gid); + 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 +355,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 +417,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 +445,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 +591,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 +655,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 +715,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 +732,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 +765,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 +803,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 +818,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..d4747fd8be --- /dev/null +++ b/zig/src/families/dsv41/neighbor.zig @@ -0,0 +1,308 @@ +//! 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"); + 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; +} +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 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 { + 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, 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; + 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 = 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); + 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 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" { + 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..6ecef491f3 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"); @@ -143,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" @@ -167,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(); } @@ -183,6 +188,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 +221,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 +293,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 +329,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 +353,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 +387,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 +436,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; @@ -447,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" { @@ -455,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))); @@ -499,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)); +} 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().*); +}