Skip to content

DeepSeek-V4.1: performance hooks and host checks - #707

Open
jayleaton wants to merge 56 commits into
ashhart:mainfrom
jayleaton:stack/13a-performance-host
Open

jayleaton wants to merge 56 commits into
ashhart:mainfrom
jayleaton:stack/13a-performance-host

Conversation

@jayleaton

@jayleaton jayleaton commented Oct 11, 2026 •

Copy link
Copy Markdown
Contributor

What this changes

DeepSeek-V4.1: performance hooks and host checks. Depends on #703. This cumulative PR targets main; the incremental range is 59af121b..3c93b75f.

#703 → #707 → later DeepSeek-V4.1 layers

Performance hooks remain. In-tree Python checks become native checks against stored fixtures. No new speed or hardware-parity claim is made.

Commits in this layer:

  • 3c93b75f DeepSeek-V4.1 performance extras, host side: side streams, the Engram gate, L2 prefetch, profile, device arena.

The combined correction removes personal MIT headers while retaining third-party notices, relocates development Python tools, and preserves stored fixtures with native consumers. Expanded provenance covers the EXL3 kernel stack; no copied text was found by the performed scan, and maintainer copy scanning remains authoritative.

Receipt

Before: public head ec25c086. After: exact corrected head 3c93b75f07c09536a25e44c64bd08d84217ee942, gated on 2026-10-12, Linux x86_64 and Zig 0.17.0.

Executed gate output:

G707: commit 3c93b75f07c09536a25e44c64bd08d84217ee942 tree 07547ce1de58a8ec4c3376af7708c7ca65696c6c dirty 0 2026-10-12T06:11:38Z
install rc=0 Build Summary: 28/28 steps succeeded
test rc=0 Build Summary: 120/120 steps succeeded; 943/951 tests passed (8 skipped)
test-tp rc=0 Build Summary: 4/4 steps succeeded
test-dsv41 rc=0 Build Summary: 56/56 steps succeeded
test-kv rc=0 Build Summary: 8/8 steps succeeded
test-dsv41-gm2pf rc=0 Build Summary: 6/6 steps succeeded
test-dsv41-longpf rc=0 Build Summary: 6/6 steps succeeded
test-dsv41-sampling rc=0 Build Summary: 4/4 steps succeeded
test-dsv41-draft rc=0 Build Summary: 3/3 steps succeeded
test-dsv41-rows rc=0 Build Summary: 6/6 steps succeeded
test-dsv41-serve rc=0 Build Summary: 3/3 steps succeeded
test-server-cpu rc=0 Build Summary: 5/5 steps succeeded
test-dsv41-knobs rc=0 Build Summary: 6/6 steps succeeded
test-dsv41-wide rc=0 Build Summary: 6/6 steps succeeded
test-dsv41-slot-drafts rc=0 Build Summary: 6/6 steps succeeded
tf-kv-split-test rc=0 Build Summary: 4/4 steps succeeded
receipt receipt_cli.sh rc=0 RESULT 0 failed
receipt receipt_cli12b.sh rc=0 RESULT 0 failed
receipt receipt_cli13.sh rc=0 RESULT 0 failed
receipt receipt_model.sh rc=0 RESULT rc=0
zig fmt --check rc=0
git diff --check rc=0
3c93b75f lean: 1506 files, 534 problems | new: 0 | gone: 0
FAILED STEPS 0

Lean uses pinned main e7c523eb: 534 inherited findings, zero new or removed identities. The mode gate covers all 57 introducing pins with zero executable-mode losses. Family suites and CLI receipts run where introduced. Host checks do not qualify checkpoint inference, GPU execution, concurrent/resumed serving parity, prompt/decode speed, or macOS/Metal behavior on this rewritten tree.

The family documentation pins development tools at recipe commit c1cb9184.

Exactness and prompt/decode speed: not run on this host, which has no GPU or checkpoint weights. No new performance claim. MLX/PyTorch/checkpoint revision: not applicable to the host gate.

Checklist

  • Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off.
  • Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo.
  • No precision traded for speed. If the bits change, the description says which sums change and why.
  • Prompt processing is no slower than the last release.
  • The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
  • New tests fail before the change, pass after it, and skip cleanly without their dependency.
  • Comments and docstrings are one line. No measurements or history in the source.
  • No personal data, machine names, internal hosts or local paths. No AI attribution lines.

Known limitation carried from #702: the missing-tokenizer error-message allocation leak remains unresolved.

Merge Danger

Door: two-way. Reverting restores the parent without a data migration.

Blast Radius: engine. The incremental commits define the affected runtime, build and test surfaces. The production branch is unchanged.

jayleaton and others added 30 commits October 11, 2026 17:13
…alls them.

graph_cache.Cache captures one executable graph per caller-named key and replays it when the key comes back; ranks agree on every capture, eviction and fallback through an Agree hook, and a memory floor turns misses eager.
arena.Arena sub-allocates DeviceBuffer.alloc from a few large chunks (VMM or plain cuMemAlloc) through memory.setHook; borrowed spans and pointers it does not own are freed as before.
…; CUresult names fall back to a table when the driver cannot answer.

cuFuncGetName is optional (CUDA 12.3+), so an older driver still opens. During
the exit's teardown cuGetErrorName fails too, which printed
CUDA_ERROR_DEINITIALIZED as "CUDA_ERROR_UNKNOWN_CODE (4)"; code 4 now says the
process is exiting.
Use CUDA device attribute 110 for GPUDirect RDMA with VMM support. Attribute 126 reports the memory sync domain count and can incorrectly enable the RDMA allocation flag. Add a regression for the external ABI contract.
…r stays under the max, a max of 0 is eager

- A capture whose agreement throws, or whose insert runs out of memory, no longer leaks its executable: the map's
  slot is reserved before the capture, and the executable is owned here until the insert, which cannot fail.
- A host map that cannot grow counts as no room in the existing room vote, so every rank declines that miss together.
- `min_keep` is clamped to `max`, so under the memory floor the cap never rises above the configured max. A max of 0
  runs every step eagerly, with no collective.
- A cap lowered to 0 under the floor goes back to max once every rank has room, instead of staying eager for good.
- Hits make no collective: with `Agree` set, keys, settings and fingerprint changes must already match on every rank.
  The docs now say so: the fingerprint changes at the same step on every rank, and `fingerprintOf` is for one rank.

Co-authored-by: Jonathan Regan-Beasley <197533392+jregan-beasley-smc@users.noreply.github.com>
The arena's `live` (small ranges) and `large` (own chunks) maps drop an entry on every free and gain one on every
allocation for the life of the process. std's tombstones never clear, so in time they take every free slot, and each
allocation, free and miss then probes the whole map. Each map now counts its removals and is rehashed in place once
they reach a tenth of its capacity.

Test: 20,000 random allocations and frees (small ranges and large chunks) keep a tenth of each map free at every step;
without the rehash it fails.
…d RoCE collectives, plan link, fail-fast.

zig/src/fabric/tp brings a rank up from TF_TP_WORLD, TF_TP_RANK, TF_TP_MASTER and TF_TP_PORT, agrees the settings
digest over every rank, and gives a family all-gather, all-reduce and broadcast over NCCL (TF_COMM_BACKEND=nccl,
the default), a pinned host mailbox for ranks on one host, or RoCE verbs between two hosts (TF_TP_ROCE_*), with an
NCCL fallback when RoCE setup fails. Rank 0 sends each round's plan to the others over its own link (TF_TP_PLAN_*).
A failure or a vanished rank ends every rank shortly after (TF_TP_FAILFAST, TF_TP_FAILFAST_GRACE_S).
No family uses it yet. zig build test-tp runs the host tests; zig build tf-tp-test builds the two-GPU runner
(-Dtp-nvcc= or -Dtp-fatbin= for the mailbox kernel). The RoCE all-gather adapts b12x's RoCEnante protocol.
… the stop word with no fail-fast exit; a rank killed mid-drain still fails the other fast.
… /24 the peer also has

The automatic pick took the first N active RoCE ports in device-list order and never checked that the peer can reach
them. On a host with two neighbours, one picked port can sit on a subnet the peer does not have, and the queue pair's
RTR step then fails, which reports RoCE unavailable.

Every rank now sends its usable ports' IPv4 addresses over the bootstrap. Each rank then keeps the ports whose /24 the
peer also lists, ordered by subnet and each /24 no more often than both sides have it, so both ranks pick the same
number. With no shared /24 the old first-N order is kept, with a warning to set TF_TP_ROCE_HCA. The exchange runs on
every rank, including ranks that set TF_TP_ROCE_HCA, so the bootstrap stays in step.
…FNAME and TF_TP_MASTER are needed, with a two-host example
… kernel backends up front; run_hosts.sh passes HCA picks

- `tf-tp-test kernel` reports the embedded kernel and exits 1 without one. `suite` and `failfast` on the mailbox or
  RoCE backend stop at once with a message that names `-Dtp-nvcc`, instead of failing later on a missing kernel.
- `run_hosts.sh` stops before any step when the kernel is missing. It exports TF_TP_ROCE_HCA and NCCL_IB_HCA (derived
  from TF_TP_ROCE_HCA when unset), warns when a host has more than one RDMA device and neither is set, and records all
  three variables in host.txt.
… with another region layout are refused

roce_guard.stripes computes each HCA's share of a post (its send stripe, the peer's receive range and its flag word)
and refuses with error.OutOfRegion any range past its slot, its flag area or the registered region, before
post_send. The verbs wire and the shm wire both post through it. The shm wire moves to roce_shm.zig and can stripe
like the verbs wire (Options.shm_stripes), with an optional tap that sees every post for tests. The connection record
carries the slot and registered length, so a peer with another layout is refused at connect; its magic is bumped.

Tests: stripes tile the slot over two HCAs and refuse a post past the slot, a third slot, an unpadded or empty post,
no HCA and a short region. On fake verbs a post past its slot reaches no verb, and a full slot stays inside the
region. Over the shm wire striped as two HCAs, 2,000 variable-length gathers stay inside slot and region, bit for bit.
…would ask for one

A post of about 9.6 MB faulted a DGX Spark's HCA. roce_guard.max_write (1 MiB, a constant) is checked in stripes()
before the slot bounds, so the verbs wire, the shm wire and the proxy refuse a larger post with error.WriteTooLarge
before any verb, whatever the slot size. Roce.init refuses a slot past it. Config.parse fails boot with a clear
message when TF_TP_MAILBOX_MAX_KB would make the RoCE slot pass 1 MiB; the value is never clamped.

Tests: a 9.57 MB and a 1 MiB + 16 post are refused in a 10 MiB slot and reach no verb, a 1 MiB post goes out; boot
refuses TF_TP_MAILBOX_MAX_KB 1028, 2048 and 16384 and keeps 1024.
…sized ops, at most 2 MiB in flight

Default off. allGatherV strides past the slot, up to TF_TP_ROCE_CHUNKED_KB, go RoCE as ceil(n / slot) launches of
the tp_roce kernel: op c moves at most a slot of each shard from c x slot (RoceArgs.chunk), and an op empty on both
ranks rings no doorbell and leaves the epoch. Every RDMA write keeps the 1 MiB cap. The kernel poisons the transport
rather than stage past a slot, and its four failure-record blocks fold into fail_record (same stores and fences).

VerbsWire.post waits for completions while the bytes posted and not completed over both HCAs would pass
max_inflight (2 MiB, a constant in the post path); a wait that sees no completion for 10 s fails the proxy. Boot
refuses TF_TP_ROCE_CHUNKED_KB when the kernel image lacks the tp_roce_chunk_abi marker. The setting is part of the
digest both ranks agree on.

Tests: chunked strides from 300 KB to a 128 MiB union posted back to back on fake verbs whose completions lag keep at
most 2 MiB in flight and no write past 1 MiB, and reach the cap; the wire's own in-flight count and peak match the
fake HCAs'. Over the shm wire, 2,000 chunked variable-length gathers of up to four slots arrive bit for bit.
…region; the region is DONTFORK and NOHUGEPAGE

A local protection error can arrive while the IOMMU invalidates its mappings lazily. On such an error the old path
left the other HCA's QP ready to send (still taking the peer's writes), hung up the fate channel without telling
the peer to stop, and freed the region at once: unregistered from CUDA with no GPU sync, then QP destroy, MR
deregistration and munmap. A device write after the unmap could then reach a reused page through a stale IOTLB entry.

- Roce.fail and Roce.abort (the fail-fast hook) move every QP to the error state: no send goes out and no peer write
  lands.
- Roce.deinit quiesces first: QPs to error, every posted request reaped (flushed or not), cuCtxSynchronize. A failed
  transport, or one that does not quiesce, keeps its region registered and mapped; the process exit tears it down.
- Session.close on a failed transport goes fail-fast before closing, so the peer aborts and halts its HCAs first.
- The region is MADV_DONTFORK and MADV_NOHUGEPAGE before ibv_reg_mr; a refusal fails the start.
- A QP whose actual max_inline_data is under 4 B is refused (the flag post uses lkey 0 on a stack word).
- One log line a rank: region base and length, each HCA's lkey and rkey, the peer's region and rkeys.

Tests: on fake verbs Roce.abort halts both QPs; quiesce halts both before the first flush and reaps every posted write;
an aborted transport reports itself failed, and over the shm wire the aborted rank keeps its region while the other
frees it; smaps shows dc and nh on a prepared region; the inline bound refuses 3 B and takes 4.
…path on real HCAs, only on an idle link

- `tf-tp-test failfast V transport [N [HOLD_S]]`: N clean 4 KiB sums, the ranks meet over NCCL, each waits for its
  link to go idle (every HCA's posts completed), both agree over NCCL, then the victim takes the proxy's failure path
  (`Roce.inject`, which is `Roce.fail`) and closes its session: both ranks halt their QPs and exit 70, and the
  victim's region stays registered. When either rank's link is not idle there is no inject, a clean stop, REFUSED and
  rc 2.
- Its host-only forms (TF_TP_FAILFAST=0, or a hold) are refused unless TF_TP_ALLOW_HOST_ONLY=1 and no RDMA wire is
  in use (the shm wire, or no RoCE), so they never run on an HCA.
- `run_transport.sh MASTER PEER_SSH OUT`, driven from rank 0's host (rank 1 over ssh): prechecks on both hosts (no
  GPU process, memlock unlimited), then T0 (suite), T1 and T3 (each rank as the victim). After each step it reads
  both hosts' kernel log since the step and every HCA *err* counter; a hit, a counter that grew, an unexpected rc or a
  failure line prints STOP: and exits 1. `run_transport.sh t2 OUT` runs the host-only form on one host.
- The `halted` and `quiesced` lines carry each QP's state read back with ibv_query_qp (an optional symbol), each
  HCA's posts and completions, the completions the drain reaped and how many of them were flushes.
- `Roce.allGatherV` refuses a stride past the slot (one op; the route never sends one).
- The README states the RoCE limits and how to run the check.

Tests: on fake verbs the link reads busy with two posts out and idle after the drain, and the report line names both
QPs in the error state with every completion a flush.
The prefix index keeps one node per distinct full page of every saved entry in `by_digest`. Each saved turn adds its
new pages, and each entry dropped or evicted removes the pages no other entry passes through, for the life of the
server. std's tombstones never clear, so in time they take every free slot, and each save's new pages and each find's
walk past the cached prefix then probe the whole map. The map now counts its removals and is rehashed in place once
they reach a tenth of its capacity.

Test: 3,000 turns of up to 60 fresh pages, the oldest dropped past eight, keep a tenth of the map free after every
turn; without the rehash it fails.
A family can now hand the lane host a round planner (api.Rounds) that admits
requests, splits prompts into pieces and runs them in the same rounds as the
other streams' decode rows. The planner is opt-in: the host's own loop, with a
backend's prefill_step a chunk a loop between rounds, stays the default, and a
host with a planner runs the planner's loop instead (LaneHost.run picks one).
The backend gains optional entries for this: begin, piece, pieces (several
streams' pieces in one call), finish and join, so a round's finished prompts
get their first drafts in one backend call.

A planner's prompts open through lanes/open.zig like the whole and the chunked
pass (open.openedWith; their first drafts sized by the head's reach under node
probabilities), and a prompt cancelled at its final releases its lane as a
cancelled chunk does.

The lane core also takes a family row policy (lanes/trim.zig) in place of the
built-in allocation, proposers that price their own drafts, and a timed switch
for backends without round timers. Under a planner the host still steps the
learner while no request waits or runs, and a failed piece or piece run fails
only its own requests.

A planner over a backend without prompt pieces (such as the external drafter
wrapper) is refused at start with error.PiecesUnsupported.

The round loop lives in lane_rounds.zig, a request's stream setup in
lane_stream.zig and the pieces' steps in lanes/pieces.zig, keeping the host
and the engine under the line limit.
…oundary, and every request queued, filling or running fails with the reason.

Opened.halt (optional) lets a server stop an engine between rounds instead of
under a running one; LaneHost.halt is that hook for lane-host families. A
request still filling its prompt a chunk a round fails too, with its lane
freed, and a prompt pass the halt cancels mid-chunk reports the halt's reason
rather than a cancel.
…s waits time out

The second prompt was submitted after start(), so a loaded scheduler could let the first finish its pieces before
the second arrived: no queued admission, and the test's waits > 0 check failed. Both are queued now before the first
round. Each wait gives up after 10 s, and a timed-out wait stops the host thread before the test's frame is gone.
family.zig is an optional vtable a family installs in place of the chat template: prompt rendering, thinking and
effort fields, a streaming reply reader that parses its own call markup, and the turn it remembers. Without a
family the server runs as before.

chat.zig's reply collection moves to chat_generation.zig and openai.zig's run to openai_run.zig (pure moves), so
neither file passes 600 lines with the hooks in.
…pt loop ends, and again if it exits with connections open.
New generation requests get 503 with Retry-After (code server_restarting;
Anthropic overloaded_error) and /health answers "draining" without asking the
engine. Requests in progress get TF_DRAIN_S seconds (default 20; a second
signal cuts it short), then the engine's optional halt ends its rounds at a
boundary and what is left fails with "server restarting: retry shortly". The
engine closes before the server is freed, once.
…in check refuses the request (503) and the drain waits for it.

A tests-only seam (routes.counted_hook) delivers SIGTERM in that gap; with the
drain check moved before the count the test fails.
…ery TF_SSE_KEEPALIVE_S seconds (default 15).

Queueing and a long prefill send nothing after the head, and clients or proxies
that time out a silent body leave and retry from zero. Chat, completions and
responses streams get an SSE comment (": keepalive"); Anthropic streams get its
ping event. A keepalive that cannot be written cancels the request as a client
that left.
…ain checks are done

The listener stays open only while a generation runs. Case 1 streamed a "@short" reply (about 180 ms) and then
sent a messages request and a /health request expecting 503; when the reply ended first, the drain finished and
those requests found the listener closed. The stub engine's "@Hold" reply now waits for the test's release.
…TP weight plan and its named tensors.

zig/src/families/deepseek_v41 starts with the parts that need no GPU: the text config (config.json's text_config,
with the per-layer CSA2 topology), EXL3 group geometry and a rank's TP slices, the pack's tensor names and shapes,
the pack folder read from its safetensors headers alone, the rank's weight plan, its host images and the weights
as named tensors in the reference loader's forms. zig build test-dsv41 runs the host tests on the release config,
a synthetic pack and a 2.8 bpw pack's expert widths; TF_DSV41_HEADERS adds real header dumps.
… cached on disk for a fast boot.

prepare.zig builds the run-time forms from a rank's named host tensors: linear strips and the routed and shared
experts' tables; its test checks the dense lanes layout at every width against a reference (fixtures/lanes-ref.bin).
prepared.zig writes and reads built layer images under TF_DSV41_PREPARED (_WRITE builds them, _VERIFY refuses a
damaged payload), keyed by the loader sources that build them, so a changed loader never reads an older folder.
…ed phase (TF_DSV41_PHASES).

TF_DSV41_PHASES=1 times each named phase of a round on the host clock, =sync also waits for the stream at each
boundary, and TF_DSV41_PHASES_OUT=<file> writes the totals as JSON. Off by default; nothing reports into it yet.
…ow's tokens, a rank's heads.

engram_host.zig lays out the Engram tables of the config, hashes each window's token n-grams into table rows as the
reference hasher does (checked on the release config and tokenizer, fixtures/engram-hash-ref.bin), splits the
heads over the TP ranks and builds the synthetic index rows the tests use.
…LRU-cached, read ahead of the window.

engram_rows.zig serves a window's Engram rows from the rank's packed shards: io_uring where it works, else Linux
native AIO (TF_DSV41_ENGRAM_AIO=1, TF_DSV41_ENGRAM_POLL) or a pool of reader threads (TF_DSV41_ENGRAM_QD, 0 reads
serially); TF_DSV41_ENGRAM_URING=0 picks the threads even where io_uring works. Reads can be issued ahead and
waited for later, bulk reads bypass the LRU cache, and eviction keeps the bits. The tests compare the bf16 table
and a packed shard with the reference reader's rows (fixtures/engram-rows-ref.txt) on every read path.
…ter, EXL3 experts and linears, glue.

zig/kernels/cuda/deepseek_v41 holds the decode kernels as .cu sources. The reference extensions' kernels are copied
unchanged (each names its source and blob): CSA2 attention with its split kernel, the top-k, mHC with its parallel
tail, the router GEMV with its narrow kernel, the EXL3 routed experts, x3ld, x3gm with its v2 kernel and the EXL3
linear with its decode headers. Ours, each bit-exact with the torch ops it replaces: the window glue and the prefill
MoE's pointwise ops (--fmad=false), the decode router glue, x3gm's pass tables and the logits' row top-k.
ptxas.sh compiles them in a CUDA image with -Xptxas -v to report spills and stack. Nothing builds them yet.
…yout.

nvcc-docker.sh runs nvcc from the NGC PyTorch image for hosts without CUDA (zig build -Dnvcc=<script>). It mounts the
build root and the repository root (else the directory) of every absolute argument at the same path, so sources and
Zig caches may live anywhere, and the -MF file keeps only dependencies under a mounted root, matched on whole path
components (in bash, no python3 on the host), so a header edit rebuilds the fatbin while the image's own headers drop
out. TF_ZIG_DOCKER_MEMORY caps each container's memory and swap.

docs/recipes/deepseek-v4.1-cuda.md links the family's reference tools, pinned to a commit of
github.com/jayleaton/deepseek-v41-tensorfold-spark (tools/tensorfold-dsv41/), shows how to run the copies'
SASS-equality check from there and lists the stored fixtures the host suites check natively; later layers add theirs.
…symbols, shapes and launch arguments.

kernels_exl3.zig (with kernels_x3gm.zig) binds the routed-expert kernels, kernels_dense.zig (with kernels_lin2.zig)
the dense EXL3 linears, and kernels_ops.zig (with kernels_shapes.zig and kernels_glue_ops.zig) attention, top-k,
mHC, the router and the glue: the functions each module resolves, checked against the device's shared-memory
opt-in, and the shape checks before a launch. kernels_kv_glue.zig, kernels_replay*.zig and sampling_kernels.zig
bind the KV store, the prefill replay and the keyed sampler. kernels.zig loads every module and resolves every
instance; it is the dsv41_kernels module. The bindings name kernels whose sources later commits add, and nothing
builds fatbins yet: zig build test-dsv41 compiles and tests the bindings on any host against the CUDA runtime
without kernel images.
…he reference extensions' outputs.

zig/tests/cuda/dsv41_main.zig is tf-dsv41-test: symbols (every instance resolves), experts (the decode chain at every
setting), x3gm (fast prefill at every configuration), replay (a recorded extension call re-issued, storage by
storage) and all (every fixture directory by its manifest's kind, including the plan, top-k keys, pointwise and glue
cases). The fixtures come from the reference extensions on seeded inputs (oracle.py, and dsv41_capture.py for
replay; the family's reference tools, tools/tensorfold-dsv41/ in github.com/jayleaton/deepseek-v41-tensorfold-spark).
zig build test-dsv41 compiles the runner on any host; zig build installs it on Linux. Running it needs a GPU and every
family kernel image, which later layers add.
…gm ablations, glue and MMA rate benches.

zig/kernels/cuda/deepseek_v41/proto keeps the standalone benches the kernels were tuned with: attention with a
latency-hiding merge and with phase timestamps, x3gm timing ablations and per-rank-shape benches, the decode glue
at 1-16 rows, mma.sync issue rates, a router glue oracle, a smaller-GPU emulation header (blocker CTAs on the
other SMs) and an nsys export summary. None is part of the build.
…time forms, pinned staging, Triton cubins.

load.zig puts a rank's named weights on the GPU (one device buffer a tensor, bound by name, digested on the way up),
forms.zig builds the run-time forms from them (strips, lanes, bf16 mHC, expert tables) and forms_inplace.zig can
write the dense forms over their trellis (TF_DSV41_FORMS_INPLACE); ex_stacks.zig reads the experts' suh / svh stacks
where the weights hold them (TF_DSV41_EX_STACKS). block_calls.zig is the list of calls a block emits and buffers.zig
the device buffers they name. host_stage.zig is the pinned staging a prompt segment rewrites without draining the
stream, and triton_capture.zig loads the reference capture's Triton cubins (aot/aot.json, by hash) and packs their
launch arguments as Triton does. Compiled on the host by zig build test-dsv41.
… sparse indexer, mHC, MoE, Engram, L2 prefetch).

block.zig turns a decode window into calls (block_calls.zig's list, tensors named by role): block_attn.zig the
RoPE, the KV stores and pools, SWA and CSA2 attention; block_index.zig the indexer's scores, selections and top-k;
block_moe.zig the mHC sites, the router and experts, Engram and the head; block_prefetch.zig the L2 prefetch
(TF_DSV41_L2PF); block_wide.zig windows of more than 16 rows. knob_types.zig holds the option types block and
forward options carry (the fused dense prefill's choices, the x3gm v2 mode) and forward_types.zig the plain types
they share with the forward. Compiled on the host by zig build test-dsv41; the forward that issues the calls comes
later.
… priced floor, the round planner and sessions.

kv/ is the family's `kv` module over `sessions`, `cuda` and `tp`: layout.zig the KV families per token,
split.zig split KV's exchange (dense, packed, deduplicated, unions) with its host kernels, budget.zig the boot
budget and the measured floor, device.zig the device pool, sess.zig a slot's bounded snapshot, sched.zig the
round planner and the memory floor (pieces saturate at the budget left), progress.zig its progress log.
Host tests run the planner against sched_golden.json, sessions over 4 live slots against
session_slots_golden.json (and the split store's golden), resends from NVMe, unions and the exchange on
in-process ranks, also over RoCE's shm wire at the served settings through tp's routes.
zig/build/kv.zig adds `test-kv` (part of `zig build test`; it also compiles the device pool's runner),
`tf-kv-bench` and, on Linux, `tf-kv-gpu-test`. The page cache kept out of the floor is
TF_DSV41_ADMIT_CACHE_KEEP_GB.
…V's device exchange and the decode graphs.

forward.zig runs a window's calls on a rank (run.zig's launches, Triton through triton_call.zig and the AOT set),
with its parts beside it: forward_glue.zig the glue steps, forward_engram.zig Engram rows, forward_follow.zig the
follower's plan ops, forward_sample.zig the pick, forward_graph.zig and graphs.zig the decode graphs (buckets,
statics, the memory floor and the capture meter), kv_state.zig the paged pool and split KV under the forward.
forward_hooks.zig holds the hooks later stages install (prefill, Engram gate, round graph, images, branches);
unset, their paths stay off. kv/split_gpu.zig is `tf-kv-split-test`, split KV's exchange on two GPUs, and
kvsplit.cu / kv_glue.cu are its copies and the fused KV RMS and store, checked by tf-dsv41-test's `kvsplit` and
`kvglue` commands. The family module now imports `tp` and `kv`; zig build test-dsv41 compiles the forward and runs
the graphs' and greedy step's host tests, and test-kv compiles the split runner.
…replay, long prompts, x3gm v2 and its kernels.

- block_prefill and its attention, emit, mHC, MoE and select parts: a prompt segment's program, short segments on the
  decode window's row paths, stream top-k past TF_DSV41_INDEX_STREAM_MIN keys.
- forward_prefill with finish, glue and 4K parts: the forward's prefill hook (segments committed as keeps, CED's
  decoder replay, several slots' segments in one call), installed by engine open; ced: CED's keep and hand-off rules.
- longpf (split KV's row blocks for long prompts), gm2pf (x3gm v2 on prefill, TF_DSV41_GM_V2), rowtab and rowmode
  (row mode's table and program), serve_knobs (the served settings' readers), calls_check (the emitters' calls
  against a capture's launches, on the host).
- image_opts: the block options for image prompts as pure functions; the prefill glue's image steps go through two
  new forward_hooks.Vision entries (keep, split), unset until images are installed.
- Kernels: dense3, mhc_pf, pfdense (k8/k10/k12/k16), prefill_glue, topk_b, x3gm3, x3ld_epi, x3pf, x3seg.
- Tests: the family's host suite compiles every prefill file with the hooks unset and runs the CED, row table, image
  option, several-slot and stream read-back tests; test-dsv41-gm2pf and test-dsv41-longpf run in `zig build test`.
  The device runner gains the prefill glue and decode mHC cases and mhc-probe.
…e, its device steps and sampling.cu.

- serve/sampling: the exact sampler (rank 0 picks each row keyed by seed, position and id, so drafted equals serial),
  the first file of the family's serving module `dsv41_serve`, with its golden fixtures/serve/sampling.jsonl.
- sampling_pipe: a window's keyed choices from the ranks' logit slices (candidates, nucleus statistics, keyed pick);
  sampling_dev and sampling_gpu: its device steps and the forward's sampler (rank 0 sends each call's sampling on the
  plan link, followers run the same steps); sampling.cu: nucleus statistics and packed candidate rows.
- zig/build/sampling.zig: test-dsv41-sampling (the pipeline on in-process ranks against the golden, and the serving
  module's tests) in `zig build test`; tf-dsv41-samp (sampling_check, the device steps against the host rule on one
  GPU) installed on Linux and compiled by `zig build test` on any host. The family module imports `lanes` and
  `dsv41_serve`; its host suite compiles the pipeline and the device steps.
…, trees, lookup and the lanes backend.

- draft/: the drafting module over iface.zig's `Target` and `Pass`: dspark (the Markov chain, candidates, picks), costs
  and calib (the verify / DSpark cost table and its cache file; calib_env, calib_measure), joint and depth (the DSpark
  depth from costs), tree and branches (draft trees and their resolver), lookup (suffix drafts), spec (the next round's
  pass started after a verify window), policy (the depth as the lane engine's row policy), lanes (the lane backend over
  a target and a pass), twin and drive (a CPU twin and a driver for host tests), acclog (acceptance logs).
- Tests against the reference engine's traces: five goldens in draft/fixtures,
  plus the chain bit for bit, calibration files, lanes, branches, lookup and speculation; test-dsv41-draft runs them in
  test-dsv41 and `zig build test`.
…es and chain on the device, several slots' statics.

- dspark_emit: DSpark's drafting pass (ingest, then the pass) as calls on the block emitter; dspark_layout: its statics,
  ring rows and the overlapped ingest's scratch roles; dspark_rows: several slots' grouping, pass statics, plan-link
  messages and candidate order (TF_DSV41_SLOT_DRAFTS).
- dspark_gpu and dspark_gpu_device: draft/iface.zig's `Pass` over those launches on every rank, one slot;
  dspark_dev: candidates on the GPU in CUDA graphs (TF_DSV41_DRAFT_GRAPHS), self-checked against the host.
- Tests: slot_drafts_test, slot_drafts_chain_test and slot_drafts_fixture (one to three slot passes equal the one-slot
  pass; the overlapped ingest, the device chain and the batched main_proj on the real emitter) in
  test-dsv41-slot-drafts, part of test-dsv41 and `zig build test`; the family's host suite compiles the DSpark files.
…t, DSpark over slots and sessions on the GPU.

- batch with its graph, profile and target parts: several slots' rows in one forward call, the round's CUDA graphs,
  per-phase timing and the lanes target's batched windows; batch_engram: each slot's staged Engram rows and the warm
  jobs (op 56), reaching the Engram gate only through forward_hooks.Gate (unset until the gate is installed).
- slots and sessions_gpu: the slots' rows and saved sessions on the GPU; target: the lanes target over the forward;
  rounds: a round's admit and ids on the plan link; lockstep: the ranks' lockstep checks; round_graph: the round's
  launches (TF_DSV41_ROUND_GRAPH).
- dspark_slots with its device and ingest parts: DSpark passes over several live slots of a batch.
- gate_lanes: the lanes engine over the GPU target on rank 0, for a GPU gate (tokens equal serial greedy).
- Tests: CED's decoder pass, the index bound and the epilogue loads on the emitted calls (ced_test, ibound_test,
  x3ld_epi_test), and the rounds, lockstep and round graph tests, in the family's host suite; it also compiles every
  batched-rows file.
…ol calls and the server's family text over HTTP.

- serve/: the release tokenizer (tokenizer.json's pipeline, byte-level BPE, a streaming detokenizer), the chat
  encoding (template.zig: effort budgets, thinking, tools, response-format schema, image placeholder text) and
  DSML tool calls (scan, typed parameters, a streaming parser); golden_test checks them against committed goldens
  and, with TF_DSV41_MODEL / TF_DSV41_GOLDEN, against the reference tokenizer and generator output.
- server/deepseek.zig: the family's model_text.Text and family.Family (thinking and effort fields, tools, the
  encoding, a streaming reply reader, reasoning memory across turns); deepseek_session.zig holds one reply's state.
  main.zig loads it in place of the chat template when config.json's model_type is deepseek_v41.
- The server reads json.zig as a module so the family's serving code shares its values: every server file imports
  `json`, json.zig exports wtf8Encode, and each server module (both native servers, the server tests, tool_parse_test,
  zig/tests/server/build.sh and zig_test.sh) gets `json` and `dsv41_serve`; json's own tests run beside them.
  Linux's test-cuda-host now also compiles tensorfold-native (main.zig), which only `zig build native` built before.
  zig/build/server.zig hosts the serving module, `test-dsv41-serve` (in `test`) and `tf-dsv41-tokbench`.
- The core tokenizer exports its UTF-8 helpers (`pub const unicode`) for the family's encoder.
- Tests: deepseek_test drives the family over real sockets with a scripted engine; deepseek_wire_test checks what
  the family's request fields give the engine and the reply (efforts, template kwargs, sampling defaults and min_p,
  drafts, background, ignore_eos, tools, stops, limits, model aliases, completions and /tokenize).
- tokbench.zig. The goldens' generators are the family's reference tools (tools/tensorfold-dsv41/ in
  github.com/jayleaton/deepseek-v41-tensorfold-spark).

Images are served as a placeholder notice (TF_DSV41_IMAGES=placeholder or reject); structured output stays refused.
…ehind the server, round planning and tf-dsv41.

- model.zig: one rank's served model at boot (weights, forms, the buffer plan, rope, Engram, the KV pool and sessions,
  DSpark, batched rows, the keyed sampler) with the forward's hooks (engine_hooks.zig) and the served knobs
  (serve_knobs_apply.zig); boot_mem.zig: boot's memory figures; aot_trace.zig: TF_DSV41_AOT_TRACE attaches an AOT set's
  lookup trace, which the set owns (attached once the set's cleanup is registered; a failed attach leaves the set as
  it was).
- serve_engine.zig: the engine the native server opens: lanes over the GPU target with DSpark on rank 0; sched_host.zig
  plans the lane host's rounds through engine_api's Rounds (TF_DSV41_PREFILL_PIECES=1); calib_gpu.zig and
  phases_timed.zig: calibration and timed phases. tensorfold-dsv41 is the native server with this engine.
- tf-dsv41 (tool_main.zig): capture replay and block-forward checks (gate_replay, gate_layers), host checks (check,
  plan, kv-emit, aot-needs), GPU checks (forward-digests, kv, slots, session-slots, slot-drafts) and the served engine
  without HTTP (follow, generate; serve_cli.zig).
- Tests: the family's host suite compiles every file here; aot_needs_test and aot_trace's ownership tests run in it;
  tf-dsv41 and tensorfold-dsv41 compile on any host in `test`.

The branches side stream, the Engram gate, L2 prefetch, the profile, the device arena, images and structured output
are not in this engine yet (their knobs are not read); the decode TP check and the image subcommands are not in the
tool.
… gate, L2 prefetch, profile, device arena.

- branches.zig: TF_DSV41_BRANCHES / TF_DSV41_MHC_DEFER, a decode window's attention sublayer and mHC coefs on side
  streams, installed with the served knobs (serve_knobs_apply.zig: TF_DSV41_PF_TBO / PF_OVERLAP now find their side
  stream) and read by aot_needs.zig for the variants they launch.
- engram_gate.zig, engram_gate_fwd.zig: TF_DSV41_ENGRAM_GATE, the GPU waits for a window's Engram rows (same bits as
  the host wait); engine_hooks.zig installs its forward_hooks.Gate and the branches' table, and the served engine's
  Engram stats report the gate's rounds, syncs, warm jobs and failures.
- l2pf.zig: TF_DSV41_L2PF, the decode windows' L2 prefetch tables filled at boot; profile.zig: TF_DSV41_PROFILE, the
  eager windows' calls timed and written as JSON at close; dev_arena.zig: TF_DSV41_ARENA's device sub-allocator and
  TF_CUDA_SCHED. model.zig boots and closes each of them.
- tf-dsv41 decode-tp PACK REF OUT (gate_decode_tp.zig): a rank of TP=2 over NCCL against tools/zig/dsv41_decode_ref.py;
  knobs TF_DSV41_DECODE_TP_{LANES,PREFILL,TOPP,PROF_ROWS}.
- Tests on the real emitters: test-dsv41-rows (row mode, branches, KV split early), test-dsv41-wide (windows past 16
  rows, deferred coefs, mHC prefill-decode, the device arena) and test-dsv41-knobs (the serving knobs, 4K prefill
  segments, prefill exchange overlap and two-batch overlap), each in test-dsv41 and `zig build test`; the family's
  host suite compiles every file above.

The kernel sources these bind (engram_gate.cu, l2pace.cu, l2pf.cu, lin2.cu) and the family's device images come next;
until then the bindings stay host-only, as before. Images are not in this engine (their boot and the decode TP check's
image prompts come with the vision layer).
… TF_TP_ROCE_HCA nor NCCL_IB_HCA is set

- `tf-tp-test hca` (host only, no GPU): both ranks run the RoCE pick over the bootstrap, the same rule the wire uses,
  and each prints its own devices as `NCCL_IB_HCA=name,...` (no GIDs). It exits 1 with the reason when the host has
  no usable port or no port on a /24 the peer shares, so no rank falls back to its first ports for NCCL.
- run_hosts.sh runs it when neither variable is set, exports the names for every step and logs each rank's pick;
  a failed or malformed `hca` stops the runner (exit 3) with the reason; `hca` is capped at TF_TP_PLAN_TIMEOUT + 31 s, so
  its own reason (no peer joined in time, port in use, the peer closed, a bootstrap error) comes first. An explicit NCCL_IB_HCA, an explicit
  TF_TP_ROCE_HCA (NCCL_IB_HCA derived from it) or both keep their behaviour; the many-devices warning is gone.
  Both hosts use the same mode: with one naming its ports and the other not, the latter's `hca` finds no peer and
  stops after TF_TP_PLAN_TIMEOUT with that reason (README).
- roce_pick: `select` / `selectFrom` return how the devices were chosen (`resolve` is unchanged for the wire),
  `joinNames` and `hcaLine` give the NCCL list. Host tests: the port scan over a stub verbs table (state, Ethernet,
  the IPv4 RoCE v2 GID index), then both ranks through a real bootstrap on the two hosts ashhart#636's receipt reports
  (modelled, documentation addresses), an explicit rank, and a host without ports.
…idency budget, streamed NVMe tier.

zig/src/sessions keeps a chat's KV pages after its request ends: entries are found by a trie over chained page
digests, held in RAM within a residency budget (chats and pool bytes), and the least recently used idle chat is
parked to an NVMe tier and streamed back on resume. Ranks that split KV run every store call in the same order, so
their books stay identical. It complements core/prompt_cache, which snapshots a whole family state; a family that
uses it reports the resident bytes as Memory.cache and its budget as Info.prompt_cache_plan. No family uses it yet;
zig build test-sessions runs the host tests, the NVMe tier on real files.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants