[BlivionIaG]

Draft for upstream review; not merge-ready. This ex3-on-upstream-wip/branch/exl3-on-upstream-wip integrates the Qwen3.8-Flash-Next EXL3 path with upstream main (base d1018b318). It is a broad working draft intended to be split into focused PRs before merge. The implementation has working paths, but required correctness, capacity, and performance gates remain open; in particular, CUDA graphs have not demonstrated a repeatable speedup.

What it adds

• EXL3 weight loading and CUDA execution in TensorFold's Flash Next path.
• Routed-expert offload with per-layer GPU pools, pinned host banks, and asynchronous H2D staging.
• Opt-in persistent decoder/MTP expert-frequency priors, with private per-model/rank storage, checked reload, adaptive startup decay, pool warming and atomic graceful-stop saves.
• Expert-parallel TP=2 with split routed experts, shared-expert execution, vocabulary-sharded lm_head, and NCCL collectives.
• BF16 KV paging to host RAM, snapshots/restores, and prefix/cache reuse, with an opt-in per-rank shared block arena. Ordinary-context private/shared snapshot and served reuse gates pass; the 200k × 2 target remains unqualified.
• MTP drafting and CUDA graph capture/replay. Decode graphs cover eligible rounds/windows; prefill graphs capture dense sub-blocks, while expert scheduling/copies remain eager.

These are implemented paths in the branch, not claims that every supported shape or failure mode is validated. The evidence and remaining gates below define what has actually been checked.

Memory and execution split

Configuration shown: TP=2 with BF16 KV paging enabled (TF_FLASHNEXT_KVPAGE=1). Each request uses both GPUs. The boxes show placement, not allocation sizes.

flowchart LR
  subgraph HOST["CPU scheduling and host RAM"]
    direction TB
    RAM0["Rank 0 RAM<br/>Expert bank: 0-255<br/>MTP expert bank<br/>Full BF16 KV backing<br/>Prompt snapshots"]
    RAM1["Rank 1 RAM<br/>Expert bank: 256-511<br/>Shared weight source<br/>Full BF16 KV backing<br/>Prompt snapshots"]
  end

  subgraph VRAM["Two GPUs - one TP=2 model"]
    direction TB
    GPU0["GPU 0 / rank 0 / 16 GiB<br/>Dense/router/attention<br/>Scales + lm_head half<br/>Routed cache: 0-255<br/>MTP weights/pool/state<br/>BF16 window + blocks<br/>GDN, tables, scratch"]
    GPU1["GPU 1 / rank 1 / 16 GiB<br/>Dense/router/attention<br/>Scales + lm_head half<br/>Routed cache: 256-511<br/>Resident shared expert 512<br/>BF16 window + blocks<br/>GDN, tables, scratch"]
  end

  FREQ["Optional local disk profiles<br/>Expert counts per model + TP rank<br/>No weights, KV or graphs"]
  FREQ <-->|"Checked priors / graceful save"| RAM0
  FREQ <-->|"Checked priors / graceful save"| RAM1

  RAM0 <-->|"PCIe transfers"| GPU0
  RAM1 <-->|"PCIe transfers"| GPU1
  GPU0 <--> NCCL["NCCL collectives<br/>MoE all-reduce<br/>Logit all-gather"]
  GPU1 <--> NCCL
  classDef gpu fill:#e8f2ff,stroke:#3765a3,color:#152b46
  class GPU0,GPU1 gpu

Loading

• Replication and sharding: paging currently disables attention-head sharding, so each rank has the full attention/KV layout and its own host backing allocation. Decoder routed experts and lm_head vocabulary columns remain split. With paging disabled, a separate attention-head split is available; it is not the configuration drawn here.
• Expert movement: host-planned paths use asynchronous H2D copies. The deep-prefill ring reuses decoder pool storage across layers. In device-planned decode, GPU kernels read mapped host banks to fill cache lines; overflow experts can be read in place. Routed caches retain experts for reuse, while the shared expert has a separate resident allocation on rank 1. CPU scheduling does not mean CPU expert execution.
• KV and snapshots: selected older KV blocks are fetched into the device block cache; recent window rows are written back to RAM during window rebasing or snapshot flushing. Snapshots preserve host KV history, device windows/cache contents, pager tables, and recurrent/PLE/MTP state where present. The PLE table itself is file-backed; gathered embeddings are uploaded.
• Shared block option: TF_FLASHNEXT_KVPAGE_SHARED=1 uses one block pool per rank with fixed per-lane windows and generation-tagged ownership. Private per-sequence paging remains the default. Shared version-2 snapshots retain only owned blocks and remap them into destination slots; they do not overwrite the whole arena or deduplicate identical-prefix payloads. Capacity pooling can trade cache hits for VRAM, so it is not a throughput claim.
• Scope: the NCCL box denotes logical communication, not a direct GPU peer link. Memory budgets depend on configuration. Ordinary-context TP=2 paging/cache behavior has been exercised; 200k × 2 with paged KV remains unproven. Prefill expert scheduling/copies remain eager when dense prefill graphs are enabled.

Proposed PR segmentation

This is a proposed split plan, not a claim that the listed PRs already exist. Each future PR should own one coherent layer, add focused tests, and record its correctness and performance evidence. Until then, this mega PR remains a tracking draft and should not be merged as one piece. The current branch already contains prototype implementations for these areas; the listed gates are remaining review/acceptance criteria.

1. EXL3 foundation and resident forward: isolate format/config parsing, weight loading, and the upstream CUDA adapter. Acceptance: layer/model oracle gates for the resident path.
2. Expert pool and offload: isolate host-bank layout, device line pools, replacement policy, async copies, and pool lifecycle. Acceptance: pooled output matches resident execution; measure copy/compute overlap rather than assuming it.
3. TP=2 execution and serving: isolate expert partitioning, shared-expert placement, vocabulary-head split, collectives, and rank lifecycle. Acceptance: TP=2 logits/outputs meet a defined numerical tolerance and the server works on both GPUs.
4. BF16 KV paging and state reuse: isolate device/host page tables, paging/eviction, paged snapshots/restores, and prefix branching. Acceptance: paged/resident equivalence where comparable; exact state/logit parity for snapshot/restore and cache reuse, plus long-context tests.
5. MTP drafting: isolate draft-head input, TP=2 draft/verify, and prompt-cache interaction. Acceptance: committed drafted tokens equal serial decoding with drafting disabled.
6. CUDA graph decode: isolate stable graph inputs/outputs, TP=2 collectives, offload boundaries, and capture/replay lifecycle. Acceptance: graph/eager output equality over repeated replay and snapshot restore.
7. CUDA graph prefill and scheduling: isolate chunked capture, expert-major scheduling, staging overlap, and profiling-driven tuning. Acceptance: repeatable graph-on improvement in controlled prefill/decode measurements without correctness or memory regressions.
8. Final capacity and product integration: run repeated 200k-context × 2-concurrency TP=2 tests with BF16 paging/cache reuse, resource margin, and failure cleanup; then consider family/recipe registration. Do not register until preceding gates pass.

Proposed dependencies: the foundation precedes the other slices; expert offload and TP=2 build on it; paging needs stable TP=2 allocation/lifecycle; MTP needs TP=2 and cache semantics; graph work depends on stable paging/offload and eager behavior. Decode and prefill graph slices can be developed in parallel once interfaces are fixed, but should be integrated and reviewed in dependency order. Avoid duplicate implementations: use stacked branches or clean cherry-picks, and keep this umbrella draft as an index until the sequence is complete.

Benchmark report

The detailed tables, exact tested commit, model/checkpoint identity, runtime configuration and methodology are in the 21 October benchmark comment. These historical results apply to source ca2b13c88315c9c1c7a9ce17c002226f3585f73a, not subsequent candidate changes. The single matched graph/eager run did not establish a speedup.

Evidence and remaining gates

• Prior replay/cache correctness checkpoint: runtime 4afc6169 passes 192 EXL3 host tests, real native/GPU-runner builds, private/shared buffer-reuse/snapshot checks on both cards, and exact TP2 prefill origin/grid rebinding through dense/sparse and paging-split transitions. The native MTP/graph/shared-BF16 cold/warm/branch/rewind gate retains all eight reference hashes. Exact sources, binaries, model, methodology and scope. Correctness-run timings overlapping compilation are excluded; isolated graph speedup remains unproven. Retaining buffers does not retain old-owner graphs.

• Prior replay/cache telemetry checkpoint: runtime 975c1df9 passes 177 EXL3 host tests, the native CUDA build, diagnostics-enabled TP2 shared-BF16 cache reuse with separate valid expert/KV counters, and diagnostics-disabled c8 hash/replay qualification. Source b3a7b565 also passes the focused real TP2 width-3/4/16 rebind fixture. Full source, model, methodology and bounded measurements. Those receipts do not qualify subsequent runtime changes. No repeatable speedup, independent golden fidelity or 200k capacity is inferred.

• New shared BF16 checkpoint: historical runtime ccb32fc3; these direct fixture results are not relabeled as current-head tests. 153 EXL3 host tests and real native/GPU-test CUDA builds pass. Matched TP=2 private/shared sparse-paging fixtures pass exact restored-state checks, owned-entry remapping, recycled/asymmetric generations, foreign-request isolation and actual prefill/decode graph replay. Two 16,424-token HTTP histories with MTP and graphs pass cold/warm/branch/rewind reuse against the private reference. The full qualification comment separates these tests from the earlier same-binary 0de41d37 allocation measurement (920 MiB lower sampled peaks per GPU). Shared mode stays opt-in; none of these results establishes independent model fidelity, repeatable graph speedup or 200k × 2 capacity.

• Served BF16 paging/cache gate on ca2b13c8: two concurrent HTTP histories, each with 16,424 prompt tokens and 32 generated tokens, passed cold, warm, branch and rewind phases with graphs and MTP enabled. Full response hashes matched a separately started cache-disabled reference for every phase. Cache-enabled requests reused 16,407–16,417 tokens on warm/branch/rewind phases; warm native prefill times were 0.944 / 0.974 seconds. This proves ordinary-length prefix reuse and continuation parity, not long-context capacity, independent model fidelity, or exact internal-state parity at this new head.

• Completed matched graph/eager gate on ca2b13c8: all four exact 1024/512 and 16384/1024 cases at c1/c8 passed in each mode (36 valid responses, 18/18 matching full token/text hashes). Binary and per-rank expert geometry were identical. Both ranks replayed decoder graphs while all eight requests decoded in each c8 cell, and replayed prefill graphs in the long-input cells. Graph-run physical GPU peaks were 14,958 / 14,996 MiB; automatic expert pools were 10 / 23 lines per layer. This is an ordinary-context functional/capacity pass, not a 200k claim.

• 21 October checkpoint: source ca2b13c8 is integrated without rewriting the draft history. The parent EXL3 host gate passes 90/90 tests, scoped formatting/diff checks pass, all tracked source hashes match the Linux build tree, and the native CUDA sm_120 build succeeds. This adds workload-scaled graph budgets, class-aware local graph charges, selective coordinated TP LRU retirement, and a tested local-reclamation fallback. Graph memory charges remain heuristic, not measured exec sizes.

• New snapshot hardware receipt: on the preceding d3662d1b runtime with test-only upstream loader compatibility fixes, both RTX 5060 Ti ranks passed the compiled two-lane BF16-paged snapshot fixture. Restored canonical KV history, recurrent/index/PLE state and continuation logits matched exactly. Each graph iteration recorded 27 prefill and 5 decode replays per rank. This uses 24 expert lines per layer and no MTP head; it is not an eight-lane capacity, served prefix-cache, or 200k test. The new checkpoint separately passes the served cache-reuse gate below; this older direct state/logit fixture is not relabeled as a current-tree run.

• Performance remains an acceptance blocker: graph decode was 3.4–5.4% slower across the four cells in this single matched pair. No repeatable speedup is established. Cache-key churn and the graph/expert/KV memory partition remain optimization work; increasing reserved memory is not itself a performance fix. The earlier interrupted diagnostic overlapped compilation and is not a controlled performance receipt.

• On earlier source-tree revisions, TP=2 serving, draft-on/off output parity, BF16 paging, snapshot/cache reuse, and graph/eager correctness were exercised on development hardware. These results are not a full validation of the exact current PR head.

• CUDA graphs currently match eager in tested paths, but controlled profiling has not established a graph-on speedup. That is a hard performance gate, not a claimed win.

• The strict golden-prompt near-tie and remaining TP numerical-parity questions still require closure.

• Ordinary-context TP=2 private/shared paging, scoped state/logit restoration, graph replay and served cache reuse now pass on runtime ccb32fc3 as described in the qualification comment. Cancellation/overflow/live-rank failure injection, repeated performance qualification and the 200k × 2 BF16 paging target remain required before merge/registration. The earlier near-200k × 2 retrieval test used resident BF16 KV and is not evidence for paged-KV capacity.

• Fresh tools/zig/lean_check.py run: 1,762 findings on the current source versus 444 on pinned upstream d1018b318 (1,318 more). This remains a CONTRIBUTING blocker, not a clean result.

• CPU expert execution is out of scope. BF16 remains the required KV-cache format; int8 KV is not part of the serving path.

Branch: exl3-on-upstream-wip
Current fork head: 0272aafe266299c078976c74fbe5347bdf484f7c (qualified target-layer history, DMA coalescing, opt-in transfer provenance and actual-layer GDN graph keys; bounded TP2 BF16 cache/parity and screening measurements pass)
Upstream base: d1018b318fb51981c4a18f12e40269a4d3ee9713

Persistent expert-profile checkpoint (22 October)

Native source babc9762e2dbc999c9b68df252d619830203d693, binary SHA-256 86a7b20b6bac2bc5b2b78f9c89654a21cd1c0bd2ad9e6806ee1dc70abe110907, passes a real TP2 graceful stop/start gate with graphs, MTP and shared BF16 paging: two waves of two concurrent 1,024-input/64-output requests in each fresh process pair, prefix cache off, fixed pools 32/48. Both ranks save, reload, warm and update their profiles; response hashes match across restart. The leader is signalled and the existing TP stop protocol releases the follower. Startup warms 1,568 / 2,304 existing lines, with no pool enlargement. The profile checkpoint also passes 211 host tests and 15 focused Linux Debug/ReleaseSafe persistence tests; a separate source-53ef CUDA fixture verifies actual bytes/tables and resized pool restart on both GPUs.

Enable with TF_FLASHNEXT_EXL3_PROFILE_DIR. The advisory fingerprint covers routed trellis plus config/index bytes, not every checkpoint tensor; hashing costs about 16.6 seconds per rank at startup in this gate. Counts are decayed once and saved by replacement, not added twice. The separate deep-prefill ring histogram is not persisted; prefill may discard startup residency. Abrupt follower/process kills can lose unsaved updates. These are placement priors, not KV snapshots or elastic budgeting. A separate short-prompt ABBA with drafts off measured 9.0% lower TTFT and 6.5% lower decode latency for warm versus cold persisted profiles (TP=2, graphs and BF16 paging on; 128-token input / 64-token greedy output; identical full output hashes). This narrow result is not evidence for long-context or drafted-serving gains, nor a graph speedup; profile fingerprinting adds about 16.6 seconds of startup per rank. Full methodology and values are in the final report. No independent golden-fidelity, long-prefill-speed or 200k×2 gate is closed by this test.

Latest prefill efficiency candidate

ce10271d starts initial expert copies before CPU member-list construction and coalesces copy-stream dependencies for equal/older release epochs, preserving newer-reader and event-wrap guards. It passes 215 EXL3 host tests, native CUDA host/forward-body compilation, formatting/diff checks and zero new lean findings. Native TP2 BF16 cold/warm/branch/rewind parity now passes for the specified ce10271 workload; the matched ce/babc revision comparison is effectively flat. Independent model fidelity and 200k capacity are not inferred. Previous hardware receipts apply only to their named sources. Copy-grid-only tuning (144/36/36/144, frozen native babc9762, TP2 c2, 1k/256, two waves/server) was 0.3–0.5% slower in both ordered pairs with exact response parity; the default stays unchanged.

Latest measured follow-up

a25fa226 passes a real native TP2 gate with two concurrent 16,424-token histories, MTP, prefill/decode graph replay, BF16 RAM paging and prefix reuse at 4096-row prefill. All eight cold/warm/branch/rewind outputs match the separate cold reference. Demand/initial-prefetch/cross-layer-prefetch copy and requested-byte categories reconcile on both ranks. Native SHA256: 939f5ef182c9bfefb1efe2a3847715c32011b734f703427a52a311ada61df5bf; 222 EXL3 +218 runtime host tests pass, zero new lean findings.

The previous ce10271d revision showed no meaningful speedup against its baseline. Profiling identified expert-transfer volume and many separate projection DMA submissions. The new direct-copy helper combines projections only with proven common pinned-bank provenance and exact adjacency, with guarded fallback; no CUDA arithmetic or graph ownership checks were weakened. The matched ce/a25 code comparison at identical4096-row settings is complete and shows no repeatable gain; one pair has worse decode/ITL. A subsequent matching-target-layer-history candidate passes native cache parity and all16 benchmark hashes/replay checks in four fresh A/B/B/A servers, with about2.1% aggregate decode-service improvement in both screening pairs and only0.2–0.3% effective-prefill improvement. It is integrated with unchanged mandatory loading and graph/BF16-paging guards;227 host tests and native compilation pass. Two pairs do not establish statistical significance or causal graph/eager speedup; diagnostic total transfer volume is effectively unchanged. Tested source and binary, exact ratios and the separate current-default CUDA profile are recorded in the linked measurement comment.

A separate four-fresh-server setting comparison shows 4096 versus forced2048 rows improves TTFT about39% and effective aggregate-prefill about64% on the stated TP2/c2 workload, with exact hashes and actual graph reuse. 4096 is already the engine default; the inherited benchmark harness forced2048. This is a setting comparison, not a new-code gain. The8192 pilot is not accepted because cached-prefill replay failed its unchanged gate. Cross-request graph retention remains disabled pending complete captured-allocation lifetime proof. Full methodology, exact tested commits, units and limits are in the existing measurement comment.

Dense EXL3 prompt-kernel update (22 October; pushed head 5a94359c07c8f716255ff9b0de7a7257dec1432f)

A Blackwell sm120+ mul1/C2 ldmatrix.x4 input-fragment path is now included. The four fresh-server A/B/B/A TP2 screening pairs with 16,384 input / 256 output, c2, graphs/MTP, BF16 RAM paging, 4096-row prefill and pools 32/48 show effective aggregate-prefill +2.27%/+2.23% and TTFT −2.28%/−2.24%; aggregate decode-service +0.77%/+1.23%. All response hashes match. This is a small two-pair screening result, not statistical significance or graph/eager causality. The canonical candidate passes the BF16 paging/cache cold/warm/branch/rewind gate (8/8 responses) and changed-kernel CUDA/resource checks. Full kernel-contract timed out at 300 seconds on both GPUs; broad contract pass is not claimed. Exact methods, hashes, limits and assessment of the reviewer’s cache/page-table pointers are in the latest measurement comment.

[ashhart]

Thanks @BlivionIaG. A note while this is a draft, since it no longer rebases onto main (the first of 119 commits conflicts in 8 files after main's batch-code restructure and Living Weights changes):

• Main's Flash Next CUDA already takes KV pages from the core pool (TF_FLASHNEXT_KV_PAGES), so a second BF16 paging system beside it won't land; paging should extend the core pool. Main also has the CUDA start-up budget and the shared prompt cache, which cover the reserve and much of the cuda_snap rewrite. CUDA graphs for Flash Next are new and welcome once they show a repeatable speedup.
• Rules we'd hold the split PRs to: drafts can't be switched off when experts are offloaded (cuda_exl3_pass.zig, cuda_lanes.zig); TP=2 must be bit-identical to TP=1, not within a tolerance (cuda_exl3_heads.zig); new kernels under the kernel contract with records; far fewer env switches (about 60 new TF_FLASHNEXT_* now); files within 600 lines and lean clean (1375 new findings); no AI co-author trailers, no local user paths (tools/exl3_nsys_step_breakdown.py, the golden JSON); and the 1373-line report kept out of docs/plans.
  Small, separately-measured PRs on current main (graphs first) would be the quickest way in.

[BlivionIaG]

Thanks @BlivionIaG. A note while this is a draft, since it no longer rebases onto main (the first of 119 commits conflicts in 8 files after main's batch-code restructure and Living Weights changes):

• Main's Flash Next CUDA already takes KV pages from the core pool (TF_FLASHNEXT_KV_PAGES), so a second BF16 paging system beside it won't land; paging should extend the core pool. Main also has the CUDA start-up budget and the shared prompt cache, which cover the reserve and much of the cuda_snap rewrite. CUDA graphs for Flash Next are new and welcome once they show a repeatable speedup.
• Rules we'd hold the split PRs to: drafts can't be switched off when experts are offloaded (cuda_exl3_pass.zig, cuda_lanes.zig); TP=2 must be bit-identical to TP=1, not within a tolerance (cuda_exl3_heads.zig); new kernels under the kernel contract with records; far fewer env switches (about 60 new TF_FLASHNEXT_* now); files within 600 lines and lean clean (1375 new findings); no AI co-author trailers, no local user paths (tools/exl3_nsys_step_breakdown.py, the golden JSON); and the 1373-line report kept out of docs/plans.
  Small, separately-measured PRs on current main (graphs first) would be the quickest way in.

Hi this is indeed meant to be a draft that will be split into pr later, I still have some drafting related things to test and improve.
I have started to reduce the amount of env vars, it's mainly there in this draft to allow tuning and test, final PRs will trim them down to the essential.

[ashhart]

Thanks @BlivionIaG. The per-part reviews are now on #1441 to #1446, and #1447 has its own. #1419 says that chain
replaces this draft, and today's re-push of the chain only added Windows build guards, so those reviews stand. This
branch no longer merges with main (17 files conflict), so we suggest closing it and continuing on the chain.