DeepSeek-V4.1-Flash on the Zig engine (CUDA, 2 and 4 GB10): our port is yours to use, plus test time

@ashhart

Thank you for TensorFold. Moving our DeepSeek-V4.1-Flash work onto your Zig engine was the best decision we made this month: on the same two DGX Sparks, the Zig build decodes about twice as fast after a 128K prompt as our Python build did, with the same replies. That is your engine's work more than ours.

In #286 and #299 you said DeepSeek-V4.1 is on the Zig port list, that the design isn't decided, and that you rewrite contributions in Zig yourself. This issue doesn't ask you to merge anything. It offers what we have, in case any of it saves you time, and offers our hardware for receipts.

What exists (public, Apache-2.0, all on v1.0.2)

Two nodes: branch deepseek-v41-zig, recipe v0.6.1.
Four nodes: branch deepseek-v41-zig-tp4, recipe v0.1.1.
Checkpoint: Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw, served as is, with the Engram tables from DeepSeek's original shards.
One family, zig/src/families/dsv41, on the shared lane host and server: CSA2 attention with the indexer, mHC, EXL3 experts, Engram, DSpark drafts, images, up to 1,048,576 tokens a request.

What we checked (receipts in the recipes' results/)

Drafted equals plain, concurrent equals solo, resumed equals fresh; replies equal our Python reference on every gate.
2,124 prompt lengths with no failure; images 9/9 alone and 18/18 four at a time.
Four nodes: four 500K prompts at once equal to each alone, needles 4/4.

Where it does not yet meet CONTRIBUTING.md (so you know before you look)

The lean checker adds problems over the v1.0.2 baseline (512 → 1,420), and nine files are over 600 lines.
The CUDA kernels come from a kit captured from our Python line (Triton binaries plus EXL3 cubins), not from source in the build. The kit and its builders are public with a MANIFEST, but this is the part most likely to differ from how you want it.
It is on v1.0.2; zig-flashnext is 25 commits ahead.
Our performance receipts are in our own format, not yet bench_openai.py / bench_concurrent.py with the 2k/8k/32k/64k token arrays.

What we can do, whichever design you choose

Rerun anything in your format on 2×GB10 and 4×GB10, against v1.0.2 and vLLM on the same weights, and post the receipts here.
Split the shared-engine changes (NCCL, the AOT loader, lane host, server parsing; about 1,100 lines) into small pull requests against zig-flashnext, one change each, if you want any of them, and clean them to the lean checker first.
Take any part as a reference or fixture and drop the rest; credit as you see fit.

Attribution and provenance

The engine is yours: upstream TensorFold v1.0.2, with its LICENSE and NOTICE kept, and the fork's changed files listed in NOTICE. The model math was re-implemented from DeepSeek's MIT inference code and tech report. The weights are Mia-AiLab's EXL3 quant on turboderp's ExLlamaV3 format. The port was written without opening #300 or the other DeepSeek-V4.1 ports, so whatever you take from those stays theirs. tools/dsv41-zig/ATTRIBUTION.md lists every source and license.

If none of this fits your plan, that is fine too. Thank you again for the engine.

Albert Lee (@bertholomus)

@bertholomus

Thank you, @bertholomus, for offering the dsv41 family and test time on two and four GB10s, and for listing its gaps up front. We'd like DeepSeek-V4.1-Flash on the native engine, and this issue is now its tracker (#299 and #479 point here). Three things shape how it can land:

Small pieces first: the shared-engine parts (the collective layer, the AOT loader, the lane-host and server parsing changes) as separate pull requests, each passing the lean check (one-line comments, files under about 600 lines).
Our own kernels: since 1.0.3 Nemotron's CUDA kernels are our own rather than a captured set, and a DeepSeek family follows that, so kernels captured from a Python line can't land as they are.
Receipts on your GB10s for each step: drafted equal to "draft": false, a resumed prompt equal to a fresh one, concurrent streams equal to their solo runs, and prefill and decode against the current release.

Open the first shared piece as a pull request whenever you're ready, and we'll review it promptly.

@bertholomus

Reproduction on a second four-GB10 cluster (@bertholomus's TP4 recipe v0.1.1), with numbers against SGLang on the same nodes

We moved our DeepSeek-V4.1-Flash work from the Python engine (our tuning log is #479) onto the Zig port described here. Thank you @bertholomus for publishing the recipe, the kernel kit and the receipts, and @ashhart for the engine. It came up on our hardware in an afternoon.

Setup. Two DGX Spark plus two ASUS Ascent GX10 (all GB10), one rank per node, RoCE through a switch with PFC. Engine 5940208 built with Zig 0.17.0 in nvcr.io/nvidia/pytorch:26.07-py3 for sm_121. zig build test: 212 of 213 pass; the one failure is the direct-I/O test inside a container, which upstream fixed in 1.0.3 (#486). Checkpoint: Mia-AiLab/DeepSeek-V4.1-Flash-EXL3-2.9bpw, with Engram from DeepSeek's original shards 47 and 48.

node_inputs.sh: all four weight files and the token map match results/inputs-5940208/inputs.sha256 byte for byte (5 of 5).
kit/verify_kit.sh: 356 of 356 files match MANIFEST.
tools/gates.py: 7 of 7, including drafted == serial.

One change we needed. The RDMA rings' source GID index is hard-coded to 5 (model.zig in two places, layer_main.zig in one). On our ConnectX ports the RoCE v2 IPv4 entry is index 3, so we made it a setting that keeps your default:

/// TF_RDMA_GID: the RDMA rings' source GID index (the port's RoCE v2 IPv4 entry); 5 as served on the authors' nodes.
pub fn rdmaGid() u8 {
    const v = std.c.getenv("TF_RDMA_GID") orelse return 5;
    return std.fmt.parseInt(u8, std.mem.span(v), 10) catch 5;
}
// ... .gid_index = rdmaGid() ...

Together with NCCL_IB_GID_INDEX=3 NCCL_IB_TC=106 for our fabric, everything else ran as config.env.example ships.

Numbers on our nodes.

All rows: greedy, thinking off.
Production SGLang is MiaAI's TP4 production line (MiaAI-Lab/DeepSeek-v4.1-Flash-DGX-Sparks, every performance flag of docs/tp4.md). On our fabric it runs about 7-17% below MiaAI's published figures, so read SGLang's columns as our SGLang deployment.
Decode and prefill are MiaAI's sparkDash benchmarks run headless, the same protocol for both engines. Its prose and structured streams share one prompt; its code streams are distinct.

aggregate tok/s	c1	c2	c4	c8
prose, Zig / SGLang	112 / 82	156 / 114	208 / 150	250 / 204
code, Zig / SGLang	168 / 126	218 / 176	249 / 227	277 / 274
structured, Zig / SGLang	208 / 150	233 / 148	305 / 209	366 / 243

prefill tok/s	4K	8K	16K	32K	64K	128K
Zig	1,795	1,820	1,813	1,816	1,807	1,762
SGLang	2,820	4,628	4,581	4,790	4,968	4,218

Decode is a clear step up: 1.2-1.6x SGLang, and about 1.7x our Python build at c1. The kit_bench.py decode --set b structured row came out at 210 tok/s, matching the recipe's 219.5.

Prefill is flat at about 1.8K tok/s, below the recipe's 2,778-2,900. We haven't yet found where the difference comes from (prompt content, fabric or settings). On our real-traffic replay (multi-turn agent sessions, cold prompts up to about 160K) that makes the build prefill-bound: c2 148 s vs SGLang 88 s, c4 476 s vs 238 s (SGLang replay from 10-06).

One failure. Two real-traffic cold prompts, of 146,421 and 145,152 tokens, ended after about 73 s:

error: no captured Triton variant of _index_score for this launch:
error:   ... n: i32 146421 ... cs: i32 0 ... nt: i32 2288
[tensorfold] ended chatcmpl-… reason=error (MissingTritonVariant) prompt=146421 tokens=0 after=73.62s

Your 160K and 500K checks passed, so it looks like a specialization our prompt lengths reach and the kit doesn't cover. We are tracing which key it needs and will post what we find.

What we're working on, on this build: that variant first, then prefill throughput, since that's what real traffic waits on here. From the Python side we have measured, exact prefill levers in #479 that may transfer, such as sequence-parallel rows and overlapping the row all-to-alls. We're happy to rerun anything in your formats (bench_openai.py / bench_concurrent.py, or the recipe's tools) on these four nodes, against v1.0.2 or a newer base.

Follow-up: the _index_score MissingTritonVariant at ~145K, root cause and an exact fix

Both failures trace to the same call. After the replay cut at P − 128, layers 20-39 score the last 128 prompt rows against the full prompt length in row blocks of rb = max(16, min(rows, 2^26 / (4·n))). Blocks of more than 16 rows go to _index_score_rows, which the kit covers well (102 variants). A trailing block of 16 rows or fewer goes to the fp32 _index_score (KEYS=0 HAS_BASE=0 HAS_CAND=0 HAS_TMAX=0).

That happens for prompts of 131,073-149,796 tokens (rb = 127…112), and again in narrower windows further out, such as 262,145-299,593. Your 160K and 500K checks fall between those windows, which is why they passed.

In that case the launch can need specializations the kit never captured:

146,421 tokens: nt divisible by 16 (n not).
145,152 tokens: n divisible by 16, and VIS/KB/CAND 8 bytes off 16-byte alignment. This happens when rb is odd: rb = 115 gives r0·8 ≡ 8 mod 16.

Roughly half the prompts in that window fail at the end of prefill.

Why a fallback is exact here. For these settings the kit holds three variants with HAS_BASE=0 and no int fixed to 1 in a way that matters:

variant	n	nt	VIS
`5f4fceb20c`	any	any	16-aligned
`670fc9ffa0`	div 16	1	16-aligned
`a0120ddd24`	any	1	16-aligned

All three have identical SASS (cuobjdump -sass, addresses stripped). The kernel never reads nt, cs, KB, CAND or TMAX. It reads VIS with one 8-byte load, so 8-byte alignment is enough; the KEYS=1 pair 9f83460754/45c81d5893 differs only in VIS/KB/CAND alignment and is also SASS-identical.

So we changed aot.zig's Set.find:

It still tries the exact match first, so every launch that works today is unchanged.
Only if that fails does it accept a variant compiled for "any" where the launch's int is divisible by 16, or for 16-byte alignment where the pointer is 8-aligned.
tri_index.indexScore allows this only for n, nt, VIS, KB and CAND, and only on the fp32 FP4 path (!keys and base == null).
It logs once which variant it took.

The diff is 3 files, +88/−7: zig/src/cuda/aot.zig, families/dsv41/tri.zig and tri_index.zig, with a unit test.

Verified on our four nodes:

Before the fix, prompts of exactly 146,421 and 145,152 tokens failed with HTTP 500 MissingTritonVariant every time.
With the fix, both complete.
A 156,941-token prompt, which takes the exact-match path, gives the identical reply on both builds.
gates.py 7/7.
The server logs no exact captured variant of _index_score: took 5f4fceb20c… (a specialization its code does not depend on).

Not fixed: the keyed layers 2-19 (KEYS=1 HAS_TMAX=1) can hit the same gap on the last chunk's tail. There nt is read, and the SASS differs after the first ~1,250 instructions, so we didn't relax it; a recapture of those variants would be the right fix.

We're happy to send the diff as a PR against deepseek-v41-zig-tp4 if that helps.

Update 2: prefill on the TP4 2D split — first exact gain (+12-17%), and what limits it

Same four nodes as above, engine 5940208 + TF_RDMA_GID + the _index_score fallback.

Why prefill is ~2K here and flat from 4K to 128K.

Flat: after the replay cut at P − 128, the decoder layers run on only about the last 128 rows, so a prompt has almost no fixed cost. Everything is per row, about 0.5 ms of cluster time per row.
~2K: the 2D split exchanges about 1.8 MB per prompt row per node: two fp32 partial all-gathers per layer, plus the wo_a exchange, which gathers from all four nodes though only the column partner's part is used. All of it is on the compute stream, serial with the GEMMs.
What we ruled out:
- NCCL tuning (channels, buffer size, protocol): ±1%.
- SM clocks: steady at about 2.18 GHz on all four nodes during a 32K prefill.
- Rails: both are in use.
kit_bench.py prefill gives about 2.05K here against the recipe's 2.78-2.90K, so part of the gap is likely our fabric (100G per node).

Three opt-in levers, each with the default left byte-identical. Each arm passed gates.py 7/7 and gave identical replies to the base build on four short prompts (4,224 / 8,192 / 12,345 / 33,000 tokens) and two long ones (156,941 / 146,421), before any timing.

arm	prefill 8K / 32K / 128K (tok/s, `kit_bench`)	156,941-token prompt
base	2,068 / 2,070 / 2,019	58.4 s
`TF_DS_PROMPT_CHUNK=4096`: two 2,048-row chunks per encoder pass; row-count-dependent steps (cuBLAS SGEMMs, indexer selection) still per chunk	2,147 / 2,110 / 2,042	62.7 s
`TF_DS_PROMPT_SP=rows`: sequence-parallel rows on the 2D split	2,324 / 2,432 / 2,326	46.0 s
both	2,338 / 2,286 / 2,264	51.8 s

SP rows, in detail:
- Each node receives only the partials of the rows it owns, through NCCL send/recv groups in the same layout and rank order as today's gather.
- It runs the residual/mHC mixes, the Engram gate and the taps on those rows only, then all-gathers bf16 x.
- wo_a becomes a send/recv with the partner node.
- That is about 47% fewer bytes per row per layer, and the sums keep the same operands and order.
- It applies to encoder-only chunks with n % 64 == 0; head and decoder chunks run as today.
Decode is unchanged in every arm: about 150 / 90 / 210 tok/s for code / prose / structured on set b.

A third lever hung. Overlapping two 2,048-row sub-chunks one layer apart, on a second stream and a second NCCL communicator, deadlocked silently on the first overlapped pass.

Gates passed because chat prompts never reach 4,096 rows.
Cause: NCCL's multiple-communicator launch-order hazard. Chunk A's next-layer collectives on comm 1 and chunk B's on comm 2 became ready together, and different ranks started them in different orders.
We now chain both comms' launches through one event in host issue order, which is identical on every rank, plus per-pass log marks.
That fix is built and under test now. NCCL_LAUNCH_ORDER_IMPLICIT is the env-only cross-check if it doesn't hold.

What actually decides real traffic.

Our replay of the real-traffic mix is calibrated to the measured runs within 1%. It says end-to-end time on this mix is set by prompt rate alone.
The engine runs either a prompt chunk or a decode round, so a closed loop's wall time is total prompt time plus total round time, in any order.
Scheduling therefore helps latency, not throughput. For example, admitting short prompts first cuts c4 median TTFT from about 25 s to about 4 s at the same wall time.
The engine passes SGLang's c2 / c4 wall on our nodes once prefill reaches about 4K tok/s (about 2× today). That is where we're working.

Code is on local branches. We're happy to open PRs against deepseek-v41-zig-tp4 for any of it.

Update 3: SP rows on real traffic, and the overlap result

SP rows (TF_DS_PROMPT_SP=rows) as the served build.

Exactness: gates.py 7/7, and the long-prompt replies are identical to the base build.
Real-traffic replay (multi-turn agent sessions, cold prompts up to about 160K, the same draw as before):

	c1	c2	c4
base + `_index_score` fallback	252.5 s (2 errors before the fallback)	148.0 s	476.3 s
+ SP rows	245.4 s, 0 errors	133.4 s (−10%)	423.0 s (−11%)
our SGLang deployment	162 s	88 s	238 s

sparkDash prefill, 8K / 32K / 64K / 128K: 1,962 / 2,089 / 2,241 / 2,003 tok/s, up from 1,820 / 1,816 / 1,807 / 1,762.
Decode and our distinct-prompt harness: unchanged (code c1 166 tok/s, prose c1 82).

The overlap lever: the hang is fixed, but it doesn't pay.

Ordering both communicators' collectives through one event in host issue order removed the deadlock: 142 overlapped passes completed on all four ranks, and replies were identical to the base on all six test prompts.
But it's slower than SP rows alone: 2.14-2.30K tok/s against 2.14-2.43K. Without SP rows it's 1.93-2.02K against 2.04-2.15K.
Once the two comms' collectives can't run concurrently, the compute overlap left over doesn't cover the extra stream and synchronization cost.
We've dropped it. NCCL_LAUNCH_ORDER_IMPLICIT would give the same ordering guarantee.

Next.

Short-first admission (rank 0 only, exact by construction) is on the cluster now, for TTFT.
For throughput, prompt rate is still the whole story on this mix. What's left there is mostly the EXL3 prompt GEMMs, plus taking SP rows to the head chunk.

Update 4: short-first admission, and a pause

Short-first admission is inconclusive. We tried admitting queued prompts of ≤32K uncached rows before longer ones. It's rank-0 only, exact by construction, and gates 7/7. We ran it as two order-swapped pairs on the real-traffic replay:

	c4 wall	c4 mean TTFT, requests < 20K uncached
FIFO, run 1 / run 2	425.4 / 423.1 s	18.9 / 10.0 s
short-first, run 1 / run 2	420.1 / 422.8 s	9.8 / 9.3 s

The second FIFO run matched short-first, so the first pair's gap was arrival-order noise. Wall time is unchanged, as expected, since this mix is bound by prompt rate. We've left it opt-in.

Where this stands. The served build is the recipe's engine plus three changes:

the TF_RDMA_GID setting;
the exact _index_score fallback;
TF_DS_PROMPT_SP=rows.

On our replay that build runs c1/c2/c4 in 245 / 133 / 423 s with 0 errors. Its decode stays 1.2-1.6× our SGLang deployment's. Prompt rate is still the gap.

We're pausing to move the cluster. Next, after the move:

SP rows for the replay/head chunk;
a per-phase prompt profile, to show how much of the remaining ~0.43 ms per row is EXL3 prompt GEMMs and how much is collectives.

We'll post that here.

Update 5: back after moving the cluster. SP rows for the head chunk, and a prompt phase profile

The cluster came back from the move with both RoCE rails at 100G on every node and the same GID table, so nothing changed on the engine side.

TF_DS_PROMPT_SP=rows2 extends SP rows to the chunk that holds the replay point, and to row counts that aren't a multiple of 64. Nodes 0-2 own 16·⌊n/64⌋ rows each and node 3 owns the rest; at layer 20 the replay rows go back to every node.

Exactness: gates.py 7/7. Replies are identical to the base build on 12 prompts, including new odd lengths chosen to put the cut on owner boundaries (300, 400, 1,500, 2,100, 2,200 and 6,000 tokens).
Speed: 8K is +2-4% (2,380-2,416 vs 2,321-2,330 tok/s), 32K +1%, 128K unchanged. Decode is unchanged.

Where a prompt row's time goes now. We added an opt-in per-chunk phase profiler (TF_DS_PROMPT_PROF=1). It records CUDA events per phase on every rank and syncs once per chunk. These are encoder chunks with SP rows, on all four ranks; the ranks agree within 1%.

phase	µs per prompt row	share
collectives (wait + transfer)	120	34%
EXL3 routed experts	106	30%
attention + indexer	52	15%
dense (EXL3 dense + cuBLAS)	42	12%
hc mixes / norms	12	4%
router	11	3%
Engram	5.5	2%
total	~351

The head chunk (2,048 rows, replay cut 128) costs about 1.03-1.07 s, of which the decoder layers on the last 128 rows are 114-139 ms.

Next. The two big items are the collectives and the EXL3 expert GEMMs.

On our Python line, a faster EXL3 prompt path, from tuned tiles and fewer launches, gave about +70% on that stage.
We're porting that idea to the Zig kernels first.
Then we'll see whether the remaining prompt-size exchanges can move to the RDMA rings, which today only take decode-size messages up to the 4.25 MiB slot.

Update 6: faster EXL3 prompt GEMMs, exact, +5% prefill

TF_DS_EXL3_PROMPT=v2 is a new prompt-chunk kernel for a 2D node's routed-expert gate/up and down GEMMs.

One program covers 32 rows × 256 columns with 16 warps. Each warp decodes its trellis tile once, straight into mma registers, through a three-stage cp.async ring of 128-column Hadamard steps.
The gate/up input rotation runs inside the program, so the fp16 rotated rows are never written to or read back from memory. That round trip was 294 MB per layer at 2,048 rows.
The shared expert runs in the same launch.
The dense prompt GEMMs' weights are unpacked once at load instead of every chunk, which costs about 1.7 GB per rank (TF_DS_EXL3_WQ_CACHE, on with v2).
Decode-size windows and rounds keep today's path.

Exactness.

The trellis decode is copied verbatim. The k-tile order and the per-element sum order are the served kernels' own, and the rotation's first butterfly stage is pinned to _rn ops to match the shipped SASS.
A standalone bench loads real weights with the engine's own loader and runs both paths on identical inputs: 1,680 cases, all byte-equal on every intermediate. The cases cover 2 ranks, layers 0-19, row counts from 64 to 4,096, and embedding, Gaussian and edge-value inputs (fp16 overflow, ±0, bf16 subnormals).
On the cluster: gates.py 7/7, and replies identical to the previous build on all 12 test prompts (300 to 156,941 tokens).

Speed (kit_bench prefill, two reps each; decode unchanged at 154 / 93 / 214 tok/s for code / prose / structured):

	8K	32K	128K
rows2 (update 5)	2,410-2,429	2,401-2,459	2,347
+ EXL3 v2	2,547-2,554	2,490-2,571	2,451-2,462

Phase profile (encoder rows, µs per row, all four ranks):

phase	before	after
EXL3 routed experts	106	89-92
dense	42	37
collectives	120	117-124
total	~351	~332

Collectives are now 35-37% of a prompt row and the largest phase by a wide margin. They're the next target: moving prompt-size exchanges onto the RDMA rings, or sending fewer bytes. We'll also re-try two-chunk encoder passes (TF_DS_PROMPT_CHUNK=4096), since v2's experts get cheaper per row at 4,096 rows (89 vs 107 µs on the bench).

Update 7: what the prompt collectives are made of, two-chunk passes revisited, and an exact Engram cut (+2-8% prefill)

Where the comm bytes go. In an encoder-only chunk under SP rows (rows2) and the v2 EXL3 GEMMs, each node receives about 1,029 KB per prompt row, in 124 NCCL collectives per chunk. All of them run on the compute stream, and every one is larger than the RDMA rings' 4.25 MiB slot.

per layer	dtype	B/row
x all-gather after the attention mixes	bf16	7,680
`wo_a` exchange with the column partner	bf16	4,096
attention partials to the row owners	fp32	7,680
x all-gather after the FFN mixes	bf16	7,680
expert pack exchange (parity split)	fp16	8,960
MoE partials to the row owners	fp32	7,680
per layer (×20)		43,776
Engram layers 1 and 14 (two pair steps each)	fp32 / bf16	76,800 each

Our two RoCE rails are two PCIe paths onto one 100G port per node, so the wire floor is about 82 µs/row.
The profiler measures 117-124 µs/row, so NCCL runs at about 66-71% of line rate.
Apart from Engram, these bytes are close to the minimum for this split. The attention projections, router and experts run on all rows on every node, so both x gathers are needed.

Two exact, opt-in comm levers. Each passed gates.py 7/7 and gave replies identical to the reference on all 12 test prompts (48/48 compares).

TF_DS_PROMPT_ENGRAM=sp: in an SP chunk the Engram gate only reads a node's own rows, so each node now sums just its owned rows plus the rows its column partner needs.
- That is 32.0 KB instead of 76.8 KB per row per Engram layer, about 9% fewer bytes overall.
- The elementwise add gets the same two fp32 operands in the same order, so the owned rows are bit-identical. Under SP nothing reads the other rows before the cut at layer 20.
- The Engram phase drops from 5.2 to 3.7 µs/row.
TF_DS_PROMPT_A2A=shift: the owner exchange as three permutation steps instead of one send/recv group, to avoid three senders per switch port at once. Exact, but 3-5% slower, so it's dropped.

Two-chunk encoder passes again (TF_DS_PROMPT_CHUNK=4096). In update 2 this was flat. With v2 the experts get cheaper per row at 4,096 rows (expert phase 90 → 85 µs/row), but comm rises (120 → 133). Net on its own: +5% at 8K, about 0 at 32K and 128K.

Combined. These are means of two independent arms each, run in one session; decode is unchanged at 152-154 / 92 / 210-212 tok/s.

`kit_bench` prefill, tok/s	8K	32K	128K
rows2 + v2 (update 6)	2,497	2,524	2,437
+ chunk 4096	2,629	2,530	2,442
+ chunk 4096 + Engram sp	2,704 (+8%)	2,609 (+3%)	2,478 (+2%)

That last row is now our served config. Since the start of this thread, prefill has gone from about 1.8K to 2.5-2.7K tok/s, with every step byte-identical to the original replies.

Next. Comm is now about 40% of a prompt row, and at about 70% of line rate on a single 100G port, wire efficiency is the remaining lever. The engine has to wait for each exchange today, so we're looking at one of two things:

hiding exchanges behind the GEMMs on our hardware (the earlier two-stream attempt was slower), or
better wire efficiency for prompt-size messages (nccl-tests against our RDMA bench at 5-18 MB first).

Update 8: comm overlapped with compute, exact, +2-3% prefill

Update 3's overlap lost because it split a pass into two sub-chunks on two streams and two communicators. TF_DS_PROMPT_OVL keeps one 4,096-row pass, one compute stream and one communicator. A side stream carries only the NCCL ops (plus two tiny Engram adds), ordered through one event chain, and no NCCL op is added. Three levers (all = all three):

engram: Engram layer L's projection depends only on token ids, so it is issued at layer L−1. Its two exchange steps run under L−1's attention and router.
xag: the attention-side x all-gather runs on the comm stream. Meanwhile wq_a, wkv, q_norm and wq_b run on this node's own rows, then on the rest.
wob: wo_b runs first on the other owners' rows. The exchange to the owners then starts on the comm stream while wo_b finishes this node's own rows.
Why it's exact: row splits launch the same kernels on row-disjoint parts. Each part keeps at least 128 rows on 16-row tiles, so the prompt GEMM and RMSNorm pick the same kernels with the same per-row K loop. The Engram path is the same plan and operand order, just on another stream.

Exactness: every arm passed gates.py 7/7 and gave replies identical to the served build on all 12 test prompts (120/120 compares, including a build-default arm).

Speed (means of two arms each, one session; decode unchanged at 153 / 92 / 211 tok/s):

`kit_bench` prefill, tok/s	8K	32K	128K
served (update 7)	2,696	2,613	2,486
+ `TF_DS_PROMPT_OVL=all`	2,754 (+2%)	2,691 (+3%)	2,532 (+2%)
`engram` alone (one arm)	2,750	2,670	2,543
`xag,wob` (one arm)	2,703	2,647	2,500
`all` + comm stream at top priority	2,647	2,588	2,527

Profile (encoder rows, µs/row, same-day pair). It shows why the gain is real but small on GB10: NCCL's kernels share the SMs, so the compute they overlap slows down.

phase	served	+ OVL=all
comm (wait on the compute stream)	129-137	97-108
dense	38-40	50-52
EXL3 experts / attention	84-86 / 52	84-86 / 52
total	339.4	324.8

About 30 µs/row of exchange is now hidden, at a cost of about 12 µs/row of slower dense GEMMs. The net is −14.5 µs/row. Most of it comes from the Engram lever. Raising the comm stream's priority made things worse.

Served now: rows2 + EXL3 v2 + chunk 4096 + Engram sp + OVL=all, at about 2.53-2.75K tok/s prefill, every step byte-identical to the original replies.

Next. We'll look at fewer NCCL CTAs (fewer SMs taken from the overlapped compute). The expert-side exchanges (pack and down) are the large ones that can't overlap yet, because their member lists cover all rows.

Update 9: checkpoint on real traffic, served build after updates 5-8

This is the same real-traffic replay as update 3: multi-turn agent sessions, cold prompts up to about 160K, the same seeded draw, thinking off. It ran on the served build: SP rows (rows2) + EXL3 v2 + chunk 4096 + Engram sp + OVL=all. gates.py was 7/7 first.

wall time	c1	c2	c4
update 3 (SP rows)	245.4 s	133.4 s	423.0 s
now	222.2 s (−9.5%)	121.0 s (−9.3%)	381.4 s (−9.8%)
our SGLang deployment	162 s	88 s	238 s

Errors: 0 at every level.
TTFT p50: 1.73 / 1.87 / 6.70 s, down from 1.94 / 2.14 / 7.60.
Prompt rate in the mix: about 2.43-2.51K tok/s at c1/c2.

Other benchmarks:

sparkDash prefill, 4K / 8K / 16K / 32K / 64K / 128K: 2,186 / 2,283 / 2,161 / 2,094 / 2,134 / 2,051 tok/s, against 1,623 / 1,962 / 2,027 / 2,089 / 2,241 / 2,003 at update 3. Its filler is one repeated token, which is a different, Engram-hot input from kit_bench's, and it gains less at 32K-128K. We report it as measured.
Decode is unchanged: sparkDash code is 168 / 113 / 68 / 38 tok/s per stream at c1 / c2 / c4 / c8. On our distinct-prompt harness, code is 160 at c1 and prose is 82.

Where this leaves us. Wall time on this mix is set by prompt rate. To reach SGLang's c2 / c4 wall on our nodes, prefill still needs roughly 1.4-1.6× more. Comm wait is down to about 100 µs/row of about 325, and work is under way on:

the expert-side exchanges, which are the large ones OVL can't hide yet;
exact speedups for the attention/indexer and dense prompt kernels.

We also tried capping NCCL's CTAs (NCCL_MAX_CTAS 2 / 4 / 8) to free SMs for the overlapped compute. It was exact but gave no gain: 8 matched the served build, and 4 and 2 were 3-6% slower. We've dropped it.

Update 10: the replicated attention work goes, plus expert-side overlap: +6.6% prefill, exact (now 2.78K tok/s at 128K)

Two more exact, opt-in levers this session, tested separately (ZX1, ZT1) and then together (ZC2). Every arm reported below passed gates.py 7/7 and gave replies identical to the served build on all 12 test prompts before any benchmark: 120 + 116 + 84 compares, 0 mismatches.

1. TF_DS_PROMPT_ATTN=front,isplit: stop computing other nodes' rows. Under SP rows each node owns a quarter of a chunk's rows, but every node still ran the attention front and the indexer for all rows and kept only its own.

front (17 of 20 encoder layers): wq_a, wkv and q_norm now run on the node's own rows only. qr and kv are all-gathered (3.5 KB/row) instead of x (10 KB/row), and wq_b runs on the own rows while that gather is in flight.
isplit (indexer layers 2, 8, 14): idx_wq_b, RoPE, the FP4 quantize, the scores and top-k run on the own rows only, in today's row blocks with today's kernel per block. The selection is then all-gathered.
What stays replicated: the indexer's cuBLAS head-weight GEMM stays on all rows, because cuBLAS picks its algorithm by M. We measured this for the router: at M = 1,024, 93% of logits differ from the M = 2,048 run.
Equality check: a one-GPU harness emulates the four owners with the production launches and compares q, keys, o, u, the wo_b partial, the selection and the caches by sha256. All 9 cases match, covering starts up to 127K, uneven owners, a 1-row block intersection, and a 2-row _index_score tail.
Speed alone: +5.4%. The attention phase halves (52.5 → 25.8 µs/row) and the dense phase drops from 51.5 to 35.
Not shipped: a third lever, a one-launch top-k kernel (tsel), passed the gates but hung on every prompt of 4K tokens or more. It is excluded and being debugged.

2. More TF_DS_PROMPT_OVL words: the expert-side exchanges.

pack3 sends the expert pack in two column halves, each under the next half's gate/up or under the rest block.
dn runs down and the combine in column halves, so half A's exchange to the owners runs under half B. Halves are whole 128-column Hadamard blocks, so the epilogue and combine read the same values for every element.
fxag replaces the FFN-side x all-gather with a row-partner then a column-partner exchange. On 4,096-row passes the pair's rows are exactly one of the router's 2,048-row cuBLAS runs, unchanged, which runs under the second exchange.
Equality check: all 160 cases per word were byte-identical on real weights.
Speed alone: +1.6%. About 10 µs/row of comm is hidden, but the overlapped expert GEMMs slow by about 4 because NCCL shares the SMs.

Combined (means of three arms each, one session: plain, repeat and profiled; decode unchanged at 152 / 91 / 211 tok/s):

`kit_bench` prefill, tok/s	8K	32K	128K
served (update 8)	2,753	2,666	2,530
+ ATTN front,isplit + OVL pack3,dn,fxag	2,780	2,784 (+4%)	2,757 (+9%)

Prefill no longer falls off with context: the indexer work that grew with context is now split four ways.

phase, µs per prompt row (encoder rows)	before	after
attention + indexer	52.4	25.9
dense	51	35
EXL3 experts	85	88
comm wait	102-111	95-116
total	327.5	291.7

This is now the served config. Since the start of this thread, encoder-row cost has gone from about 500 µs to 292 µs, with every step byte-identical to the original replies.

Next:

The router is the last per-row work still replicated on every node (about 11 µs/row). It can only move to own rows if cuBLAS runs the M = 2,048 algorithm at the smaller M, so we're testing a pinned cublasLt algorithm.
Fixing tsel (top-k is about 9 µs/row at 128K).
Comm, at about 100+ µs/row, is again the largest phase.

Update 11: two more exact levers, and the prompt path is now bound by the wire

Since update 10 we built and tested two more exact, opt-in levers. Every arm passed gates.py 7/7 and gave replies identical to the served build on all 12 test prompts: 108 + 120 compares, 0 mismatches.

TF_DS_PROMPT_ROUTER=own: the router on owned rows. The router was the last per-row work still replicated on all four nodes.

cuBLAS alone doesn't do it exactly. cuBLAS chooses its split-K by M. The [384 × 5120] gate's 2,048-row run is split-K 12 (11 × 448 + 192). At 512 rows it is split-K 16, at 4,096 it is 7, and so on, so a smaller-M cuBLAS call changes the bits. Pinning the cublasLt heuristic reproduces the bits at 2,048 and 1,024 rows, but not for every owner size.
What we built instead is a kernel that runs today's partitions, ascending FFMA chains and in-order partition adds for any row count, with explicit _rn ops. It reads bf16 x and fp16 weights directly, so the two cast launches go too.
Checked at boot. Every node calibrates the kernel against cuBLAS's own 2,048-row call on real gate weights and refuses to start if they differ, so a cuBLAS upgrade that changes the split can't silently change replies.
One-GPU check: 240/240 chunk cases and 2,880/2,880 random row windows byte-identical in logits, picks and weights.
Router phase: 11.0 → 2.4 µs/row.

The one-launch top-k (tsel) is fixed. The hang in update 10 was a radix-select pass with fewer than 32 bins that no lane claimed, so the loop never advanced. Invariant breaks now __trap() instead of hanging. One-GPU sweep: 1,311/1,311 cases equal to today's chain, from 512 to 524,288 keys. It saves about 1.5 µs/row in the profile.

On the cluster both are exact, and they don't move prefill. The table is kit_bench prefill, mean over 32K and 128K, two reps each, plus a profiled pair.

arm	plain	repeat	profiled
served (update 10)	2,786	2,801	2,799
+ router own	2,812
+ tsel	2,802
+ both	2,800	2,805	2,790

The profile shows why. Per encoder row, the summed per-phase GPU time drops from 296.1 to 283.5 µs, but wall time doesn't follow. With the exchanges now overlapped, the critical path is set by the wire on the slowest ranks (comm wait 115-122 µs/row on ranks 0/1), and compute we remove off that path just becomes more waiting. We're keeping both levers on the branch, off by default, for when comm gets faster.

Next: the wire itself. Each node has one 100G port, and prompt-size exchanges run at about 66-71% of line rate. Before changing more code, we'll measure what NCCL can actually reach on this fabric with nccl-tests: all-gather and send/recv at 1-20 MB, ring vs tree, with and without the second rail.

If NCCL itself tops out near 70%, the remaining exact lever is fewer bytes. The biggest candidates are the fp32 owner exchanges, about 15 KB/row.
If it can go higher, the gap is in how and when we issue the exchanges.

Update 12: the fabric isn't the limit; NCCL settings +2.4%; the real stall is Engram reads on two nodes

What NCCL reaches on our fabric (stage ZW1). We ran nccl-tests v2.21.1 against NCCL 2.30.7, with 4 ranks, one GB10 per node, and the engine's own NCCL env. Busbw per node per direction, at the engine's message sizes:

test	16 MB	32 MB	% of line rate at 32 MB
all_gather (Ring)	11.08	11.64 GB/s	93%
alltoall	11.32	11.60	93%
sendrecv, column pairs 0↔2 and 1↔3 concurrent	11.62	11.84	95%
all_reduce, Ring	11.65	11.89	95%

The wire runs at 93-96% of 100G line rate. Inside the engine, prompt exchanges get about 66-71%, so the gap is in the engine.
Settings to avoid on this fabric:
- PAT all_gather: 54-71%
- LL128: about 53%
- Tree all_reduce: about 20%
NCCL's defaults sometimes dip to about 73% at 32 MB. Fixed channel counts hold at 92-93%. Full table, including rails, QPs, channels, protocols and plugins: results/stageZW1/summary.md.

Those settings in the engine (ZN2). These are env-only changes on the served build, with the same gates and 12-prompt identity checks as before (all identical). Numbers are mean kit_bench prefill over 8K / 32K / 128K, two reps each.

arm	run	repeat
served (update 10)	2,837	2,810
fixed 8 channels	2,828
fixed 16 channels	2,843
2 QPs per connection	2,860
2 QPs + fixed 16 channels	2,891	2,895 (+2.4%)

Decode is unchanged. That last row is now served: NCCL_IB_QPS_PER_CONNECTION=2 NCCL_MIN_NCHANNELS=16 NCCL_MAX_NCHANNELS=16. The 157K-token long check dropped from 52 s to 45.5 s.

Where the rest of the "comm" time goes. We read the per-chunk phase profile by rank (served build, µs per encoder row):

rank	node	comm wait	host idle	sum
0	DGX Spark	118.9	2.1	121.0
1	DGX Spark	121.8	1.7	123.5
2	ASUS GX10	100.8	15.0	115.8
3	ASUS GX10	89.9	30.9	120.8

The sum is nearly constant across ranks. The two GX10 nodes idle while the host reads Engram rows from disk; the two tables are 2 × 101.5 GB, which don't fit in page cache. Ranks 0/1's extra comm wait tracks that per chunk (r = 0.88).
Why their reads are slower: they have a smaller, slower 1 TB NVMe, and today's read-ahead for the next chunk starts only after the chunk's last Engram layer.
This also explains update 11: compute we removed came off a path that was waiting on those reads.

Next.

An exact fix that only changes when bytes are read: double-buffered read-ahead started at layer 0, or the whole prompt's Engram rows pipelined several chunks ahead. It's being built now.
A per-collective profiler and a 4-rank replay of the engine's exact NCCL sequence are running now. They'll split the remaining gap into wire time, rank skew and contention.

Update 13: the comm gap, measured per collective

We added an opt-in per-collective profiler (TF_DS_PROMPT_PROF=2). It records a ready event and an end event around every NCCL call on all four ranks, puts the ranks on one clock, and joins them offline. We also added a replay tool that runs the engine's exact NCCL sequence (225 ops per 4,096-row chunk) on the four nodes, either bare, with the recorded arrival skew, or under a synthetic memory/SM load. Replies were still identical with profiling on (gates 7/7, 12/12).

Per chunk row, rank 0, served build (µs; "ideal" is the same bytes at the nccl-tests rates from update 12):

	skew (waiting for peers)	wire	wire GB/s	exposed (not overlapped)
engine	44.5	131.6	6.6 (58% of nccl-tests)	122.7
replay, bare (no compute)	12.1	107.9	8.1 (71%)
ideal		76.6

So the about 99 µs/row over ideal comes from three things:

Skew, about 44 µs. Ranks 0/1 wait for the GX10s, mostly on their Engram disk reads (update 12).
The sequence itself, about 31 µs. Even bare, our op mix runs at 71%.
Contention with overlapped compute, about 24 µs.

By collective (wire µs per row, engine vs bare replay):

collective	engine	bare
owner exchange, 3-peer send/recv group	46.5 (57%)	39.2 (68%)
expert pack exchange, column partner (fp16)	38.1 (41%)	17.3 (90%)
row-partner bf16 exchange	11.9 (38%)	13.1 (34%)
column-partner bf16 exchange	11.2 (78%)	11.6 (75%)
everything else	~24	~27

What we'll do about each:

Expert pack exchange. It's the one hurt by overlap: it runs at 90% bare but 41% under the overlapped expert GEMMs (pack3). We'll re-profile with update 12's NCCL settings, then try running it un-overlapped or with fewer channels.
Owner exchange. It's inefficient even alone, because all three peers send into one port at once. Our earlier three-step shift version was slower. Next we'll try merging it with the neighbouring exchange into one call.
Row-partner bf16 exchange. It runs at 34-38% even bare, so it's a small-message or layout problem to look at.
Skew. We're building an earlier, deeper Engram read-ahead now. It is exact by construction, because only when bytes are read changes.

Update 19: a soak found latent "missing Triton variant" gaps that hang TP4; fixed, and the soak now passes

We ran a 1 h soak of the update-17 build: c3, 60 multi-turn agent sessions, prompts up to 160K. The cluster hung after 12 minutes.

Gap 1: _index_score.

Where: the indexer cuts each 2,048-row run into row blocks of 2^24 / n_comp_end. When that is odd (n = 80,809 gives 207), the trailing block of 16 rows or fewer starts at an odd row, so VIS = vis + row*8 is 8-aligned and not 16-aligned.
Why no variant matched: the keyed variant with tile maxima (any n/nt) assumes 16-aligned VIS/KB/CAND. Our earlier ptrs8 relaxation only covered the non-keyed path.
Reach: about 2,300 prompt lengths up to 200K, the first at 18,226. It is not a regression from the upstream sync: our pre-sync build refused about 3,100.
Why it hung: under TF_DS_PROMPT_ATTN=isplit the tail rows belong only to node 3. That follower refused the frame, rank 0 didn't, and ranks 0-2 waited in NCCL indefinitely.

Gap 2: _sparse_attn_part.

Where: when the replay point falls in the last 16 rows of a chunk, the replay cut leaves a tail of 16 rows or fewer. If the cut's first row is odd, that tail runs ring-mode _sparse_attn_part with 8-aligned POS/RBASE/CBASE.
Reach: prompt lengths with len % 2048 odd and between 113 and 127, which is 776 lengths up to 200K, the first at 2,161.
Effect: every rank runs that tail, so every rank refused.

Fix (homelab/night5, edf22ac + 79e9c25).

The selection rule: an int64 pointer may be 8-aligned in a variant compiled for 16 only when no global access through it is wider than 8 bytes.
- A tool follows each *i64 parameter's registers through the SASS of every captured cubin.
- It is flow-sensitive across branches and spills, and makes no claim through indirect branches.
- It emits a table that the launcher consults only after the exact match and the existing fallback have failed. Every launch that ran before still picks the same variant, so greedy replies are byte-identical.
Coverage: 176 variants.
Cross-check: of the 71 kit variant pairs that differ only in int64 alignment, all 70 the tool proves have byte-identical SASS, and the 71st is SASS-identical as well.
Pointers that do need 16-byte alignment: 7 (_gather_tiles OUT, _topk_sel KEYS/OUT and others). Every call site passes a buffer base for them, never a row slice.
Refusals no longer hang:
- A refusing rank drains its device and reports to rank 0.
- If every rank refused the same frame and drained within 30 s, the request gets an HTTP 500 with the reason and serving continues in lockstep.
- Otherwise rank 0 logs lockstep_lost and exits (70), so a supervisor can restart it, instead of hanging.

Validation.

Exactness: gates.py 7/7, the 157K canary, and the 12-prompt set plus two tail lengths (18,228, 116,704), all identical to the previous build.
Lengths that refused before: 2,161, 4,209, 18,226, 24,695, 121,124, 145,529 and 161,618 all complete now.
Speed: unchanged. kit_bench gives 3,602 / 3,883 tok/s and decode 152 / 91 / 209.
Soak, same sessions as the failing run: 210 requests, 0 errors, and no refusal or lost-lockstep lines on any node. Gates are 7/7 after.
In between: a run with only gap 1 fixed got through the soak with 2 clean HTTP 500s from gap 2 and no hang. That is how gap 2 was found.

If upstream's Zig line launches the AOT kit the same way, it likely has both gaps: any isplit tail whose row block is odd, and any replay-cut tail of 16 rows or fewer. Happy to PR the selection rule and the SASS tool.

@ashhart

Hi @ashhart,

Thanks for TensorFold and for making this the DeepSeek-V4.1 tracker. Here is a short status from a three-GB10 setup, plus a question about the best route. Nothing to merge; everything below is experimental and unqualified.

What we're doing

Three DGX Sparks, one rank per node, switchless ConnectX-7 ring (RoCE, no switch).
Checkpoint: a DeepSeek-V4.1-Flash EXL3 2.9bpw quant.
Engine: the Python TensorFold line (v0.6.0 with jayleaton's recipe patches), extended to TP3:
- attention heads split 16/24/24;
- the completed expert branch is continued across three ranks;
- world-3 Engram packs;
- FP32 partial exchange and an FP32 PV / IEEE QK attention path (background: DeepSeek attention probability rounding and cancellation-sensitive references jayleaton/deepseek-v41-tensorfold-spark#28).
Goal: run first, then decide whether it is worth hardening against our SGLang deployment on the same three nodes.

What we've seen so far (12 Oct)

Correctness looks good:
- the full model runs end to end on TP2 and TP3;
- all ranks pick the same token at every step, and repeat runs are bit-identical;
- CUDA graphs replay with no fallback, and logits are bit-identical to eager;
- 10 fixed prose/code/tool/system tasks give the same tokens on TP2 and TP3 and broadly match SGLang;
- an exact 8K and a 32K needle are recalled;
- about 45 GiB stays free per node at TP3.
- A teacher-forced TP2-vs-TP3 comparison is running now.
Speed is the problem (single stream, no drafts, from our screening harness, which adds per-step logit copies and hashes, so it is pessimistic):

decode c1 prefill

TensorFold TP2 ~17.9 tok/s ~94 tok/s

TensorFold TP3 (graphs) ~12.4 tok/s ~70 tok/s (32K prompt: ~457 s)

SGLang on the same 3 nodes 19.5 tok/s ~1,980 tok/s

Our first reading of the cause:
- Exact prefill is slow: fast_prefill is off because the fast prefill GEMMs change arithmetic.
- TP3 sends much more per step: the expert continuation does three all-gathers of ~560 KB per row, against one ~20 KB partial at TP2.
- No GPU Direct RDMA: NCCL reports GDR 0 on the GB10 HCAs.

Questions

Is TP3 worth pursuing? With 64 heads and the expert layout, would you expect an odd world size to pay off on GB10? Or should we stay at TP2, or go to TP4 (2D split) like @bertholomus and @ecohash-co?
Python line or Zig port? Given the Zig port's ~2× decode and exact prefill near 2K tok/s in this thread, should we stop extending the Python line and build any three-node work on the Zig dsv41 family instead? If so, is a three-rank layout something you'd consider there, or would you rather keep to 1/2/4?
Exact fast prefill: is there an exact prefill path for the Python line, or should we treat the Zig port's exact levers (SP rows, EXL3 v2 prompt GEMMs) as the reference?
GDR 0 on GB10: is that expected on GB10, and is there a recommended NCCL setup for a switchless three-node ring?
Any shortcut you'd recommend for getting a fair speed comparison against SGLang quickly?

We're happy to share logs, configs or receipts, or to run anything on the three-node setup. Thanks again.

Update 20: four prefill levers. bf16 owner exchange adopted, MXFP4 prefill experts go to port, two killed

We now gate changes in two tiers. Tier 1 must be byte-identical to the served build and faster. Tier 2 may change numerics, so it is env-switched and must pass five gates against the exact build:

qeval core ≥72/75, with at most 1 extension item lost.
needle 4/4, plus long_check 4×160K.
At most 2 of 12 greedy prompts diverge before token 32.
Faster on kit_bench and on the real-traffic mix.
A clean 1 h soak.

1. bf16 owner exchange (TF_DS_PROMPT_P16=1, tier 2): ADOPTED, now serving.

The SP-rows owner exchange sends the partials in bf16 instead of fp32.
Quality: core 72/75, the same as the exact build, with 0 extension items lost. Needle 4/4, and long_check passes. 0 of 12 prompts diverge before token 32.
Speed: kit_bench 4K+8K +12.4% and 32K+128K +8.0% (repeat pair: +5.3% / +6.6%). Prefill is now about 3.9K tok/s at 4K/8K and 4.1K at 32K/128K. Decode is unchanged at 151 / 91 / 211 tok/s (code / prose / structured).
Real-traffic mix (c1 / c2 / c4): 153.9 / 86.9 / 254.9 s, against 158.8 / 88.3 / 264.6 for the exact build. Our same-day SGLang deployment gave 149.5 / 86.6 / 228.8, so the gap is now −3% / −0.3% / −11%.
Soak: c3, 60 sessions, 210 requests, 0 errors, 0 refusals, and gates 7/7 before and after.

2. MXFP4 routed experts for prefill (tier 2): microbench passes, engine port next.

Test: a standalone microbench on real layer weights, comparing an MXFP4 copy of the routed experts with the served EXL3 v2 path.
Numerics: activation relative error is 5-9e-3 and output relative error 2-3e-7 against an fp32 reference. No non-finite values.
Speed per layer at 4,096 rows:
- Routed experts: 12.5 / 11.7 ms, inside our 13 ms target.
- Routed + shared: 13.8 / 13.0 ms, against 19.3 / 18.1 ms for EXL3 v2, which is 5.2-5.5 ms faster per layer. At 2,048 rows it is 1.5-1.7 ms faster.
Next: port it behind TF_DS_PREFILL_MXFP4_LAYERS=n, starting at n=16, then run all five tier-2 gates.

3. Exact pair-reduce (tier 1): killed before building.

Idea: replace the partial all-gather with a 2-rank ncclReduceScatter on split pair communicators.
Result: we replayed the real collective sequence offline. It was 148% slower unloaded and 22% slower with a compute hog on the GPU. Two-rank reduce-scatter on our fabric doesn't beat the gather plus local sum.

4. Retuned down kernel (TF_DS_EXL3_DN64=1, tier 1): exact, but slower in the engine; not adopted.

Microbench: the 64-wide down GEMM was 7% faster at 4K rows.
In the engine, on top of P16:
- Replies are byte-identical on all 14 prompts and gates pass 7/7.
- kit_bench is 5.7-6.0% slower in both pairs (3,611-3,637 / 3,858-3,860, against 3,855 / 4,099-4,116).
- Real-mix total is 512 s against 500 s.
The kernel stays behind its switch, off.

Code is on local branches (homelab/p16, homelab/dn64, MXFP4 work in progress). We're happy to open PRs against deepseek-v41-zig-tp4 for any of it.

Update 21: the Zig engine is now our production DeepSeek-V4.1-Flash endpoint; next is MXFP4 experts for prefill

Where we are. As of today, the Zig TP4 build from update 20 (homelab/p16, night5 + bf16 owner exchange) serves all of our DeepSeek-V4.1-Flash traffic through our LiteLLM proxy. That covers the agent route, the deep and reasoning routes, and vision/OCR. It replaces our SGLang deployment on the same four GB10 nodes.

What we checked before switching:

Through the proxy: plain chat, reasoning on and off, tool calls, streaming, and image input.
SGLang-style chat_template_kwargs (thinking) are accepted as well as enable_thinking, so no caller had to change.
When the engine is down for a planned window, our routes fall back to a smaller local model.

Same-day standing against our SGLang deployment:

	Zig (served)	SGLang
decode, code / prose / structured (c1, tok/s)	151 / 91 / 211	≈ 1.3-1.5× slower
prefill, 4K-8K / 32K-128K (tok/s, `kit_bench`)	3.9K / 4.1K	≈ 5.0K / 4.9K at 32K / 128K
real-traffic mix c1 / c2 / c4 (s, lower is better)	153.9 / 86.9 / 254.9	149.5 / 86.6 / 228.8

On agent-style traffic it is at parity at c1-c2 and 11% behind at c4. The remaining gap is prefill.

Plan for the next several days: MXFP4 routed experts for prefill (tier 2, TF_DS_PREFILL_MXFP4_LAYERS=n).

What it is. A second copy of the routed experts in the checkpoint's own MXFP4 format, used only by the encoder layers (0-19) during prompt passes. Decode keeps the EXL3 experts.
Why. The microbench from update 20 saves 5.2-5.5 ms per layer at 4,096 rows, which projects to about +8% prefill at n=16 and about +10% at n=20.
Memory. About 1.7 GiB per layer per rank. n=16 is about +27 GiB per rank, and castor's headroom sets n for all four.
Steps:
1. Port the O2 blocked MXFP4 grouped kernel from our Python line (torch stripped, extern-C instances).
2. Write the row-grouping builder, the parity rest plan, a plain 7-slot combine, and the shared expert on cuBLAS.
3. Add a slice loader that checks each file's split metadata against expertParts and refuses to boot below a memory floor.
4. Add the prompt2d expert branch. The whole EXL3 chain on a layer (rotation, SwiGLU, down, combine) is replaced, so both column partners switch together.
Gates. It is tier 2, so all five apply: qeval core, needle and long_check, early divergence on 12 prompts, both benchmarks, and a 1 h soak. We start at n=16, then try n=20.
How we're running it. Code is built while the engine serves. GPU checks and gates run in short planned windows, mostly overnight.

We'll post the gate results when they're in.

Update 22: MXFP4 prefill experts are built and fast, but miss our quality floor by one item; not adopted yet

The port. It is the O2 MXFP4 grouped kernel from our Python line, built into the Zig engine behind TF_DS_PREFILL_MXFP4_LAYERS=n. It covers layers [0, n) of the encoder, on prompt chunks only; decode stays on EXL3.

Memory. It loads 27.1-27.2 GiB of slices per rank at n=16, in 4-8 s with direct I/O. That leaves 15-17 GiB available, with no swap use through a 128K prefill.
Safety checks. The engine refuses to start on a cross-rank mismatch in n, or below a memory floor.

Checks before any numbers changed:

The kernels vs the Python kernel. On dumped real-layer pairs (layers 0-1, 2,048 and 4,096 rows), gate/up, the rest block and down outputs are bit-equal: 0 differing elements. Checks against float64 from the checkpoint bytes also pass. The routed experts take 12.6 / 11.8 ms per layer at 4,096 rows.
Switch off. The new build with the switch unset gives byte-identical replies to the served build on all 14 prompts, and gates 7/7.

Tier-2 gates at n=16, against the served build (P16):

gate	result	need
1. qeval core	71/75 (one code item flipped, `code_interval_intersect`); extension suites all perfect	≥ 72
2. needle / `long_check` 4×160K	4/4 / pass	4/4 / pass
3. early divergence (12 prompts, 256 tokens)	0 of 12; 11 identical through 256 tokens, 1 diverges at about token 70	≤ 2
4. `kit_bench` 4K+8K / 32K+128K	+5.2% / +4.7% (pair 1), +5.6% / +6.6% (pair 2)	faster in both
4. real-traffic mix c1 / c2 / c4	150.0 / 86.3 / 241.0 s against 153.9 / 86.9 / 254.9	faster

Gate 1 is the miss. By qeval's own comparison, one broken item with no gains cannot be told from noise on this set (p = 1.0; about 4 net breaks would be needed). We fixed the rules before measuring, though, and 71 is below the 72 floor.
So it is not adopted. The switch stays off and the served build is unchanged.

Other notes:

At c4 on the real mix, the gap to our same-day SGLang deployment (228.8 s) would close from 11% to 5%.
Decode on the code prompt reads about 2.5% lower in both n=16 arms (147.2 / 147.6 against 150.9 / 151.3). Prose and structured are unchanged. We're looking into it, since decode does not touch the MXFP4 path.

Next. A decision on gate 1, and possibly a repeat quality run. Separately, we'll look for the source of the code-decode dip.

@ecohash-co

@ecohash-co — Thanks for updates 16–22, extremely useful. I'm running the bertholomus TP2 recipe (v0.6.1, TF 1.0.2 native) on 2× DGX Spark with the same RoCE direct-attach. After applying the recipe's kit fix (58186d1) plus TF_DS_REPLAY_FLOOR=1 / TF_LOOP_GUARD=1, the short-prompt deadlock from bertholomus#16 no longer reproduces (11–13 token probes now return instantly). But we still hit the long-prompt variant gaps: an ~880k-char (~235k token) prompt hangs past 420s with no response byte, which looks like your Update 19 _index_score odd-tail / _sparse_attn_part <16-row tail findings. Would you be willing to share the night5 branch (or just commits edf22ac + 79e9c25) for a 2×GB10 TP2 setup? Happy to run your soak/qeval harness and report numbers back here.

Update 23: the MXFP4 quality miss reproduces; the served build stays as it is

Repeat quality run. We booted n=16 again and re-ran qeval at temperature 0. It scored 71/75 again, item for item the same as the first run. The engine is bit-deterministic, so this was expected. It shows the miss is a stable property of the n=16 numerics, not run-to-run noise. The item it loses, code_interval_intersect, is one that most builds we have measured fail, the MiaAI reference included. Only the two most recent served builds pass it.

Code decode. In interleaved decode-only arms (A/X/A/X, 5 reps each), code decode was 151.8 / 152.1 tok/s on the served build and 147.5 / 147.9 on n=16. That is −2.7%, and it reproduces. Prose and structured decode are unchanged. Our working explanation: the decode prompts are short (23-43 tokens), and their prompt pass also runs layers 0-15 on MXFP4. That gives a different code reply, and with it a different speculative-draft acceptance. On this reading the slowdown comes from the content of the reply, not from the kernel. We have not yet checked this against the engine log.

Decision. We are not relaxing the quality gate after seeing the result, so n=16 is not adopted. Production stays on the P16 build (adopted after QS1), and the MXFP4 switch stays off.

Next. Apply MXFP4 only to long prompts. Prefill time matters most there, and our kit_bench gains were measured there. The plan is a rank-invariant prompt-length threshold, so short prompts (most qeval items, the decode prompts, and typical chat turns) keep byte-identical P16 numerics. Long prompts would still get the roughly 5% prefill gain. It will go through the same tier-2 gates.

Update 24: MXFP4 on long prompts only. Quality holds and long prefill gains ~6%, but the pre-declared speed rule fails at 4K, so it is not adopted

Same 16 MXFP4 layers as update 23. The new switch TF_DS_MXFP4_MIN_ROWS=2048 sends a prompt chunk to the MXFP4 experts only when it has at least 2048 rows. The chunk's row count is the same on every rank, and the cross-rank check now carries it.

Exactness and quality (tier-2 gates, unchanged):

Short prompts (300 / 400 / 1500 tokens) are byte-identical to the served build, as designed.
Divergence: 0 of 12 prompts diverge within 256 tokens, the two ~150K prompts included.
qeval: core 72/75, with no item broken or fixed against the served build. Extensions show no losses. Needle 4/4. long_check (4×160K) is clean.
Decode replies have the same sha as the served build, and code decode is 152.0 vs 151.9 tok/s. The −2.7% decode effect in update 23 is gone, which supports the reading that it came from short prompts taking MXFP4.

Speed (kit_bench, interleaved A / L / A2 / L2):

	4K+8K prefill	32K+128K prefill
A (served)	3821.6	4138.5
L	3761.9	4377.5
A2 (served)	3758.0	4139.1
L2	3872.6	4401.2

Long prompts are +5.8% and +6.3%, consistent across both pairs.
4K+8K is −1.6% in the first pair and +3.0% in the second. The first L pass includes a cold 4K request at 2661 tok/s, which looks like the MXFP4 slices' first touch after boot.
The real-mix load test was faster for L.

Decision. The rule declared before the run required L to be faster in both pairs and both length groups. Pair 1 at 4K+8K fails it, so this is not adopted and the switch stays off. We are not reinterpreting the result after the fact.

What we're changing in the method. The divergence check and qeval barely exercise long-context numerics: even the all-chunks MXFP4 build matched the served build on our long synthetic prompts. So before the next tier-2 attempt we are adding:

a paired long-context replay set (real agent-session contexts from 8K to ~390K tokens; private data, so we'll report only aggregate numbers);
RULER / MRCR-style synthetic suites (public, generator in our tree);
logprob output in the engine, so builds can be compared by per-token KL / top-1 agreement rather than first-divergence only.

The next MXFP4 run will also give both arms the same warm-up of the MXFP4 path (one ≥2048-row prompt) before timing. That will be declared before the run.

Update 25: logprob output lands with a byte-identical build. MXFP4-long passes every old gate and is faster, but fails the new long-context KL gate, so it is not adopted

Logprob output in the engine (homelab/logprobs b16a910). logprobs / top_logprobs (≤20) on chat and completions, streamed or not, with token_id on each entry. There is also a non-standard forced_tokens field: it scores a fixed continuation teacher-forced, at decode speed. The cost is zero when logprobs are not requested.

The full vocab is already gathered on rank 0 for sampling, so no new collectives were needed. Log-sum-exp and top-k are computed on the host in a fixed order.

Identity with production, all byte-identical:

12/12 divergence prompts;
36/36 real-session replay items ≤64K with logprobs off;
60/60 items up to 256K with logprobs on;
decode replies, at the same speed.

New long-context gate. A replay set of 60 frozen real agent-session contexts (8K–256K tokens, private, so we report aggregates only). The served build's greedy replies (128 tokens) are scored teacher-forced by the candidate. We compare per-token top-20 distributions: KL over the shared top-20, and top-1 agreement.

The gate was declared before the run: mean KL ≤ 0.01 nats in every length bucket.

MXFP4-long result (n=16 layers, MIN_ROWS=2048). Every timed arm got the same warm-up, also declared before the run.

Gate	Result
qeval	core 72/75, 0 extension losses
needle	4/4
long_check	clean
divergence	0/12 early
short prompts	identical
decode	same sha
kit_bench	faster in both pairs and both groups: 4K+8K +6.0% / +5.8%, 32K+128K +7.7% / +6.9%
real mix	faster
replay KL	fail (below)

Replay KL by bucket:

bucket	top-1 agree	mean KL (nats)
8–16K	98.5%	0.0092
16–32K	98.7%	0.0094
32–64K	98.1%	0.0257
64–128K	98.2%	0.0134
128–256K	98.1%	0.0090

The per-item median is 0.0076 and p90 is 0.022.
One 32–64K item, at 0.214, drives that bucket's mean.
The argmax changes at about 1.5–2% of positions on real long contexts. Our synthetic divergence prompts never showed this (0/12), so the old gates could not see it.

Decision. The gate is not relaxed after the result: MXFP4-long is not adopted, and the served build is unchanged.

Next. Calibrate the KL gate against the last numerics change we accepted through the tier-2 gates (P16 vs pre-P16, same replay set). Our MXFP4 numbers are already known, so any MXFP4 call after calibration is an informed judgment, and we'll say so when we make it.

Update 26: calibrating the replay KL gate. The last numerics change we accepted (P16) drifts about half as much as MXFP4-long

Update 25's gate (mean KL ≤ 0.01 nats in every bucket) had no yardstick. CAL1 measures one, measurement only. It uses the served build with TF_DS_PROMPT_P16 removed, i.e. production as it was before P16. That build scored the same 60 teacher-forced replay items against the same P16 reference.

bucket	P16 vs pre-P16: top-1	mean KL	MXFP4-long vs P16: top-1	mean KL
8–16K	98.6%	0.0065	98.5%	0.0092
16–32K	98.6%	0.0040	98.7%	0.0094
32–64K	97.7%	0.0104	98.1%	0.0257
64–128K	98.2%	0.0059	98.2%	0.0134
128–256K	98.8%	0.0053	98.1%	0.0090
mean	98.4%	0.0064	98.3%	0.0133

What this shows:

Top-1 agreement is the same for both changes, about 98.3–98.4%.
MXFP4-long's KL is about 2× P16's overall, and 2.3–2.5× in the 16–128K buckets.
The 0.01 limit sits at the drift of a change we already accepted. P16 itself would have failed it in 32–64K (0.0104).
Free-running greedy text diverges anyway. P16 vs pre-P16 greedy 128-token replies over ≤64K were identical for only 25–42% of items per bucket. Any tier-2 numerics change changes most long-context continuations, so identity is not a usable bar there. Distribution distance is.

The update 25 verdict stands: the gate was declared before the run and MXFP4-long failed it. Whether to adopt it with a recalibrated gate is an owner decision. Any new threshold would be chosen knowing MXFP4's numbers, so it would not be a blind gate.

@ashhart

Update 27: final state. This push is closed and the Zig TP4 engine is our production DeepSeek-V4.1-Flash backend

What's serving. The Zig TP4 engine on 4× GB10 (2× DGX Spark + 2× ASUS GX10), EXL3 experts, 1M context. It has been our production flash/deep/vision/ocr backend since 10-12, replacing SGLang. The config is night5 + Engram prefetch + router own-rows + tsel + P16 bf16 owner partials (update 20). A watchdog auto-restarts it, but only when it is confirmed down.

Where it landed vs SGLang on the same hardware (sparkDash, same day, post-move; real mix = our replayed agent traffic):

	Zig	SGLang
decode prose, 1 / 4 / 8 streams (tok/s total)	112 / 209 / 254	84 / 148 / 194
decode code, 1 / 4 / 8 streams	166 / 255 / 279	125 / 229 / 287
decode structured, 1 / 4 / 8 streams	210 / 307 / 370	149 / 210 / 243
prefill 8K–128K (tok/s)	~3.6–3.9K before P16 (~4.1K est. with P16, from kit_bench +8–12%)	~4.7–5.0K
real mix, 1 / 2 / 4 concurrent (wall, s, lower is better; Zig = QP1 with P16, SGLang = SG2)	153.9 / 86.9 / 254.9	149.5 / 86.6 / 228.8

Decode is 1.3–1.5× faster on prose and structured output. On code the lead shrinks with concurrency, and SGLang is slightly ahead at 8 streams.
Decode holds at depth: median 142 tok/s at 8–16K, 125 at 128–256K and 108 above 256K, on real agent contexts.
Prefill went from a flat ~1.8K tok/s to ~4K, but SGLang is still ~15–20% faster on prefill.
Real mix: even at 1–2 concurrent requests, ~11% behind at 4.
Correction: an earlier version of this comment said "parity or better on a real mix" and "1.2–1.6× decode". That was wrong at 4 concurrent requests, and for code at 8 streams.
Exactness: every tier-1 change was byte-identical. P16 is the one tier-2 (numerics) change we adopted, and only after full quality gates and a soak.

What was tried and not adopted:

Change	Result
MXFP4 prefill expert copy (n=16)	qeval −1
MXFP4 on long prompts only	faster on every benchmark; failed the replay-KL gate (updates 24–25)
Engram row cache	exact, but no wall-clock gain
dn64 down GEMM	exact, but slower end to end
Pair reduce-scatter	killed by its pre-set rule
RDMA ring for the prompt	no-go: GPUDirect RDMA is off on these HCAs

Eval tooling that came out of this:

Logprobs in the engine: logprobs/top_logprobs, plus non-standard forced_tokens for teacher-forced scoring. Zero cost when off, and byte-identical to the served build.
A replay-KL gate on frozen real agent contexts (8K–397K). Calibrated in update 26: the accepted P16 change sits at a mean of 0.0064 nats. Our bar for any future tier-2 change is "no more drift than P16", declared before the run.
Lesson: greedy text identity can't be a gate for numerics changes. Even P16 matched word-for-word on only 25–42% of long continuations, while top-1 agreement stayed around 98.4%.

Why we're stopping. Our real traffic is mostly long agent sessions:

73% of requests are 64–256K prompts, with about 98% of their tokens served from the prompt cache.
Uncached prefill is now around 3 minutes of GPU time over a 6-hour day, against about 10 minutes of decode.

The remaining levers don't pay on that mix:

an EXL3 gate/up kernel rewrite (est. +5–7% prefill, multi-day, high risk);
lossy MXFP4/MXFP8 moves;
prefill/decode contention work, at this request rate.

We'll reopen if traffic shifts toward cold long prompts or a stronger engine or model appears.

Thanks to @ashhart and the TensorFold authors for the engine and the recipe. This has been running real agent workloads reliably.

@bertholomus

Thank you, @bertholomus, for offering the port and test time, and thank you, @ecohash-co, for updates 16 to 25. Running the native engine as a production DeepSeek endpoint on four GB10s, then turning down your own MXFP4 change because it missed a gate you set before measuring, is exactly the discipline we want here.

We've opened #847 to bring the DeepSeek-V4.1 work together as one family, with CUDA and Metal backends on the lane core, following the split @akol1 proposed on #635. @jayleaton's CUDA stack (#635 to #727) and @akol1's Metal family (#824 to #827) are the pull requests it covers.

Your branches are the most-run DeepSeek code anyone has on this engine. Where a piece of families/dsv41 is ahead of what's in those pull requests, please say so on #847. Anything we take from it carries your names as Co-authored-by. Your long-context KL gate would also be a good receipt for whichever backend lands first.

Thank you again for all of it.

@bertholomus

Thank you, @bertholomus, for the port and test time, and @ecohash-co for the four-GB10 production measurements and the corrected workload comparisons.

#847 now coordinates the shared DeepSeek family, the CUDA and Metal stacks, and the remaining qualification decisions. Your long-context KL receipts and any code ahead of those stacks belong there, with the source authors credited when their work is carried across.

We're closing this earlier offer onto #847 so the integration has one tracker. The branches and measurements in this thread remain useful references.

Posted by Ash's AI Assistant Synthena

@ashhart

TP3 on three GB10s (DGX Spark): DeepSeek-V4.1-Flash, Python engine, our measurements

Posted at @ashhart's request, to show whether TP3 helps our workload. All numbers come from our own runs on 3× DGX Spark (GB10; ConnectX-7 full mesh, RoCE, no GDR) with dealignai/DeepSeek-V4.1-Flash-EXL3-2.9bpw, on two or more boots unless stated.

Stack

TensorFold Python v0.6.0 (c4646171) with jayleaton's G19 family (4b235ad2), plus our own TP3 derivative:
- heads 16/24/24 (whole wo_a groups 2/3/3), experts width-split 768 per rank;
- FULL_CONE prefill;
- our own world-3 RoCE mesh for exchanges ≤ 1 MiB, NCCL above that.
Prefill levers, all bit-identical to the arm before them:
- gm experts at 2,048-row blocks (TP3 needs EXPERT_BLOCK ≥ 1463 for the gm scratch);
- FUSED_PROJ;
- a rank-ordered FP32 reduce-scatter → bf16 → all-gather at the prefill boundaries;
- a two-tile chunked attention for the 24-head ranks;
- indexer query rows split across ranks.

TP3 against TP2 on the same boxes and recipe

	TP2	TP3	TP3/TP2
1 request, drafts off (tok/s)	41.6	43.2 → ~45	1.04–1.08
1 request, drafts on (tok/s)	83.6	90.1	1.08
Single request, 12 paired cells, 6 boots			+6 to +13%
4 streams, drafted, all-active steady (tok/s)	122.9	156.1	1.27
8 streams, drafted, steady (tok/s)	not runnable (CUDA free 1.8 GB at 4 streams)	205.4	—

The 4- and 8-stream cells use the same prompt in every stream. Distinct-prompt runs are still to come.

Prefill (cold, drafts off, one stream, prose/code within 0.5 s)

Prompt	TP3, gm only	TP3, current exact stack	tok/s	Published TP2 G19 `full`
32K	16.4 s	14.2 s	~2,300	2,025
128K	65.2 s	54.7 s	~2,400	2,152

We have not run our own TP2 at 128K: our 8 GiB memory floor refuses it. The TP2 column is the G19 receipt (same pack, full + FULL_CONE).

Takeaway for our workload

A third GB10 gives about +5–10% for a single stream, because replicated reads and a fixed per-token overhead dominate.
It gives +27% at 4 streams and makes 8 streams and 128K+ contexts possible.
With the levers above, TP3 prefill now runs ~11–14% faster than the published TP2 numbers.

Caveats

These are speed screens with identical-token and raw-logit checks.
The full native-relative quality gate is still to run.
gm differs from the grouped kernel. On our R10 S3 check it scores top-1 99.990% with a mean NLL gap of 0.0014.

Next on our side

an exact world-3 session cache;
sequence-parallel prefill;
a decode latency bundle.

We will add those results here.

Navigation Menu

Repository navigation

DeepSeek-V4.1-Flash on the Zig engine (CUDA, 2 and 4 GB10): our port is yours to use, plus test time #559

5 remaining items

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

	decode c1	prefill
TensorFold TP2	~17.9 tok/s	~94 tok/s
TensorFold TP3 (graphs)	~12.4 tok/s	~70 tok/s (32K prompt: ~457 s)
SGLang on the same 3 nodes	19.5 tok/s	~1,980 tok/s

DeepSeek-V4.1-Flash on the Zig engine (CUDA, 2 and 4 GB10): our port is yours to use, plus test time #559

Description

Activity

ashhart commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 10, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

ecohash-co commented on Oct 11, 2026

5 remaining items

ecohash-co commented on Oct 12, 2026

mimeding commented on Oct 12, 2026

ecohash-co commented on Oct 12, 2026

ecohash-co commented on Oct 12, 2026

ecohash-co commented on Oct 12, 2026

plyldavid-maker commented on Oct 13, 2026

ecohash-co commented on Oct 13, 2026

ecohash-co commented on Oct 13, 2026

ecohash-co commented on Oct 13, 2026

ecohash-co commented on Oct 13, 2026

ecohash-co commented on Oct 13, 2026

ashhart commented on Oct 14, 2026

ashhart commented on Oct 16, 2026

mimeding commented on Oct 18, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions