[BlivionIaG]

Current head: 9489d0b443. Main 90b4fae62d, bafe3e07cb, 948e748584, d001c28f29 and release 09a9eb3317 was merged after the checks and device measurements below; those device measurements were not repeated on this head. CI is running on this head.

What this changes

EXL3 single-rank Flash Next loading, offloaded experts and the sm_120 kernel contracts. Expert leaves are float64 with fixed-tree merges, rounded once to float32 at the root.

• Issue: Refs [Draft] EXL3 Flash Next on CUDA: TP=2, BF16 paging, cache reuse, graphs #1113 (part of the agreed chain; does not complete it).
• Stacked on: no parent PR; this is the base of the chain (EXL3 Flash Next: expert-parallel TP=2 (same tokens as TP=1) and MTP drafts (2/8) #1442, EXL3 Flash Next: deep prefill (3/8) #1443, EXL3 Flash Next: KV pages from the core pool (4/8) #1444, EXL3 Flash Next: host-resident K/V on the tier (6/8) #1522, EXL3 Flash Next: CUDA graphs and the device planner (7/8) #1445, EXL3 Flash Next: draft lane guard for offloaded experts (8/8) #1446 stack on it).
• Based on main 6c5cef8941, merged through 484437b6ac.
• Overlap: atc overlap --json against four known peers (Flash Next TP=2 on CUDA: head and rollback graphs replay, and one-row split lanes skip the interleave copies #2837, flashnext cuda: the opt-in --tp 2 decode graph (TF_FLASHNEXT_DECODE_GRAPH=1) serves again instead of failing its first plain round #2980, flashnext cuda: the opt-in --tp 2 decode graph takes the dense split and keeps its graph across drafted verifies #2982, flashnext: the MTP head's streams by row through the lease's pages, no per-lane block outside the pool (#2444 P3, F2 part c) #2992) only, run on the integrated stack before the latest merges of main; not a global clearance. Conflicting files: Flash Next TP=2 on CUDA: head and rollback graphs replay, and one-row split lanes skip the interleave copies #2837 2, flashnext cuda: the opt-in --tp 2 decode graph (TF_FLASHNEXT_DECODE_GRAPH=1) serves again instead of failing its first plain round #2980 6, flashnext cuda: the opt-in --tp 2 decode graph takes the dense split and keeps its graph across drafted verifies #2982 8, flashnext: the MTP head's streams by row through the lease's pages, no per-lane block outside the pool (#2444 P3, F2 part c) #2992 6. Owner composition agreements exist for flashnext cuda: the opt-in --tp 2 decode graph (TF_FLASHNEXT_DECODE_GRAPH=1) serves again instead of failing its first plain round #2980, flashnext cuda: the opt-in --tp 2 decode graph takes the dense split and keeps its graph across drafted verifies #2982, flashnext: the MTP head's streams by row through the lease's pages, no per-lane block outside the pool (#2444 P3, F2 part c) #2992; my notice on Flash Next TP=2 on CUDA: head and rollback graphs replay, and one-row split lanes skip the interleave copies #2837 is unanswered.

Checks before the release merge

• Head d954f1232d. The device results below ran on 7cd3e412c41a; this head differs from it by the merges of main at 0591c4e85d, c34ff562ca, a3aecd8d0f and 484437b6ac (conflicts resolved in cuda_batch_moe.zig, cuda_batch_state.zig, the CUDA runner help, CHANGELOG.md, the fixture and one upstream test that calls passRules) and the pin fixture gradient-serving.tsv. The device cells were not repeated on this head.
• python3 tools/pr_check.py --base <main> --pr 1441: exit 0, 0 errors, 0 warnings; lean 339 vs 341 on main, 0 new. zig fmt --check zig: clean.
• Host tests: test-cuda-host, test-exl3-cuda-host, test-flashnext-exl3-host, test-exl3-prompt-host, compile-flashnext-host, test-flashnext-memory-host, test-qwen-serving-host, test-lane-api-host and test-lanes-see pass on Darwin arm64 on the EXL3 Flash Next: draft lane guard for offloaded experts (8/8) #1446 head, which contains this head; the Linux host suite is CI's.
• Linux x86_64 build with CUDA (-Dtarget=x86_64-linux-gnu): exit 0 on the EXL3 Flash Next: draft lane guard for offloaded experts (8/8) #1446 head, which contains this head (build only, no GPU executed on this head).
• CI: see the checks on this PR.

Device results on 7cd3e412c41a

Two 16 GB sm_120 GPUs, Linux x86_64, checkpoint above. One run per cell unless stated.

Cell	Result
kernel-contract exl3 flashnext (filtered)	61/61 kernels pass gates 1-4, 22/22 controls red as planted, 137 limit problems; all 137 are also reported by pinned main on the same GPU with identical registers, shared bytes and recorded values
1441-serve-tp1	concurrent equals solo, 12/12 rows
1441-tf, -pack, -layer, -pool	pass
1441-tfc, 1441-gen	differ from the golden only at the documented near-tie positions (70/72 top tokens; teacher-forced ids 16 and 17)
Chat turn resumed at 3072 and 4096 rows	same 48 token ids as a cold server
Decode and first token	about 15 tok/s greedy at one stream, 41 tok/s aggregate at four; cold 6144-token first token 73.45 s. No main comparison: main refuses this checkpoint with UnsupportedCheckpoint.

Not run or open

• The unfiltered kernel contract on this head. The filtered run covers every kernel this chain registers; the unfiltered one is about 630 kernels of other families, unchanged from main. On an older head it gave 632 kernels pass gates 1-4, 421/421 controls red, 564 limit problems (inherited; exit 1). A --jobs 4 unfiltered run on test(cuda): opt-in --jobs and --shard for kernel-contract #3093's head (main plus the runner change) gave 625 kernels pass, 410/410 controls red, 546 limit problems.
• Main against this head (main refuses this checkpoint) and the NVFP4 comparison (not on my hardware).

Receipt

• Boxes it reaches: single-rank EXL3 Flash Next loading and decode.
• New switch: none in this update.
• Environment: Zig 0.17.0. Host checks on Darwin arm64. Device runs on Linux x86_64 with two 16 GB sm_120 GPUs. Checkpoint turboderp/Qwen3.8-Flash-Next-exl3 at 69e33439ae950f17bcbe95c98f117d80f759ab6d.
• Kernel contract: Filtered contract above. This PR adds the EXL3 expert and prompt kernels and their controls.
• Model fidelity, measured on an older head and not repeated: native against an independent reference, 70/72 top tokens; the two differences are teacher-forced ids 16 (step 11) and 17, native margins 0.25 and 0.5 against reference 0.125 and 0.25. No tie waiver and no perfect-equivalence claim.
• Licence of new weights or drafter: n/a (Qwen Community License 1.0 on the checkpoint; no weights distributed, no new drafter).

Review requests implemented in source

• Main's test environment: video refusal describes the token budget without advertising a storage maximum is back byte for byte beside the unsigned-budget test.
• The layer-test helper reads the canonical global moe_f32 branch as production does.

Checklist

• The title and description match the head commit, including the receipt and the stack.
• New tests fail before the change, pass after it, and skip cleanly without their dependency. Host tests pass in CI and the device cells above pass; I did not run the new tests against the parent to show they fail before the change.
• Comments and docstrings are one line. No measurements or history in the source. Lean check (one-line comments within 120 columns): 0 new problems. I scanned every added comment line in the stack for measured results and history wording; the hits are worked examples and limits, not recorded results.
• No personal data, machine names, internal hosts or local paths in the diff or its commits. No AI attribution lines. Scanned every added line in the stack diff and every commit message for host and user names, addresses, local paths, AI trailers and tool names: none. No binary files are added. Receipts are not in the repository.
• Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off. No direct demonstration. The concurrent and drafted device cells run through the lane rounds.
• Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. Drafted: n/a, no drafts before EXL3 Flash Next: expert-parallel TP=2 (same tokens as TP=1) and MTP drafts (2/8) #1442. Resumed equals cold: a chat turn resumed at 3072 and 4096 rows gives the same 48 token ids. Concurrent equals solo: 12/12 rows.
• No precision traded for speed. If the bits change, the description says which sums change and why. Expert leaves are float64 with fixed-tree merges and one rounding to float32 at the root, so cancellation does not lose low-order terms; this is new arithmetic against the NVFP4 path, not a trade for speed.
• Decode and prompt processing are no slower than main at one stream and the widest admitted; with --drafter, no slower than without it. n/a against main (main refuses this checkpoint). Measured: decode about 15 tok/s greedy at one stream, 41 aggregate at four.

[ashhart]

Thank you for splitting #1113 into a chain and for the clear write-up. A pilot gated to sm_120 with a named refusal is
a good shape. A few things need to change before this one can land:

1. Please remove the int8 KV cache entirely. That means flashnext_exl3_kv.cu, exl3/kv/cuda_exl3_kv8.zig, the int8
   parts of exl3/kv/cuda_exl3_kv.zig, its seven contract records, and the hooks in cuda_batch.zig:358,
   cuda_snap.zig and cuda_prompt.zig. We don't take lossy formats, even opt-in. As written it also reaches every
   Flash Next checkpoint: Options.kv8 (cuda_lanes.zig:61) lifts main's --kv-cache fp8 refusal
   (native/cuda_refusal.zig:48), and setup() enables int8 before its if (!exl3) return
   (exl3/kv/cuda_exl3_kv.zig:151-154). The same applies to TF_FLASHNEXT_PLE_EMB=bf16 (cuda_exl3_ple.zig:25).
2. Please move the shared EXL3 changes into their own PRs, each with a device byte-equality receipt against main and
   re-derived contract records. They are the ldmatrix path in exl3_candidate.cu:127 (it runs on sm_121 too),
   prompt_chunk 256 to 2048 in exl3_candidate.zig:47 (8x the scratch for every user), and the counter reset inside
   tf_exl3_rot_in (exl3_linear.cu:112, exl3.zig:55-69). Nothing in this PR uses the first two.
   tools/zig/exl3_prompt_ldmatrix_fixture.py looks like the right harness for the first.
3. Please drop the comment-only edits to the 15 files outside the family (DeepSeek V4.1, GLM, H3, Drama3, HIP,
   tree_attention.cu, and others). EXL3 Flash Next: CUDA graphs and the device planner (7/8) #1445 reverts them anyway.
4. cuda_prompt.zig:183 moves every BF16 embedding table, NVFP4's included, to mapped host memory. Please limit that to
   EXL3, or measure it on NVFP4.
5. The new kernels need entries in the CUDA contract runner for gates 1 to 4. test-kernel-contract only checks
   their launch records.
6. Please remove the local paths from the tools and from zig/tests/fixtures/exl3_ref_*.json. Nothing reads the
   fixtures. One row in them needs an explanation: at position 27 the GPU picks 5269, while your f64 oracle and the
   reference both pick 6278 with a 0.95 gap, and the row is labelled a tie.

Once those are done, what remains is the pending repo-client receipts and a concurrent-equals-solo run on this slice.

[ashhart]

Thanks @BlivionIaG for the update. This slice is much tighter now. As far as we can see, the int8 KV cache, the shared EXL3 changes, the out-of-family comment edits and the local paths are gone, and the mapped-host embedding table is limited to EXL3 packs. A few things remain:

1. fn_exl3_split_f16 still has only a launch record. Please add it to the CUDA contract runner for gates 1 to 4.
2. Removing the fixtures takes the position-27 row out of the repo, but we'd still like to understand it: the GPU picked 5269 while the f64 oracle and the reference both picked 6278 with a 0.95 gap.
3. On EXL3 the prompt arithmetic depends on the chunk cuts (cuda_lanes.zig:87), so the 1024-row chunk and TF_FLASHNEXT_CHUNK change tokens. Prompt rows should match serial decode whatever the chunk. This is the same point as on EXL3 Flash Next: deep prefill (3/8) #1443.
4. Small cleanups: the bf16 Emb option in cuda_exl3_ple.zig, and the comments that cite an internal lever number and a work-in-progress branch name.

After that, we still need the repo-client receipts and a concurrent-equals-solo run on this slice.

[BlivionIaG]

Kernel-contract result on this exact head (7cd3e412), nvcc 13.4, sm_120, filtered to names containing exl3 or flashnext:

(filtered) SUMMARY 61 kernels: 61 pass gates 1-4, 0 fail, 0 known findings; 22 controls: 22 red as planted, 0 missed; 137 limit problems; gates 493768 ms, limits 260 ms

All 22 planted controls go red. The 137 limit problems are also reported by pinned main on the same GPU with identical registers, shared bytes and recorded values, so none comes from this PR. Filtered because the full contract runs about 630 kernels of other families and took 3.2 h with the GPU mostly idle; those kernels are unchanged from main. The last unfiltered run was on an older head and has not been repeated on this one.

Also restored main's video-budget refusal test unchanged beside the unsigned-budget test, and merged current main (6c5cef89).

[ashhart]

Thanks @BlivionIaG. The contract result on 7cd3e412 answers the ask: 61 of 61 EXL3 and Flash Next kernels pass gates 1–4, all 22 planted controls go red, and the 137 limit problems match pinned main's on the same GPU with the same recorded values. Filtering to this PR's kernels is fine, since the rest are unchanged from main. Thanks also for putting main's video-budget refusal test back beside the new one.

What's left is what your description lists under remaining gates, on this head:

• Exactness: tools/bench_concurrent.py --alone --serial, or token_sha against "draft": false. Drafted equals undrafted, a resumed prompt equals a fresh one, and concurrent equals solo.
• Main against this branch on your sm_120 card, same checkpoint: decode from tools/bench_openai.py, and the cold first request's ttft_s on a long prompt.

Once those are posted and CI on this push is green, we'll run main against this head on Qwen3.8-Flash-Next NVFP4 on a GB10, to check the default CUDA path is unchanged.

• Synthena (Ash's Agent)

[BlivionIaG]

Device results on this exact head (7cd3e412), same checkpoint, TP=1 on one sm_120 card, tools/bench_concurrent.py --alone --serial and tools/bench_openai.py:

Concurrent equals solo. --alone --serial, greedy and sampled (temperature 1.0, seeded), code and chat prompts, 1, 2 and 4 streams: 12 of 12 rows with 0 unequal and 0 failed in both the alone and the serial comparison.

Decode and TTFT.

	greedy	sampled
bench_openai decode tok/s, median of 3 (code prompt)	15.1	8.4
bench_openai decode tok/s, median of 3 (chat prompt)	15.2	8.3
bench_concurrent aggregate tok/s, 1 / 2 / 4 streams (code)	14.7 / 25.7 / 41.1	8.9 / 9.2 / 9.1
bench_concurrent aggregate tok/s, 1 / 2 / 4 streams (chat)	14.8 / 25.8 / 39.9	9.1 / 9.3 / 8.9

Cold first request, 6144-token prompt: ttft_s 73.45 s, then 12.5 tok/s decode (64 tokens).

Model checks on this head. 1441-tf, pack, layer and pool pass. tfc and gen show the same two positions as in the description's fidelity record (teacher-forced 22 of 24 on the first prompt, 24 of 24 on the other two; the greedy run differs from the golden at step 11 of the first prompt, golden margin 0.125), unchanged.

Not covered here:

• Drafted against "draft": false: this head has no drafts for EXL3 (they arrive in EXL3 Flash Next: expert-parallel TP=2 (same tokens as TP=1) and MTP drafts (2/8) #1442).
• A resumed prompt against a fresh one: my first run timed out at its 700 s cap after the cold request; I'm rerunning it with a larger cap and will post the row.
• Main against this branch: main does not load this checkpoint, so there is no main number for it. I'll attach main's refusal from a probe run.

[BlivionIaG]

Resume row for 7cd3e412 (same checkpoint, TP=1, --prompt-cache-gib 2, --no-drafts; the cell is 1441-resume-only):

• A 4099-token prompt, 64 greedy tokens, sent four times (the first, two repeats, one concurrent request): all four produce the same 64 tokens (token_sha a6dbe3fe1965, identical output SHA-256).
• None of them resumed. Every response reports cached_tokens 0, and the server log reads prompt cache: 4099 tokens, resumed at 0, kept 0; 0 states ... (hits 0, misses 4 ...). The repeats took a full prefill each (TTFT 53.6, 49.8, 49.7 and 50.1 s).
• The cell therefore fails its own --require-reuse check. I am not reporting "resumed equals fresh" for this head: at 7cd3e412 the EXL3 engine keeps no prompt states, so there is nothing to resume and nothing to compare. The same cell is run on the later heads in the chain, where the resume paths are added, and those results will be on those PRs.

I'm not claiming a main comparison for this row either: main refuses this checkpoint, so there is nothing to put next to it.

[BlivionIaG]

Correction and a real resume result for 7cd3e412. My earlier row used a raw token-id prompt, and at this head those never keep a prompt state by design: Flash Next EXL3 declares a 1024-row prefill grid (families/flashnext/cuda_lanes_load.zig, .snap_grid = chunk), which turns off the identical-prompt repeat marks, so only chat turns with history keep states at a chunk start. So I ran a chat turn instead (1442-chat-resume-1441, this head's binary, TP=1, --no-drafts, --prompt-cache-gib 2; greedy, 48 tokens, thinking off, ignore_eos so that every request returns exactly 48 token ids):

request	prompt tokens	cached_tokens	token SHA-256 (first 12)	TTFT s
turn 1 (about 4.1k tokens of context)	4097	0	bbe5d1fbbefd	35.0
turn 2 (turn 1, its reply, a new question), cache on	4322	3072	2dae136c9834	10.8
turn 2 again, cache on	4322	4096	2dae136c9834	2.6
turn 2 on a fresh server, --prompt-cache-gib 0	4322	0	2dae136c9834	37.0

Resumed equals fresh: the same 48 token ids on the resumed turn, the re-hit and the cold server. The model ends its short answer early and the remaining ids follow the end token, so the comparison holds for those too but says less about the content than a long reply would.

[BlivionIaG]

Main on this checkpoint, as promised: the pinned main binary exits before it is ready with error: flashnext: config.json's quantization_config.quant_method differs from the supported Flash Next CUDA layout / error: UnsupportedCheckpoint (receipt from the #1444 probe cell, same checkpoint). So there is no main decode or TTFT number to put next to this head's; the comparison on NVFP4 on your side is the one that exists.

[BlivionIaG]

Merged main at 0591c4e85d into this branch (be5274affd) after two new main commits (58bd8ac44e, 77b02600cd: the TP=2 prompt expert exchange) conflicted with it. The later PRs carry the same merge. Resolutions, in the files that conflicted:

• cuda_batch_moe.zig: both kept. moeAdds now skips the slot sum when either the EXL3 canonical branch (canon_f32) or main's rank exchange (pending_f32) already left the sum in moe_f32: if (!canon and !self.pending_f32) try sumSlots(self, 0, n). main's sumSlots sits beside this branch's leafTree, leafPartial and kept.
• cuda_batch_state.zig: this branch's allocationSizes list keeps its shape; it takes main's two size changes (y is (pairs + slots) * dims * 4 and moe_f32 is (cap + 1) * dims * 4, the spare row for the odd-row pad).
• tests/cuda/main.zig: this branch keeps the help text in runner_usage.zig; main's one changed help line (flashnext-rank ... exchange) is applied there.
• gradient-serving.tsv: raw pins refreshed to the merged bytes.

main's exchange is reached only from the NVFP4 native MoE (cuda_moe_native.zig), not from the EXL3 experts path. Checked on the merged heads: test-gradient-serving, test-exl3-cuda-host, test-flashnext-exl3-host, test-exl3-prompt-host, compile-flashnext-host, test-flashnext-memory-host, test-qwen-serving-host and the x86_64 Linux build pass. The device results on these PRs are from the heads before this merge; I have not repeated them on it and each body says so.

[Nipale-ai]

Hi @BlivionIaG, hi @ashhart,

we ran Qwen3.8-Flash-Next EXL3 3.05 bpw on one RTX 5090 (32 GB) on the chain at 06bdb53 (head of #1443 on Nov 11; the branch has moved since, so the patches may need a rebase), experts in pinned host RAM. From an nsys profile we wrote four patches. All are environment switches; unset keeps the head's behavior.

• TF_FLASHNEXT_TABLE_BATCH=1: writes expert table entries to a host shadow and uploads the table once per stage, instead of 18 memsets per reloaded expert.
• TF_FLASHNEXT_EVICT=lru: evicts the least recently used line instead of the lowest lifetime count.
• TF_FLASHNEXT_EXL3_HIT_FIRST=1: computes held experts while the missing ones copy; same plan, copies and victims.
• TF_FLASHNEXT_PLE_PREFETCH=1: madvise(WILLNEED) on the n-gram rows of the coming window, so the gather does not wait on page faults.

Decode, our speed test (1 slot, 32K window, MTP on, thinking off, 512 tokens): 81.3 tok/s with the #1441 script settings (95 expert lines); 130.5 with 253 lines and no patches; 157.2 with 253 lines, the four patches and LINE_DELTAS (page cache emptied before the start). The deltas were fitted on traces of the same prompt set, so that part is in-sample. Not tested on a 16 GB card.

Equality: 114 of 114 concurrent-vs-alone comparisons and 125 of 125 token hashes against the unpatched head match. Host tests, all switches unset: 6,352 of 6,382 passed, 30 skipped, 0 failed. One test expects a clean environment and fails when TABLE_BATCH=1 is set globally (cuda_exl3_pool_order_test.zig:258).

Details: https://niklaslenz.com/v2/qwen38-flash-next-5090-exl3-3-05bpw-tensorfold-pr-06bdb53c-our-patch3-blivion/

Patches: https://gist.github.com/Nipale-ai/e2aaf18acd3fbad93f0a3a6eb5916cdc (P0 is a trace-only helper). Happy to open a PR against your branch.

[BlivionIaG]

Hi @BlivionIaG, hi @ashhart,

we ran Qwen3.8-Flash-Next EXL3 3.05 bpw on one RTX 5090 (32 GB) on the chain at 06bdb53 (head of #1443 on Nov 11; the branch has moved since, so the patches may need a rebase), experts in pinned host RAM. From an nsys profile we wrote four patches. All are environment switches; unset keeps the head's behavior.

• TF_FLASHNEXT_TABLE_BATCH=1: writes expert table entries to a host shadow and uploads the table once per stage, instead of 18 memsets per reloaded expert.
• TF_FLASHNEXT_EVICT=lru: evicts the least recently used line instead of the lowest lifetime count.
• TF_FLASHNEXT_EXL3_HIT_FIRST=1: computes held experts while the missing ones copy; same plan, copies and victims.
• TF_FLASHNEXT_PLE_PREFETCH=1: madvise(WILLNEED) on the n-gram rows of the coming window, so the gather does not wait on page faults.

Decode, our speed test (1 slot, 32K window, MTP on, thinking off, 512 tokens): 81.3 tok/s with the #1441 script settings (95 expert lines); 130.5 with 253 lines and no patches; 157.2 with 253 lines, the four patches and LINE_DELTAS (page cache emptied before the start). The deltas were fitted on traces of the same prompt set, so that part is in-sample. Not tested on a 16 GB card.

Equality: 114 of 114 concurrent-vs-alone comparisons and 125 of 125 token hashes against the unpatched head match. Host tests, all switches unset: 6,352 of 6,382 passed, 30 skipped, 0 failed. One test expects a clean environment and fails when TABLE_BATCH=1 is set globally (cuda_exl3_pool_order_test.zig:258).

Details: https://niklaslenz.com/v2/qwen38-flash-next-5090-exl3-3-05bpw-tensorfold-pr-06bdb53c-our-patch3-blivion/

Patches: https://gist.github.com/Nipale-ai/e2aaf18acd3fbad93f0a3a6eb5916cdc (P0 is a trace-only helper). Happy to open a PR against your branch.

Feel free to open a pr, that's great, on my end I deferred pushing more patches/improvements for a while since it's difficult to retest and rebase the whole chain every time.

[BlivionIaG]

Second merge of main into this branch (fb21c01673, main at c34ff562ca, carried through the later PRs). One conflict beyond the pin fixture: cuda_lanes.zig, where main added snap_intermediate = snap.ladder next to the snap_save/snap_restore/snap_drop lines this branch renames to copies. I kept both. The ladder is off unless TF_PROMPT_INTERMEDIATE=1 (cuda_snaps.zig), so a default EXL3 run is unchanged by it; I have not run it with the ladder on for an EXL3 checkpoint, and its rungs on EXL3 prompt states are untested here.

Checked on the merged heads: test-gradient-serving, test-exl3-cuda-host, test-flashnext-exl3-host, test-exl3-prompt-host, compile-flashnext-host, test-flashnext-memory-host, test-qwen-serving-host, test-lane-api-host, test-lanes-see, the lean check (0 new) and the x86_64 Linux build pass. The device results are still from the earlier heads, as each body says.