Skip to content

MiMo-V2.6-Flash on CUDA: EXL3, DFlash, two GPUs - #728

Open
ajensenwaud wants to merge 12 commits into
ashhart:mainfrom
ajensenwaud:mimo-v2
Open

ajensenwaud wants to merge 12 commits into
ashhart:mainfrom
ajensenwaud:mimo-v2

Conversation

@ajensenwaud

@ajensenwaud ajensenwaud commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

MiMo-V2.6-Flash on the native CUDA engine, on two DGX Sparks as tensor-parallel ranks, with the pack's DFlash drafter. Proposal and discussion: #725.

Admission is what was validated: GB10 (sm_121) at two ranks. open refuses any other GPU by name (MimoV2RequiresSm121), as DeepSeek-V4.1's kernel load does, and refuses one rank or any other world (MimoV2RequiresTwoRanks), at the weight share as well as at open. tf-mimo refuses both too. Host tests cover each refusal. One rank, and the 2.27 bpw pack that fits one GB10, wait for their own receipts.

Based on current main (15a0b2a9). Twelve commits:

  1. c656d4b3 lanes: a geometric draft depth whose round cost grows with context. About 100 lines in core/lanes, plus a host test. It is off unless a family sets geometric_depth, so other families keep the per-depth rule unchanged. It can become its own PR if you prefer.
  2. 849f2dbc exl3 cuda: the common prompt path takes 3-bit mul1 experts over fp16 rows. The prompt kernels in exl3_prompt.cuh / exl3_experts.cu take the half-bit width and codebook as template arguments: tf_exl3_prompt_gu / _down stay the 4-bit mcg instance, _k6_mul1 is the 3-bit mul1 one. exl3_prompt.zig picks the instance by layer and takes the rows' dtype; GLM's call passes .bf16. On sm_121 the 4-bit gate/up SASS is unchanged and the down kernel's differs in one mask constant (the same value at 256 threads).
  3. e4e0f5ee MiMo-V2.6-Flash on CUDA: native kernels on the shared EXL3 operators, DFlash, two GPUs, prompt cache.
    • No captured Triton set and no EXL3 copies: cuda/mimo_exl3.zig and the vendored kernels are gone. Projections and decode experts run on cuda.exl3.Linear and cuda.exl3_experts, prompt experts on the common prompt kernels (commit 2).
    • The family's kernels (zig/kernels/cuda/mimo_v2): ops.cu (RMS norms, SwiGLU, fp32 router and top 8, the prompt experts' slot sum), attention.cu (queries and caches, chunked decode attention with sinks, and a tensor-core prompt kernel for the 9 full layers), glue.cu, dflash.cu.
    • Every prompt token runs on the prompt path, in chunks that end at the prompt cache's marks; a row's bits do not depend on its chunk (the prompt kernel's 64-key tiles start at 0), so a resumed prompt reads as a fresh one. The 256-token decode/prompt cut is gone.
    • Kept prompt states (kept.zig): the slot's full-layer rows stay pinned, the sliding and drafter windows are copied (33 MiB a rank); a resume binds to the state's slot or copies the rows into a free one. Rank 1 follows every cache step over the plan link.
    • Build: zig/build/mimo.zig (module, tf-mimo, test-mimo), laid out like kolibri.zig; host tests in cuda_family_host.zig beside Dory's, so zig/build/cuda.zig stays at 600 lines. native/cuda.zig: one line, the registry entry; admission uses main's weightShare hook.
    • Recipe page: docs/recipes/mimo-v2.6-flash.md.
  4. cadb88d8 kernel contract: every MiMo-V2 kernel through gates 1-4, its launches for gate 5. contract_mimo*.zig (28 kernel forms against float64 or their stated rules, at the served shapes, launched as the family launches them) and mimo_cuda_records.zig (sm_121, nvcc 13.3, including the two mul1 prompt instances).
  5. 289fe1da MiMo-V2 decode: attention rows on tensor cores, norms in registers, a shallower top-8 tree. Brings one-stream decode back level with the captured-set build.
  6. c28dc32b MiMo-V2: a refused load frees only the layers it filled.
  7. 98562b40 MiMo-V2: a bad TF_TP_ refuses the weight share instead of sizing one GPU.*
  8. 0e068f33 tf-mimo follow: refuse one rank before reading a session.
  9. 5d01f1da MiMo-V2.6-Flash recipe: licences and sources of the model, the pack and the drafter.
  10. 48468450 MiMo-V2: refuse every GPU but GB10 (sm_121) by name. The check is the first line of open, and of tf-mimo. A host test runs open on a fake driver at sm_80 to sm_120, all refused, and at sm_121, which passes the check.
  11. 799c87f3 MiMo-V2: refuse one rank by name; only two were validated. rankView, the weight share and open refuse TF_TP_WORLD unset, 1, or anything but 2. The recipe drops the one-GPU and 2.27 bpw serving claims.
  12. 021eb4c3 MiMo-V2.6-Flash recipe: full revisions of the model and both packs.
    • Model XiaomiMiMo/MiMo-V2.6-Flash-RL at 5711b268169967567844e1e560e8a3966da959b1, the revision both packs were converted from and the one validated. The repository's head since then, dcd3fd83, changes only the vision preprocessor_config.json.
    • Pack ajensenwaud/MiMo-V2.6-Flash-RL-exl3-3.0bpw at 139d9a4e1eb32193890036fdddd521fe37033a30.
    • Pack benthecarman/MiMo-V2.6-Flash-RL-exl3 (2.27 bpw, the drafter's source) at 5ae87b830eb950c705967aea6130601a8e1f74a5.

The grouped experts load-ahead that was commit 5 is now its own PR, #1505, with a GLM-5.3-Flash A/B. The gate below was measured with it in: without it MiMo's teacher-forced logits are byte-equal, so the identity and quality results stand, but decode tok/s is lower by the grouped kernel's difference (about 109 against 96.6 µs a launch) until #1505 lands.

The 8-bit KV cache stays out (fp16 only), as you asked.

Gate (the previous head with the load-ahead, two DGX Sparks)

GB10 (sm_121), checkpoint ajensenwaud/MiMo-V2.6-Flash-RL-exl3-3.0bpw, --parallel 4 --context 608256 --top-k 20 --top-p 0.95, RoCE between ranks.

  • Kernel contract (full CUDA runner on GB10, this head): SUMMARY 446 kernels: 446 pass gates 1-4, 0 fail, 5 known findings; 312 controls: 312 red as planted, 0 missed; 214 limit problems. None of the limit problems are MiMo's or the shared experts': every MiMo launch and the three tf_exl3_grouped launches match their nvcc 13.3 records live on sm_121.

  • Lane contracts (greedy, by the server's token_sha): drafted == "draft": false 6/6 (prompts of 29 to 10,806 tokens); concurrent == solo 24/24 (three groups of four, drafted and plain).

  • Resumed == fresh (prompt cache on, 3.8 GiB): 117/117 replies identical to the cold serve's, 63 resumed (102,602 of 250,302 prompt tokens from the cache), including resume points at 63/64/65, 127/128/129, 255/256/257, 1,023/1,024/1,025, 2,047/2,048/2,049 and 4,095/4,096/4,097 tokens (18/18 resumed at the mark). With "draft": false, 117/117 identical (16 resumed).

  • Teacher-forced (tf-mimo m1): against exllamav3 on the same pack, see Fidelity against exllamav3 below. Prompt path: one chunk == two chunks on the last row 8/8; against the decode rows, KL mean 0.0073 (different kernels, so not bit-equal).

  • Quality (cache on): GSM8K 477/500, MMLU 190/228, passkey 10/10 (32K and 120K, five depths each).

  • Prefill, tok/s (spark-bench depth, cold):

    8K 32K 64K
    this PR 1,106 1,037 962
    first native build of this rework (decode experts on the prompt path) 523 406 310

    The recipe's release build reads 16K at 1,309 and 64K at 1,183.

  • Decode, tok/s (spark-bench --preset full, 512 tokens, thinking off, temperature 0; chat / code / agent / report):

    1 stream 4 streams, aggregate
    this PR 42.0 / 65.8 / 54.2 / 45.8 75.6 / 108.1 / 89.1 / 82.6
    recipe release build (16-bit, an earlier session) 49.8 / 64.3 / 55.7 / 53.4 81.6 / 115.6 / 90.6 / 80.6

    The bench tags each prompt at random, so one-stream figures move with the text. On fixed prompts (512 tokens, one stream, temperature 0, both builds served the same day), tok/s:

    "draft": false drafted
    this PR 41.5 / 41.3 / 41.0 / 41.3 47.8 / 59.8 / 61.2 / 46.9
    captured-set build 42.7 / 43.0 / 43.5 / 43.0 47.3 / 64.3 / 62.5 / 49.9

Fidelity against exllamav3

The reference is exllamav3 1.6.0, the build that wrote the 3.0 bpw pack, run on the same pack. It is teacher-forced over the same 1,905 positions: 8 texts, 254 prompt rows and 1,651 generated rows. It runs its no-cache path (flash_attn_nc), streamed one module at a time on one GB10, and its logits are its default fp16 head output. Ours come from tf-mimo m1 at TP=2 on this PR's kernels.

top token = exllamav3's KL(exllamav3 ‖ engine), nats: mean / p50 / p99 / max teacher-token NLL
this PR 1,879/1,905: 26 differ, plus 2 exact ties 0.0043 / 0.00031 / 0.053 / 0.34 0.4957
the earlier captured-set build 1,881/1,905: 24 differ, plus the same 2 ties 0.0050 / 0.00032 / 0.055 / 1.41 0.4961
exllamav3 against itself (prefix run vs full run, 2,719 rows) 31 flips (1.1%) 0.0019 mean
  • Counting. In the table, a position counts as a match when the two engines share a top token. An exact tie is a position where exllamav3's top two logits are equal in its fp16 output and our top token is one of the two. Both ties are in the short text, at steps 48 and 59.
  • Why 1,878 and 28 before. The review reply counted those two ties two different ways:
    • Its count, 1,878 (27 misses), broke a tie by the lower token id. That matched ours at step 48 and not at step 59.
    • Its table's sort broke both ties against ours, so the table listed 28.
    • Both numbers describe the same positions: the 26 that differ plus the 2 ties.
  • The 26 that differ. At 25 of them this PR ranks exllamav3's top token second, and third once (edit, step 50). exllamav3's lead over its second token there runs from 0.016 to 0.31 logits, median 0.13. The full list is below.
  • exllamav3's fp32 head as the reference. Taking exllamav3's fp32 head output instead removes the ties: 1,876/1,905 (29 differ, no ties), with KL 0.0043 / 0.00030 / 0.054 / 0.34.
  • Teacher-token NLL covers the 1,897 rows that have a next token; exllamav3's own is 0.4984.
Every position where the top tokens differ, with exllamav3's margin

Margin = exllamav3's top-1 minus top-2 (logit, prob). "rank" = where this PR puts exllamav3's top-1. Token = idtext. Next = the teacher token at that position.

text step KL next exllamav3 top-1 (logit, p) exllamav3 margin (logit, p) this PR's top-1 (p) this PR's p of exllamav3 top-1, rank exllamav3 p of this PR's top-1
en_story 18 (prompt) 0.000636 1943 message 25382 mysterious (16.8438, 0.1630) 0.0312, 0.0050 14888 strange (0.1622) 0.1615, 2 0.1580
en_story 37 0.00272 5001 coming 5001 coming (18.8438, 0.2429) 0.0781, 0.0183 13458 storm (0.2410) 0.2324, 2 0.2246
en_story 40 0.0059 576 The 1260 He (18.0312, 0.3508) 0.1875, 0.0600 576 The (0.3144) 0.3092, 2 0.2908
en_story 87 0.00398 45035 gratitude 5775 wonder (22.4688, 0.4402) 0.0625, 0.0267 45035 gratitude (0.4338) 0.4308, 2 0.4135
en_story 95 0.0598 264 a 264 a (17.0625, 0.1783) 0.2500, 0.0394 504 from (0.2353) 0.1526, 2 0.1389
en_story 161 0.0255 326 l 1943 message (17.3125, 0.3852) 0.1719, 0.0608 326 l (0.3987) 0.3000, 2 0.3244
en_story 177 0.00457 1035 would 1035 would (23.4062, 0.4174) 0.1562, 0.0604 44419 climbed (0.3757) 0.3751, 2 0.3570
en_story 216 0.0362 73632 lodged 73632 lodged (22.2188, 0.2552) 0.1094, 0.0264 4279 half (0.2813) 0.2623, 2 0.2287
en_story 263 0.00522 11, 304 in (25.7188, 0.4495) 0.0938, 0.0402 11, (0.4278) 0.4165, 2 0.4093
zh_explain 63 0.0199 101068无法 53153不能 (21.8906, 0.2774) 0.0312, 0.0085 101068无法 (0.3350) 0.2731, 2 0.2689
zh_explain 119 0.0325 117184粒子 100369两个 (22.5000, 0.4304) 0.3125, 0.1155 117184粒子 (0.4073) 0.3570, 2 0.3149
zh_explain 133 0.0398 1773。 106293的状态 (25.8438, 0.3952) 0.1250, 0.0464 1773。 (0.4682) 0.2807, 2 0.3488
code_py 145 0.0289 11, 13. (23.8594, 0.3009) 0.0625, 0.0182 11, (0.3421) 0.2576, 2 0.2783
math 60 0.0157 18611 \n\n 6771 Let (14.8438, 0.0822) 0.2578, 0.0187 18611 \n\n (0.0781) 0.0755, 2 0.0635
tool_json 22 (prompt) 0.00878 25: 510:\n (17.7812, 0.2198) 0.0781, 0.0165 25: (0.2269) 0.2251, 2 0.2033
edit 49 0.023 334** 88757Rename (16.1406, 0.1151) 0.2188, 0.0226 334** (0.0988) 0.0956, 2 0.0925
edit 50 0.035 16141Answer 18896Expected (16.0625, 0.0680) 0.1875, 0.0116 16141Answer (0.0761) 0.0622, 3 0.0563
edit 72 0.0122 1760 count 1760 count (24.5781, 0.5373) 0.1562, 0.0777 13195 cnt (0.5366) 0.4598, 2 0.4596
edit 279 0.0293 25: 982):\n (30.0312, 0.5592) 0.2656, 0.1305 25: (0.5490) 0.4393, 2 0.4288
list 106 0.00615 576 The 27360 Rot (22.9062, 0.4174) 0.1406, 0.0548 576 The (0.4054) 0.3734, 2 0.3626
short 0 (prompt) 0.029 0! 358 I (8.7500, 0.0129) 0.1953, 0.0023 11, (0.0173) 0.0142, 2 0.0075
short 10 0.00149 2371-g 2371-g (17.5156, 0.0956) 0.0156, 0.0015 1448-m (0.0996) 0.0948, 2 0.0941
short 16 0.00917 12- 5338First (14.1562, 0.0263) 0.0938, 0.0024 12- (0.0298) 0.0213, 2 0.0240
short 54 0.00165 358 I 358 I (20.8594, 0.3520) 0.0625, 0.0213 576 The (0.3544) 0.3427, 2 0.3307
short 208 0.00613 12440 correctly 12440 correctly (24.6719, 0.4150) 0.0625, 0.0251 6915 fine (0.4299) 0.4114, 2 0.3898
short 237 0.00629 4460 trying 4460 trying (17.5000, 0.1722) 0.1562, 0.0249 537 not (0.1535) 0.1494, 2 0.1473

The two exact ties, both in the short text. exllamav3's top two are equal, and our top token is one of them:

step KL next exllamav3's tied pair (logit) this PR's top-1 (p)
48 0.00195 21556 Sprite 21556 Sprite and 51548 Collision (16.4219) 21556 Sprite (0.1737)
59 0.00176 17051 sprite 2781 player and 17051 sprite (18.4219) 17051 sprite (0.3132)

Host (this head, on main 15a0b2a9): test-mimo, test-kernel-contract, test-native-options-host, test-kernel-symbols-host, test-shared-lane-regressions-host and test-tp-cli pass; tf-mimo, native, compile-config-server and compile-cuda-tests build. test-cuda-host passes except two core prompt-cache spill tests, rows_async_test ("the job permit bounds metadata ...") and rows_close_test ("spill-only flushes more than eight RAM states ..."). They fail the same way on main 15a0b2a9 in the same build container. zig fmt --check and git diff --check are clean; tools/zig/lean_check.py finds 428 problems, the same as main.

Next

  • Decode: about 0.8 ms a token left against the captured-set build, mostly the norms and the top 8.
  • Prefill: the dense projections still run 128 rows a call.
  • One rank: refused for now. To admit it needs its own equalities and an exllamav3 run on the 2.27 bpw pack, which fits one GB10. TP=1 against TP=2 is expected to differ in the last bits, so that would be a closeness check (see the review reply).

Notes for review

satindergrewal added a commit to satindergrewal/TensorFold that referenced this pull request Oct 13, 2026
…3 experts

Checked ash hart PRs open+closed (read-only). Direct hits on our
problems, all OPEN, branch BHCC2025/TensorFold exl3-cuda-experts:
- ashhart#741 EXL3 routed experts behind experts.Grouped: gate+up+SwiGLU
  ONE kernel, down on pair rows, each expert's rows rotated by its
  own suh INLINE (register-resident rot for 1-2 pair chunks - the
  problem we had deemed impossible tile-incrementally). Rides on
  grouped.zig's Plan (in our tree). 297cu+124zig lines, with an
  exactness test harness (tf-cuda-test exl3-experts) - 367,872
  values within FP64 bounds, alone==grouped byte-equal.
- ashhart#740 EXL3 projections behind qlinear (the stacked base).
- ashhart#718 branch-free trellis window extraction (perf).
- ashhart#729 share decoded weights across prefill rows (perf).
- ashhart#728 MiMo-V2.6 CUDA EXL3 two-GPU (another full CUDA EXL3 family).
- ashhart#765 GLM CUDA multi-rank doorbells; ashhart#770/777/768 nemotron CUDA
  single-launch shared rounds + batched MTP drafting (patterns for
  our later concurrency work).
PLAN: vendored exl3_experts.cu/.zig + exl3.cuh into our tree (this
commit, unwired). Next: build wiring (fatbin + mod), port moeBlock
to Grouped.decode with the router->plan plumbing from grouped.zig,
run their exl3-experts test on our checkpoint, parity, bench. Our
own grouped kernels (stage-1 parity) stay as fallback.
@ashhart

ashhart commented Oct 17, 2026

Copy link
Copy Markdown
Owner

Thank you, @ajensenwaud, the shared-lane DFlash path and explicit optional FP8 policy give the MiMo work useful boundaries.

We checked 83fd3f2. native.zig:374 still requires captured kernels, :379 hides bad TP settings with catch null, and rankWeights:412-416 divides bytes despite replicated embeddings/drafter. We can repair refusal and placement accounting while moving execution onto owned kernels and the shared EXL3 module.

We will extract the depth rule and infrastructure first, then the family with prompt reuse and final-source numerical/peer/memory gates; 16-bit cache remains the default. Please provide the original/converted checkpoint and vision inventory with the pinned workload so image obligations and the final reference gate are explicit. Earlier pre-rebase receipts remain attached to their source.

Posted by Ash's AI Assistant Synthena

@ashhart

ashhart commented Oct 18, 2026

Copy link
Copy Markdown
Owner

Thank you, @ajensenwaud, for #728 and the 24 concurrent-versus-solo checks on two GPUs, with @jayleaton and the common EXL3 contributors. MiMo-V2.6-Flash is not included in 1.1.1. Its current CUDA path depends on a Python-captured Triton set that TensorFold does not ship; without TENSORFOLD_CUDA_KERNELS it refuses to start. Native kernels for its eight Triton launches come next, followed by full-model fidelity, drafted/plain, resumed/fresh, concurrent/solo, two-rank failure and compiler/resource gates on the integrated EXL3/DFlash path. This is held for a later 1.1.x; captured-set receipts do not qualify the new native path.

Posted by Ash's AI Assistant Synthena

@ashhart

ashhart commented Oct 22, 2026

Copy link
Copy Markdown
Owner

Thanks @ajensenwaud for MiMo-V2.6-Flash on CUDA. Main has moved a long way since this branch (it is stacked on #703), so rather than ask you to rebase, we'll port its MiMo-specific commits onto current main ourselves, keeping you as their author and saying in each message what was adapted. Then it goes to the GB10 kernel contract and served-equality runs on one and two GB10s before it lands. If you are partway through a change here, please say so, so we don't collide.

- Synthema, Ash's AI Assistant.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Thank you. Very keen to contribute. I know it's moving fast. I can rebase and recommit if that helps?

@ashhart

ashhart commented Oct 22, 2026

Copy link
Copy Markdown
Owner

@ajensenwaud, a question from the port before we combine the MiMo paths with main's shared EXL3 operators. In the branch, the arithmetic a row gets depends on how many rows or tokens arrive together: the router picks its tile by row count (BM 16 up to 16 rows, then larger tiles with BK 256 to 64 and BLOCK_E 32 to 64, zig/src/families/mimo_v2/forward.zig:404 and :456, with the same switch in prompt.zig:381); prompt projections use a separate GEMM with BK 32 while decode uses the EXL3 linear (prompt.zig:204 against forward.zig:268); experts switch to per-expert GEMMs from 3,072 rows (prompt.zig:389); and prompts of 256 tokens or more take the prompt path (native.zig:218). Our rule is that a row's result never depends on batch size or prompt chunking, so we need one arithmetic per operation, or a fixed-order equivalent with checks at 16/17 router rows, 255/256 prompt tokens and 3,071 to 3,073 expert rows. Would you be happy for main's shared EXL3 operators to define the reference, with the oracle fixtures regenerated against them, or would you rather supply fixed-order versions of the MiMo paths? The rest of the port (the geometric draft depth, which stays off unless a family opts in, and the FP8 KV option) is in hand on our side.

- Synthema, Ash's AI Assistant.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Thanks Synthema, and thanks for porting rather than asking for a rebase. Answers to both open questions.

1. Reference arithmetic. Please let main's shared EXL3 operators define the reference and regenerate the oracle fixtures against them. One arithmetic per operation is the right contract: our 24/24 concurrent-versus-solo result was measured, not guaranteed by construction, and I'd rather not keep a parallel MiMo numeric stack. Nothing MiMo-specific needs to survive in the linears or experts. The router tile switch came from the Python engine's shared._tile; one fixed tile for every row count is fine with us (the router is 4,096 x 256, 2 MB of weights a layer).

Two requests so the port keeps what the branch was for:

  • Speed. The thresholds exist for measured reasons, so please keep them as performance choices once the paths are bit-equal. Prompts under 256 tokens on decode rows took 4-stream code from 94-105 to 117-120 tok/s aggregate; grouped experts up to 3,072 rows cut a 512-row prompt from 2.62 s to 0.74 s. Bar at 83fd3f2 on two GB10s: 1 stream code 76.3 tok/s with DFlash, 4 streams 113.4; prefill 4K/16K/64K/128K 772 / 1,317 / 1,210 / 1,066 tok/s.
  • Quality. Fixtures regenerated against main prove self-consistency, not fidelity. Please keep a full-model gate against an external reference too: ours is the pinned 500-item GSM8K and MMLU-Pro sets (greedy, thinking off) plus per-position KL against exllamav3 on the same pack.

I can run both on two Sparks against your port branch and post the receipts; point me at it.

2. Checkpoint, vision inventory, workload (your 17 Oct note)

  • Original: XiaomiMiMo/MiMo-V2.6-Flash-RL @ 5711b268 (166 GB, 65 shards). Multimodal content: visual.* (28 blocks), speech_embeddings.*, audio_tokenizer/ (1.9 GB).
  • Converted: ajensenwaud/MiMo-V2.6-Flash-RL-exl3-3.0bpw @ f2d8d100, exllamav3 v1.6.0, -b 3.0 -hb 8 -mb 4 -hq -cr 250 -cc 2048 -cb mul1. Text only: no vision or audio tensors. DFlash drafter at 4 bpw from benthecarman/MiMo-V2.6-Flash-RL-exl3 @ 5ae87b83.
  • Image obligations: none claimed by this PR. Vision would be a separate PR; until then I'd suggest the family refuses image content rather than silently dropping it.
  • Pinned workload: tools/upstream_contracts.py in the recipe repo (six prompts covering all three prompt paths, greedy, 256 tokens, thinking off), results in results/zig-z7/upstream/upstream-contracts.json; eval items in results/evals/items-*.json.

Nothing in progress on the branch, so no collision risk.

@ajensenwaud ajensenwaud changed the title MiMo-V2.6-Flash on CUDA: EXL3, DFlash, two GPUs (stacked on #703) MiMo-V2.6-Flash on CUDA: EXL3, DFlash, two GPUs Oct 25, 2026
@ajensenwaud

ajensenwaud commented Oct 25, 2026 •

Copy link
Copy Markdown
Contributor Author

Thanks Synthema. To avoid duplicated effort: I've rebased #728 onto current main (2ff01910), so it no longer depends on #703. It is now three commits: the geometric depth rule (off by default), the MiMo family, and the 8-bit cache. The 8-bit cache is now opt-in through main's existing --kv-cache fp8, since main refuses --policy on every CUDA family. native/cuda.zig is down to the one registry line. The family's build sits in zig/build/mimo.zig, like Kolibri's, and its EXL3 bindings are cuda/mimo_exl3.zig, next to main's shared module.

I also fixed two points from your earlier review. Admission now uses main's weightShare hook, counting the embeddings and drafter whole on each rank. A bad TF_TP_* setting now refuses to start instead of being silently ignored.

test-cuda-host, test-mimo, the shared lanes tests, zig fmt and lean_check pass (no new problems), and CI is green. Nothing else is in progress on my side, and I'm happy for you to move it onto the shared EXL3 operators as discussed.

@ashhart ashhart left a comment

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks @ajensenwaud. Commit 026bd20 adds an fp8 KV cache (--kv-cache fp8) that the PR itself notes is not bit-equal; TensorFold takes no lossy formats, even opt-in, so that commit has to go. Main has also moved MiMo-only kernels and geometric depth to a separate MiMo track (1964746), so the MiMo EXL3 kernels here, copies beside the shared cuda/exl3.zig, need to fit that plan and be registered with the kernel contract. It's about 150 commits behind and conflicts in docs/recipes/README.md and native/cuda.zig. After a rebase we'd check on two GB10s that a resumed conversation equals a fresh run across the 256-token boundary, and TP=2 equals TP=1.

@ashhart

ashhart commented Oct 26, 2026

Copy link
Copy Markdown
Owner

Review against main e9b69ed. Commit 026bd20 adds an 8-bit key/value cache (--kv-cache fp8) that the PR itself notes is not bit-equal. TensorFold doesn't take lossy formats, even opt-in, so that commit needs to come out. Beyond that:

  • Main moved MiMo-only kernels and geometric depth to the separate MiMo track (1964746); whether MiMo lands is a maintainer decision about that track.
  • The MiMo EXL3 kernels are copies beside main's shared cuda/exl3.zig, and none are registered with the kernel contract. Please reuse the shared EXL3 path where it fits and register every launch.
  • The branch is about 150 commits behind and conflicts in docs/recipes/README.md and native/cuda.zig.
    Geometric draft depth is off by default, which is right. Device plan once rebased (two GB10s): a resumed conversation equal to a fresh run across the 256-token boundary (prompts under 256 tokens run as decode rows), and TP=2 equal to TP=1.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Yes in progress. Working on it.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Thanks Ash. Rebased onto current main (8c36359d, no conflicts), and the 8-bit KV commit is out: the cache is fp16 only. The branch is now two commits, the geometric depth rule (off by default) and the family. On the family with the 8-bit cache removed: drafted == plain 6/6 and concurrent == solo 24/24 on two GB10s. Host suites, fmt and lean are clean.

Next on this branch, from your review: native CUDA kernels for the eleven captured Triton launches, each registered with the kernel contract; then the EXL3 copies moved onto cuda.exl3 / exl3_experts and main's prompt kernels, as in #1295. After that, the device runs you listed: resumed vs fresh across the 256-token boundary, and TP=2 vs TP=1 on the 2.27 bpw pack, which fits one GB10. I understand landing is your call on the MiMo track.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Pushed the rework your review asked for, rebased on main 8c36359d:

  1. No captured Triton set and no EXL3 copies. MiMo runs on cuda.exl3.Linear and exl3_experts, and its prompt experts on the common prompt kernels, which gain a 3-bit mul1 instance in their own small commit (GLM's 4-bit instance is unchanged).
  2. Every MiMo kernel is registered with the kernel contract (28/28 pass gates 1-4 on GB10) and has its sm_121 launch records.
  3. The family keeps prompt states. Resumed == fresh holds on 117/117 replies with 63 resumes, including resume points at 255/256/257 and the other chunk and window edges. Every prompt token now runs on the prompt path, so the 256-token cut is gone.

Quality held: GSM8K 476/500, MMLU 190/228, passkey 10/10. Prefill reads 64K at 924 tok/s. Numbers and the full gate are in the description; decode is the next thing I'm working on. Happy to run anything else on our two Sparks.

@ajensenwaud
ajensenwaud marked this pull request as ready for review October 27, 2026 08:42
@ajensenwaud

Copy link
Copy Markdown
Contributor Author

This one is ready for review.

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Rebased on main 3f144ccf (the only conflict was the contract's image list). Two new commits on top of the four you reviewed, both for decode speed:

  1. exl3 cuda: the grouped experts kernel loads a tile ahead. warp_tiles now fetches the next k tile's words while it decodes the current one, as exllamav3's experts kernel does at pf = 1. The decode, the mma chain and the warp sums are unchanged, and the SASS arithmetic is the same, so the bits are the same. On GB10 MiMo's routed experts go from 110 to 97 µs a layer step; GLM's experts run the same kernel. Registers rise to 145-168. I measured the records with nvcc 13.0 and 13.3 on all five SMs. I have no 13.2, so I dropped its one stale grouped_1 override rather than guess; it needs re-recording on a 13.2 box.
  2. MiMo-V2 decode: attention rows on tensor cores, norms in registers, a shallower top-8 tree.
    • mimo_attn_rows runs a chunk's 8 or 16 query heads as one m16 tile on the prompt kernel's loads and mma, 64-key tiles at multiples of 64, and the four warps' partials fold in warp order. A row's bits still depend only on its own position and window. p is in fp16 for o, under the same contract allowance as the prompt kernel.
    • mimo_rms keeps its values in registers, and mimo_topk runs the last five levels of its tree as warp shuffles over the same pairs. Both give the same bits as before.

Where the time went (nsys, one stream, plain decode, kernel time per call):

captured-set build previous head these commits
sliding attention rows 8.1 µs 72.4 µs 8.5 µs
norms 2.4 µs 11.2 µs 5.7 µs
top 8 3.7 µs 8.9 µs 6.9 µs
routed experts (grouped) 101.7 µs 109.3 µs 96.6 µs

Decode, one stream, 512 tokens, fixed prompts, temperature 0, tok/s (chat / code / agent / report):

captured-set build previous head these commits
"draft": false 42.7 / 43.0 / 43.5 / 43.0 34.2 / 34.2 / 33.9 / 34.1 41.6 / 41.7 / 41.6 / 41.7
drafted 47.3 / 64.3 / 62.5 / 49.9 38.8 / 54.2 / 48.6 / 39.0 48.8 / 64.1 / 60.6 / 48.1

Decode is back level with the captured-set build, about 0.8 ms a token behind it, with every kernel native or shared. These figures are from one session before the rebase; the rebased head's are in the description.

Gate on this head (two Sparks), in the description: the full contract runner passes 446/446 with every control red, and MiMo's and the grouped launches match their records live; resumed == fresh 117/117 drafted and plain; GSM8K 477/500, MMLU 190/228, passkey 10/10.

@ashhart

ashhart commented Oct 27, 2026

Copy link
Copy Markdown
Owner

Thanks @ajensenwaud, this rework is a big step. The captured Triton set and the EXL3 copies are gone, and the fp8 cache is out. Every MiMo kernel is in the contract with records, and resumed equals fresh across the old 256-token boundary. We rebased it onto main b30d5b9d3 (clean). The CUDA host suites, test-mimo and both Linux compiles pass, and lean adds nothing.

A few things before our GB10 run:

  1. Shared kernels. Commit d93bf26e9 changes warp_tiles, the body of tf_exl3_grouped_0/1/2, which GLM and the other EXL3 families run. It raises them from 127-128 to 145-168 registers, depending on the SM. It also drops the nvcc 13.2 override, so 13.0's numbers stand in for 13.2, which the contract doc rules out. Could it go in its own PR with a GLM A/B (identical tokens, and tok/s)? For the templated prompt kernels in c0796d262, please attach a cuobjdump -sass comparison against main for tf_exl3_prompt_gu/_down on every SM in the fatbin.
  2. Fidelity. The teacher-forced run compares with the earlier captured-set build, which is our own engine. Please compare against exllamav3 on the same pack (the per-position KL you offered), and list each top-token miss with its contenders and margins.
  3. load.zig:220. layers is allocated uninitialised, and errdefer w.deinit() reads every layer's experts. A refusal before the layers are filled, such as EmbedNotBf16, frees garbage. Also, weightShare (native.zig:490) still treats a bad TF_TP_* as one GPU.
  4. Licence. Please state the licences of MiMo-V2.6-Flash and the DFlash drafter in the recipe, and name where the pack and the drafter come from.

The geometric depth is off unless a family opts in, which is right. Main had moved it to the MiMo track (1964746), so where it lands is our decision. On our pair we'll then run the contract, a GLM control, drafted, concurrent and resumed at TP=2, and TP=1 on the 2.27 bpw pack. TP=2 adds the two ranks' fp32 partials in rank order. Does TP=1 split its sums the same way, or should we expect TP=2 and TP=1 to differ in the last bits?

@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Thanks @ashhart. All four points are addressed; the branch is rebased on main 41abcf40 (the only conflict was the contract's image list again).

  1. Shared kernels.

    • The experts load-ahead is out of this PR and in its own, exl3 cuda: the grouped experts kernel loads a tile ahead #1505, with the GLM-5.3-Flash A/B you asked for: two GB10s, TP=2, main against main plus the commit, the same flags. 30/30 outputs identical by token SHA (drafted and "draft": false, solo and 4 streams, every speed run), decode +0.4-0.7% at 1 stream. Its records keep 13.2 pending, as the contract doc asks.
    • Without it, this PR's MiMo logits are byte-equal (the teacher-forced dumps with and without it match), so nothing here depends on it.
    • cuobjdump -sass for tf_exl3_prompt_gu/_down, main against this branch, on every SM in the fatbin plus sm_89, with nvcc 13.3 and 13.0: identical everywhere except one constant in _down on sm_120/121 (0xe0 -> 0x3e0 in a mask of threadIdx.x, which is below 256 there, so the result is the same). Full table below.
  2. Fidelity against exllamav3. exllamav3 1.6.0 (the build that wrote the 3.0 bpw pack), teacher-forced over the same 1,905 positions on the same pack, one module at a time on one GB10, fp32 logits of its head:

    top token = exllamav3's KL(exllamav3 ‖ ours), nats: mean / p50 / p99 / max
    this PR 1,878/1,905 0.0043 / 0.00031 / 0.053 / 0.34
    the captured-set build (our earlier reference) 1,880/1,905 0.0050 / 0.00032 / 0.055 / 1.41
    exllamav3 against itself (prefix run vs full run, 2,719 rows) 31 flips (1.1%) 0.0019 mean
  3. load.zig and weightShare.

    • Every layer now starts blank, without experts, so a refusal before the layers are filled frees nothing it didn't allocate. A host test refuses an fp16 embedding under the testing allocator; without the blank start it crashes in deinit.
    • weightShare returns BadSetting for a malformed TF_TP_*, as open does, and the admission reports it. TF_TP_WORLD unset is still one GPU; GLM's and Nemotron's shares pass through unchanged. Host test: unset, six malformed settings, and a two-rank share.
    • While there: tf-mimo follow now refuses a one-rank world before it reads a session.
  4. Licences. The recipe now has a licences and sources section: the model is XiaomiMiMo/MiMo-V2.6-Flash-RL, MIT; our 3.0 bpw pack ajensenwaud/MiMo-V2.6-Flash-RL-exl3-3.0bpw is MIT as its source, quantised with exllamav3 v1.6.0 (MIT) at the flags given; benthecarman's 2.27 bpw pack is MIT; the DFlash drafter is Xiaomi's, in the model repository's dflash/ folder under its MIT licence, and DFlash itself is z-lab's (MIT).

TP=1 against TP=2: expect them to differ in the last bits.

  • At TP=2, o_proj, the expert down projections and the dense down_proj each finish an fp32 partial over half of K on each rank; those are gathered as bytes and added once, rank 0 + rank 1.
  • At TP=1 one kernel reduces the full K, with different split boundaries and one output Hadamard finish, and the expert slot sum runs over full-width outputs.
  • q_proj's EXL3 plan is chosen from the per-rank shape, (4, 4) at TP=2 against (1, 8) at TP=1, so its sum order changes too.
  • Bit-equal TP=1 would need a mode that runs those projections as two virtual ranks and adds the halves the same way. Short of that, TP=1 against TP=2 is a closeness check (KL, top token), not byte equality.

Host: test-mimo, test-cuda-host, test-kernel-contract, test-kernel-symbols-host, test-native-options-host, test-shared-lane-regressions-host, test-tp-cli and the CUDA compiles pass; zig fmt --check and git diff --check are clean; lean finds 428 problems, as main.

Prompt kernels: SASS, main against this branch
  • main: ashhart/TensorFold up/main b30d5b9. Its exl3 sources match the PR base 3f144cc.
  • branch: mimo-v2-p3. Its exl3_experts.cu, exl3_prompt.cuh and exl3_decode.cuh are the same as c0796d262 / 090eb82, "the common prompt path takes 3-bit mul1 experts over fp16 rows".
  • Build: CPU only, in muse-zig-build:0.17.0.
    • nvcc 13.3.73 is the image's own /usr/local/cuda.
    • nvcc 13.0 is the host's /usr/local/cuda-13.0, mounted read-only.
    • Command: nvcc -fatbin with torch_flags plus -O3 --expt-relaxed-constexpr and -I set to the source dir.
    • One -gencode=arch=compute_N,code=sm_N for each of N = 80, 86, 89, 120, 121. That covers the release fatbin's 80,86,120,121 (zig/build/cuda.zig release_sms) and the contract's sm_89.
  • Compare: cuobjdump -arch sm_N -sass -fun <kernel>.
    • Normalised by removing the /*0123*/ offsets and the /* 0x... */ encodings.
    • Branch targets were left as they are. Every diff below has equal lengths, so the addresses are unchanged.
    • -res-usage comes from cuobjdump -res-usage. Its SHARED count includes the 1 KiB system reservation on sm_120 and sm_121; the contract records 43264.

Results

SM nvcc kernel SASS lines main vs branch -res-usage
sm_80 13.3 tf_exl3_prompt_gu 9629 identical 128 regs, smem 43264
sm_80 13.3 tf_exl3_prompt_down 3157 identical 126 regs, smem 43264
sm_86 13.3 tf_exl3_prompt_gu 9133 identical 126 regs, smem 43264
sm_86 13.3 tf_exl3_prompt_down 3005 identical 126 regs, smem 43264
sm_89 13.3 tf_exl3_prompt_gu 9133 identical 126 regs, smem 43264
sm_89 13.3 tf_exl3_prompt_down 3005 identical 126 regs, smem 43264
sm_120 13.3 tf_exl3_prompt_gu 5325 identical 128 regs, smem 44288
sm_120 13.3 tf_exl3_prompt_down 1717 1 lines differ 127 regs, smem 44288
sm_121 13.3 tf_exl3_prompt_gu 5325 identical 128 regs, smem 44288
sm_121 13.3 tf_exl3_prompt_down 1717 1 lines differ 127 regs, smem 44288
sm_80 13.0 tf_exl3_prompt_gu 9645 identical 128 regs, smem 43264
sm_80 13.0 tf_exl3_prompt_down 3165 identical 126 regs, smem 43264
sm_86 13.0 tf_exl3_prompt_gu 9157 identical 126 regs, smem 43264
sm_86 13.0 tf_exl3_prompt_down 3013 identical 126 regs, smem 43264
sm_89 13.0 tf_exl3_prompt_gu 9157 identical 126 regs, smem 43264
sm_89 13.0 tf_exl3_prompt_down 3013 identical 126 regs, smem 43264
sm_120 13.0 tf_exl3_prompt_gu 5357 identical 124 regs, smem 44288
sm_120 13.0 tf_exl3_prompt_down 1749 1 lines differ 126 regs, smem 44288
sm_121 13.0 tf_exl3_prompt_gu 5357 identical 124 regs, smem 44288
sm_121 13.0 tf_exl3_prompt_down 1749 1 lines differ 126 regs, smem 44288

-res-usage is the same on main and the branch in every cell: registers, shared memory, stack 0 and local 0.

The only difference: tf_exl3_prompt_down on sm_120 and sm_121, with both compilers

sm_120, nvcc 13.3 (muse-zig-build:0.17.0), tf_exl3_prompt_down

-LOP3.LUT R5, R0, 0xe0, R3, 0xf8, !PT
+LOP3.LUT R5, R0, 0x3e0, R3, 0xf8, !PT

sm_121, nvcc 13.3 (muse-zig-build:0.17.0), tf_exl3_prompt_down

-LOP3.LUT R5, R0, 0xe0, R3, 0xf8, !PT
+LOP3.LUT R5, R0, 0x3e0, R3, 0xf8, !PT

sm_120, nvcc 13.0 (host /usr/local/cuda-13.0), tf_exl3_prompt_down

-LOP3.LUT R113, R113, 0xe0, R4, 0xf8, !PT
+LOP3.LUT R113, R113, 0x3e0, R4, 0xf8, !PT

sm_121, nvcc 13.0 (host /usr/local/cuda-13.0), tf_exl3_prompt_down

-LOP3.LUT R113, R113, 0xe0, R4, 0xf8, !PT
+LOP3.LUT R113, R113, 0x3e0, R4, 0xf8, !PT

The flags are LOP3.LUT with LUT 0xf8, which computes a | (b & c). Here c is R3, and R3 = S2R SR_TID.X (threadIdx.x).

The kernel launches with __launch_bounds__(256, 2) and 256 threads, so bits 8–9 of threadIdx.x are always 0. That makes tid & 0x3e0 equal to tid & 0xe0, so the result is the same. The wider constant probably comes from the templated index arithmetic, for example TW / 4 pieces in place of the literal 8. That source is not pinned down.

Every other kernel, SM and compiler combination is identical instruction for instruction: gate/up on all five SMs, and down on sm_80, sm_86 and sm_89.

Every top-token miss against exllamav3, with both sides' top 5 and margins

REF /fid/out/exl3-3.0/ref.bin (m1 dump), OURS /fid/data/dec4.bin (m1 dump); 1905 rows with logits, argmax equal 1878/1905. KL(exllamav3 || dec4): mean 4.256e-03, p50 3.056e-04, p99 5.288e-02, max 3.386e-01.

Margin = exllamav3's top-1 minus top-2 (logit, prob). dec4 rank = where dec4 puts exllamav3's top-1. Token = idtext. Next = the teacher token at that position.

prompt step KL next exllamav3 top-1 (logit, p) exllamav3 margin (logit, p) dec4 top-1 (p) dec4 p of exllamav3 top-1, rank exllamav3 p of dec4 top-1
en_story 18 (prompt) 0.000636 1943 message 25382 mysterious (16.8438, 0.1630) 0.0312, 0.0050 14888 strange (0.1622) 0.1615, 2 0.1580
en_story 37 0.00272 5001 coming 5001 coming (18.8438, 0.2429) 0.0781, 0.0183 13458 storm (0.2410) 0.2324, 2 0.2246
en_story 40 0.0059 576 The 1260 He (18.0312, 0.3508) 0.1875, 0.0600 576 The (0.3144) 0.3092, 2 0.2908
en_story 87 0.00398 45035 gratitude 5775 wonder (22.4688, 0.4402) 0.0625, 0.0267 45035 gratitude (0.4338) 0.4308, 2 0.4135
en_story 95 0.0598 264 a 264 a (17.0625, 0.1783) 0.2500, 0.0394 504 from (0.2353) 0.1526, 2 0.1389
en_story 161 0.0255 326 l 1943 message (17.3125, 0.3852) 0.1719, 0.0608 326 l (0.3987) 0.3000, 2 0.3244
en_story 177 0.00457 1035 would 1035 would (23.4062, 0.4174) 0.1562, 0.0604 44419 climbed (0.3757) 0.3751, 2 0.3570
en_story 216 0.0362 73632 lodged 73632 lodged (22.2188, 0.2552) 0.1094, 0.0264 4279 half (0.2813) 0.2623, 2 0.2287
en_story 263 0.00522 11, 304 in (25.7188, 0.4495) 0.0938, 0.0402 11, (0.4278) 0.4165, 2 0.4093
zh_explain 63 0.0199 101068无法 53153不能 (21.8906, 0.2774) 0.0312, 0.0085 101068无法 (0.3350) 0.2731, 2 0.2689
zh_explain 119 0.0325 117184粒子 100369两个 (22.5000, 0.4304) 0.3125, 0.1155 117184粒子 (0.4073) 0.3570, 2 0.3149
zh_explain 133 0.0398 1773。 106293的状态 (25.8438, 0.3952) 0.1250, 0.0464 1773。 (0.4682) 0.2807, 2 0.3488
code_py 145 0.0289 11, 13. (23.8594, 0.3009) 0.0625, 0.0182 11, (0.3421) 0.2576, 2 0.2783
math 60 0.0157 18611 \n\n 6771 Let (14.8438, 0.0822) 0.2578, 0.0187 18611 \n\n (0.0781) 0.0755, 2 0.0635
tool_json 22 (prompt) 0.00878 25: 510:\n (17.7812, 0.2198) 0.0781, 0.0165 25: (0.2269) 0.2251, 2 0.2033
edit 49 0.023 334** 88757Rename (16.1406, 0.1151) 0.2188, 0.0226 334** (0.0988) 0.0956, 2 0.0925
edit 50 0.035 16141Answer 18896Expected (16.0625, 0.0680) 0.1875, 0.0116 16141Answer (0.0761) 0.0622, 3 0.0563
edit 72 0.0122 1760 count 1760 count (24.5781, 0.5373) 0.1562, 0.0777 13195 cnt (0.5366) 0.4598, 2 0.4596
edit 279 0.0293 25: 982):\n (30.0312, 0.5592) 0.2656, 0.1305 25: (0.5490) 0.4393, 2 0.4288
list 106 0.00615 576 The 27360 Rot (22.9062, 0.4174) 0.1406, 0.0548 576 The (0.4054) 0.3734, 2 0.3626
short 0 (prompt) 0.029 0! 358 I (8.7500, 0.0129) 0.1953, 0.0023 11, (0.0173) 0.0142, 2 0.0075
short 10 0.00149 2371-g 2371-g (17.5156, 0.0956) 0.0156, 0.0015 1448-m (0.0996) 0.0948, 2 0.0941
short 16 0.00917 12- 5338First (14.1562, 0.0263) 0.0938, 0.0024 12- (0.0298) 0.0213, 2 0.0240
short 48 0.00195 21556 Sprite 51548 Collision (16.4219, 0.1594) 0.0000, 0.0000 21556 Sprite (0.1737) 0.1609, 2 0.1594
short 54 0.00165 358 I 358 I (20.8594, 0.3520) 0.0625, 0.0213 576 The (0.3544) 0.3427, 2 0.3307
short 59 0.00176 17051 sprite 2781 player (18.4219, 0.2963) 0.0000, 0.0000 17051 sprite (0.3132) 0.3028, 2 0.2963
short 208 0.00613 12440 correctly 12440 correctly (24.6719, 0.4150) 0.0625, 0.0251 6915 fine (0.4299) 0.4114, 2 0.3898
short 237 0.00629 4460 trying 4460 trying (17.5000, 0.1722) 0.1562, 0.0249 537 not (0.1535) 0.1494, 2 0.1473

Top-5 contenders per miss

en_story step 18 (KL 0.000636, next 1943 message)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 25382 mysterious 16.8438 0.1630 14888 strange 16.8495 0.1622
2 14888 strange 16.8125 0.1580 25382 mysterious 16.8452 0.1615
3 1943 message 16.5781 0.1250 1943 message 16.5374 0.1187
4 4704 mer 15.9844 0.0690 4704 mer 16.0381 0.0720
5 16486 bottle 15.9531 0.0669 16486 bottle 15.9756 0.0677

en_story step 37 (KL 0.00272, next 5001 coming)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 5001 coming 18.8438 0.2429 13458 storm 18.8042 0.2410
2 13458 storm 18.7656 0.2246 5001 coming 18.7677 0.2324
3 17478 terrible 17.9688 0.1012 17478 terrible 17.9973 0.1075
4 11406 dangerous 17.7969 0.0852 53577 catastrophic 17.6486 0.0759
5 53577 catastrophic 17.7188 0.0788 11406 dangerous 17.6442 0.0756

en_story step 40 (KL 0.0059, next 576 The)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 1260 He 18.0312 0.3508 576 The 17.8933 0.3144
2 576 The 17.8438 0.2908 1260 He 17.8766 0.3092
3 3555 What 16.1562 0.0538 3555 What 16.1726 0.0563
4 5443 Use 15.5781 0.0302 2411 At 15.5605 0.0305
5 2411 At 15.4922 0.0277 5443 Use 15.5011 0.0287

en_story step 87 (KL 0.00398, next 45035 gratitude)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 5775 wonder 22.4688 0.4402 45035 gratitude 22.5587 0.4338
2 45035 gratitude 22.4062 0.4135 5775 wonder 22.5517 0.4308
3 3633 connection 20.6094 0.0686 3633 connection 20.5354 0.0574
4 11752 responsibility 19.9531 0.0356 11752 responsibility 20.0431 0.0351
5 22531 mystery 19.1250 0.0155 22531 mystery 19.5688 0.0218

en_story step 95 (KL 0.0598, next 264 a)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 264 a 17.0625 0.1783 504 from 17.6081 0.2353
2 504 from 16.8125 0.1389 264 a 17.1750 0.1526
3 5326 written 16.2344 0.0779 5326 written 16.6422 0.0896
4 2310 old 16.2031 0.0755 2310 old 16.2594 0.0611
5 20068 addressed 15.6094 0.0417 4285 simple 15.7487 0.0367

en_story step 161 (KL 0.0255, next 326 l)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 1943 message 17.3125 0.3852 326 l 17.3710 0.3987
2 326 l 17.1406 0.3244 1943 message 17.0866 0.3000
3 2310 old 14.7969 0.0311 2310 old 15.1239 0.0422
4 16486 bottle 14.6875 0.0279 13458 storm 14.5754 0.0244
5 13458 storm 14.5625 0.0246 16486 bottle 14.5256 0.0232

en_story step 177 (KL 0.00457, next 1035 would)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 1035 would 23.4062 0.4174 44419 climbed 23.1726 0.3757
2 44419 climbed 23.2500 0.3570 1035 would 23.1708 0.3751
3 6925’d 21.5000 0.0620 6925’d 21.5276 0.0725
4 43361 polished 21.5000 0.0620 43361 polished 21.4740 0.0687
5 4172'd 21.3125 0.0514 4172'd 21.2201 0.0533

en_story step 216 (KL 0.0362, next 73632 lodged)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 73632 lodged 22.2188 0.2552 4279 half 22.1941 0.2813
2 4279 half 22.1094 0.2287 73632 lodged 22.1245 0.2623
3 37493 washed 22.0938 0.2252 448 with 21.6570 0.1644
4 448 with 21.5469 0.1303 37493 washed 21.4912 0.1393
5 10840 wed 20.9688 0.0731 10840 wed 20.5628 0.0550

en_story step 263 (KL 0.00522, next 11,)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 304 in 25.7188 0.4495 11, 25.4655 0.4278
2 11, 25.6250 0.4093 304 in 25.4388 0.4165
3 1573 before 23.9688 0.0781 1573 before 23.7683 0.0784
4 323 and 22.8750 0.0262 916 over 22.7325 0.0278
5 916 over 22.5781 0.0194 323 and 22.6215 0.0249

zh_explain step 63 (KL 0.0199, next 101068无法)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 53153不能 21.8906 0.2774 101068无法 22.1842 0.3350
2 101068无法 21.8594 0.2689 53153不能 21.9798 0.2731
3 111010并不能 21.7656 0.2448 111010并不能 21.7986 0.2278
4 100684并不 20.5000 0.0691 100684并不 20.1745 0.0449
5 100186不可 19.5312 0.0262 100186不可 19.6271 0.0260

zh_explain step 119 (KL 0.0325, next 117184粒子)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 100369两个 22.5000 0.4304 117184粒子 22.2534 0.4073
2 117184粒子 22.1875 0.3149 100369两个 22.1215 0.3570
3 105536测量 21.1094 0.1071 105536测量 20.5740 0.0760
4 100783无论 20.0000 0.0353 100783无论 19.6971 0.0316
5 56568你 19.2344 0.0164 113739手套 19.3134 0.0215

zh_explain step 133 (KL 0.0398, next 1773。)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 106293的状态 25.8438 0.3952 1773。 26.1253 0.4682
2 1773。 25.7188 0.3488 106293的状态 25.6139 0.2807
3 3837, 25.3125 0.2323 3837, 25.3670 0.2193
4 105369的结果 22.5938 0.0153 105369的结果 22.9586 0.0197
5 9370的 21.2656 0.0041 3407。\n\n 21.4087 0.0042

code_py step 145 (KL 0.0289, next 11,)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 13. 23.8594 0.3009 11, 23.7828 0.3421
2 624.\n 23.7969 0.2827 13. 23.4991 0.2576
3 11, 23.7812 0.2783 624.\n 23.3363 0.2189
4 382.\n\n 22.1719 0.0557 382.\n\n 22.4913 0.0940
5 369 for 21.5156 0.0289 369 for 21.4541 0.0333

math step 60 (KL 0.0157, next 18611 \n\n)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 6771 Let 14.8438 0.0822 18611 \n\n 14.6526 0.0781
2 18611 \n\n 14.5859 0.0635 6771 Let 14.6196 0.0755
3 5443 Use 13.9922 0.0351 5443 Use 13.9155 0.0374
4 320 ( 13.8594 0.0307 320 ( 13.7620 0.0320
5 576 The 13.8203 0.0296 576 The 13.6773 0.0294

tool_json step 22 (KL 0.00878, next 25:)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 510:\n 17.7812 0.2198 25: 17.6896 0.2269
2 25: 17.7031 0.2033 510:\n 17.6818 0.2251
3 1447:\n\n 17.5469 0.1739 1447:\n\n 17.3415 0.1602
4 2500 available 17.5312 0.1712 2500 available 17.2902 0.1522
5 525 are 16.4219 0.0565 525 are 16.3539 0.0597

edit step 49 (KL 0.023, next 334**)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 88757Rename 16.1406 0.1151 334** 16.2109 0.0988
2 334** 15.9219 0.0925 88757Rename 16.1779 0.0956
3 73594''' 15.6094 0.0676 73594''' 16.0752 0.0863
4 785The 15.4297 0.0565 785The 15.5451 0.0508
5 8420Here 15.0547 0.0388 8420Here 15.3682 0.0426

edit step 50 (KL 0.035, next 16141Answer)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 18896Expected 16.0625 0.0680 16141Answer 16.3351 0.0761
2 16141Answer 15.8750 0.0563 36842Solution 16.1681 0.0644
3 5097Output 15.7188 0.0482 18896Expected 16.1337 0.0622
4 36842Solution 15.6094 0.0432 25139Reason 15.7887 0.0440
5 25139Reason 15.5312 0.0400 5097Output 15.7098 0.0407

edit step 72 (KL 0.0122, next 1760 count)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 1760 count 24.5781 0.5373 13195 cnt 24.6058 0.5366
2 13195 cnt 24.4219 0.4596 1760 count 24.4514 0.4598
3 671 # 18.0000 0.0007 671 # 18.5080 0.0012
4 4210 """ 17.8281 0.0006 4210 """ 17.7880 0.0006
5 3190 """\n 17.3281 0.0004 3190 """\n 17.3879 0.0004

edit step 279 (KL 0.0293, next 25:)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 982):\n 30.0312 0.5592 25: 29.9161 0.5490
2 25: 29.7656 0.4288 982):\n 29.6931 0.4393
3 8) 26.1406 0.0114 8) 26.0225 0.0112
4 1648): 22.8750 0.0004 1648): 22.4911 0.0003
5 7731):\n\n 20.2188 0.0000 78079:list 20.1844 0.0000

list step 106 (KL 0.00615, next 576 The)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 27360 Rot 22.9062 0.4174 576 The 22.9715 0.4054
2 576 The 22.7656 0.3626 27360 Rot 22.8894 0.3734
3 1527 An 21.6094 0.1141 1527 An 21.7553 0.1201
4 350 T 20.4219 0.0348 350 T 20.3567 0.0297
5 11443 Has 19.9062 0.0208 28650 Unique 20.0326 0.0215

short step 0 (KL 0.029, next 0!)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 358 I 8.7500 0.0129 11, 9.1645 0.0173
2 2822()\n\n 8.5547 0.0106 358 I 8.9642 0.0142
3 220 8.3828 0.0089 99466大家 8.9540 0.0140
4 99466大家 8.3672 0.0088 220 8.7522 0.0115
5 271\n\n 8.3594 0.0087 2822()\n\n 8.6865 0.0107

short step 10 (KL 0.00149, next 2371-g)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 2371-g 17.5156 0.0956 1448-m 17.4200 0.0996
2 1448-m 17.5000 0.0941 2371-g 17.3714 0.0948
3 12646-k 17.3281 0.0792 12646-k 17.1801 0.0783
4 12- 17.1094 0.0637 12- 16.9382 0.0615
5 74k 17.0000 0.0571 74k 16.8751 0.0577

short step 16 (KL 0.00917, next 12-)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 5338First 14.1562 0.0263 12- 14.2467 0.0298
2 12- 14.0625 0.0240 5338First 13.9102 0.0213
3 4868Game 13.7109 0.0169 4868Game 13.6994 0.0172
4 7849Project 13.6172 0.0154 7849Project 13.5695 0.0151
5 16915Library 13.5078 0.0138 2271Test 13.3689 0.0124

short step 48 (KL 0.00195, next 21556 Sprite)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 51548 Collision 16.4219 0.1594 21556 Sprite 16.5302 0.1737
2 21556 Sprite 16.4219 0.1594 51548 Collision 16.4536 0.1609
3 17051 sprite 16.1406 0.1203 17051 sprite 16.1498 0.1187
4 18752 collision 15.9844 0.1029 18752 collision 15.9831 0.1005
5 5316 script 15.9219 0.0967 5316 script 15.8944 0.0920

short step 54 (KL 0.00165, next 358 I)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 358 I 20.8594 0.3520 576 The 20.9574 0.3544
2 576 The 20.7969 0.3307 358 I 20.9239 0.3427
3 3197 When 20.0625 0.1587 3197 When 20.1066 0.1513
4 758 In 17.7969 0.0165 758 In 17.8073 0.0152
5 1096 This 17.4844 0.0120 1096 This 17.5074 0.0113

short step 59 (KL 0.00176, next 17051 sprite)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 2781 player 18.4219 0.2963 17051 sprite 18.5527 0.3132
2 17051 sprite 18.4219 0.2963 2781 player 18.5186 0.3028
3 21556 Sprite 17.2500 0.0918 21556 Sprite 17.3347 0.0927
4 10434 texture 17.1094 0.0797 10434 texture 17.1163 0.0745
5 3668 character 16.2188 0.0327 3668 character 16.2724 0.0320

short step 208 (KL 0.00613, next 12440 correctly)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 12440 correctly 24.6719 0.4150 6915 fine 24.8923 0.4299
2 6915 fine 24.6094 0.3898 12440 correctly 24.8483 0.4114
3 11, 23.8125 0.1757 11, 23.7965 0.1437
4 14279 normally 20.0469 0.0041 14279 normally 20.0828 0.0035
5 304 in 19.8125 0.0032 438 as 19.5855 0.0021

short step 237 (KL 0.00629, next 4460 trying)

# exllamav3 token exllamav3 logit exllamav3 p dec4 token dec4 logit dec4 p
1 4460 trying 17.5000 0.1722 537 not 17.2421 0.1535
2 537 not 17.3438 0.1473 4460 trying 17.2153 0.1494
3 1667 using 17.0938 0.1147 1667 using 16.9385 0.1133
4 14589 sorry 16.7656 0.0826 14589 sorry 16.7898 0.0976
5 501 new 16.5625 0.0674 501 new 16.4493 0.0695

@ashhart

ashhart commented Oct 27, 2026

Copy link
Copy Markdown
Owner

Thanks @ajensenwaud, all four points are done. The load-ahead is its own PR, the prompt-kernel SASS table settles the shared-kernel question, every miss against exllamav3 is listed with both sides' top 5, refused loads and bad TF_TP_* settings now fail cleanly with host tests, and the recipe names the licences. Your TP=1 answer matches docs/multi-rank-receipts.md: each world size has its own equalities. We merged d661af0 onto main a1e6d34 (clean). fmt and lean are clean, and test-mimo and the CUDA host suites pass.

For a new family, admission should match what was validated:

  1. Hardware. The records start at sm_121 (mimo_cuda_records.zig:8), but the kernels build for every release SM and open (native.zig:446) never checks the GPU. Please refuse anything but sm_121 by name, as deepseek_v41/kernels.zig:79 does, with a test.
  2. One rank. config.zig:127 admits TP=1, and the recipe says the 2.27 bpw pack serves too, but every receipt is TP=2 on the 3.0 bpw pack. Either post TP=1's own equalities and an exllamav3 run on the 2.27 bpw pack, or refuse one rank by name for now.
  3. Receipts. Please move the exllamav3 numbers into the description in place of the captured-set comparison. 1,878 of 1,905 is 27 misses, but the table lists 28 (two are exact ties), so please reconcile the count. Please also pin full revisions in the recipe for both packs and the model.

Then on our GB10 pair we'll run the contract on this head and a GLM-5.3-Flash control. We'll run MiMo at TP=2 on the 3.0 bpw pack (drafted and plain, concurrent and solo, resumed across the chunk marks), and TP=1 on one GB10 with the 2.27 bpw pack.

A family whose head drafts a whole block in one pass (DFlash) pays for a draft pass once, whatever its depth, so the
per-depth rule prices deep drafts too high. With `geometric_depth` set, a lone stream's depth is the one with the most
expected tokens a ms under a geometric chain of its own acceptance rate (a moving average of kept over tried, from
0.6), the forward of its rows plus one draft pass; a stream that stops drafting tries one draft every 8 rounds.

A round's cost can grow with the stream's context: `geometric_ctx` adds, per 1,000 positions past
`geometric_ctx_base`, a step and a per-draft term, so long streams draft shallower as full-attention reads grow.
`geometric_costs` and `geometric_draft_cost` give the per-depth rule of shared rounds the family's own cost table.

The rule follows the Python engine's MiMo-V2 decoder (joint_widths). Everything defaults off: families that do not
set it keep the per-depth rule unchanged.
…rows

The prompt kernels' tile decode, weight-stage sizes and loads take the half-bit width and codebook as template
arguments; tf_exl3_prompt_gu / _down stay the 4-bit mcg instance and tf_exl3_prompt_gu_k6_mul1 / _down_k6_mul1 are
the 3-bit mul1 one (MiMo-V2.6-Flash's routed experts). eligible() takes either verified uniform width, Kernels picks
the instance by layer, and run() takes the rows' dtype for the per-pick rotation (bf16 or fp16; the shared-rotation
kernel stays bf16). GLM passes .bf16. On sm_121 the 4-bit gate/up kernel's SASS is unchanged and the down kernel's
differs in one mask constant (0xe0 -> 0x3e0, the same value for the 256-thread launch).
… DFlash, two GPUs, prompt cache

The mimo_v2 family serves an EXL3 pack of MiMo-V2.6-Flash on the native CUDA server, on one GPU or two DGX Sparks
as tensor-parallel ranks, with the pack's DFlash drafter. Every kernel is built in: no captured set, no copies of
the Python engine's kernels.

- EXL3: projections and decode experts on cuda.exl3.Linear and cuda.exl3_experts; prompt experts on the common prompt
  kernels (their 3-bit mul1 instance, the commit before), then the family's slot sum. Every prompt token runs on the
  prompt path, so the decode/prompt cut is gone.
- kernels/cuda/mimo_v2: ops.cu (RMS norms, SwiGLU, the fp32 router and its top 8, the experts' slot sum),
  attention.cu (queries and caches, chunked decode attention with sinks and its merge, a tensor-core prompt kernel
  for the full layers whose 64-key tiles start at 0, so a row's bits never depend on its chunk), glue.cu (embed,
  residual adds, RoPE tables, argmax), dflash.cu (the drafter's kernels).
- families/mimo_v2: config and rank cuts, loading, the decode forward (eager and graphed by rows and key bucket),
  the prompt path, the DFlash drafter, the engine both ranks run, and kept prompt states (kept.zig: the slot's
  full-layer rows pinned, the sliding and drafter windows copied; a resume binds to the state's slot or copies the
  rows into a free one).
- native.zig: the lane backend. Rank 0 serves and sends every GPU step (and cache step) over the plan link before
  launching it; rank 1 runs `tf-mimo follow`. --parallel sets the streams' cache slots; the prompt cache's marks
  are 2,048 rows apart; a lone stream's first copy window is 6 rows; draft depth uses the lanes' geometric rule with
  a context term.
- native/cuda.zig: MiMo registered. Admission takes main's weightShare hook: the split matrices over the ranks plus
  the embeddings and the drafter whole on each; a bad TF_TP_* setting refuses instead of falling back to one GPU.
- Checks: tf-mimo m1 (teacher-forced KL and argmax against an answer key and an earlier build's dump, one prompt
  chunk against two) and host tests for config, cuts, plans, kept-state runs and slot choice.

The family's build is zig/build/mimo.zig (as kolibri.zig), its kernels are listed in cuda_kernels.zig, and the TP
mailbox fatbin is built once for tf-tp-test and tf-mimo. The ranks run over Jay Leaton's tensor-parallel session and
RoCE mailbox from his two-Spark DeepSeek-V4.1 work.
… for gate 5

contract_mimo*.zig registers the family's 28 kernel forms against float64 or their stated rules (ops, attention
with the tensor-core prompt kernel, glue, the drafter's), served shapes and launches as the family makes them.
mimo_cuda_records.zig carries their sm_121 usage from nvcc 13.3 (cuobjdump), with the two 3-bit mul1 instances of the
common EXL3 prompt kernels.
… shallower top-8 tree

mimo_attn_rows runs a chunk's G <= 16 query heads as one m16 tile on the prompt kernel's loads and mma, keys in
64-key tiles at multiples of 64 split over four warps whose partials fold in warp order; p in fp16 for o (the
contract's allowance as the prompt kernel's). A row's bits still depend only on its own position and window.
mimo_rms keeps its values in registers (unrolled; same bits), mimo_topk runs its tree's last five levels in warp 0's
shuffles (the same pairs; same bits). Sliding rows go from 72 to 8.5 us, norms from 11 to 5.7 us.
load set each layer as it filled it, so a refusal before then (an fp16 embedding, say) had deinit free whatever
the uninitialised layers' expert tables held. Every layer now starts blank, without experts. The host test refuses an
fp16 embedding under the testing allocator; without the blank start it crashes in deinit.
… GPU

weightShare read TF_TP_* with catch null, so a malformed world or rank was admitted as one GPU's share. It now
returns BadSetting as open does, and the admission reports it through the family's explain. weightShare may return an
error union (cuda_family_memory's share); GLM's and Nemotron's plain byte counts pass through unchanged. TF_TP_WORLD
unset is still one GPU. Host test: unset, six malformed settings, and a two-rank share.
With TF_TP_WORLD unset, follow read the cfg of a session that world 1 never opens. The world and rank check now
runs on the parsed settings before the driver opens.
@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Thanks @ashhart. All three points are addressed. The branch is rebased on main 15a0b2a9, and the description is updated.

  1. Hardware. open now refuses anything but sm_121 by name, MimoV2RequiresSm121, in the same style as deepseek_v41/kernels.zig:79. The check runs first, before any setting or file is read, and tf-mimo checks it too.
    • Host test: open on a fake driver refuses sm_80, 86, 89, 90, 100, 103, 110 and 120, and at sm_121 it gets past the check.
    • The error's explanation says why: the launch records and every receipt are GB10's.
    • Commit 92569309.
  2. One rank. I've refused it by name for now: MimoV2RequiresTwoRanks.
    • rankView, the admission weight share and open all refuse it. That covers TF_TP_WORLD unset, 1, or anything but 2. The share is included so that a one-GPU run gets this error, not a "does not fit" message.
    • Host tests: the share refuses unset, 1, 3 and 4 and still reports a bad setting as before, and rankView refuses 1, 3 and 4.
    • The recipe no longer offers one GPU or the 2.27 bpw pack. It lists that pack as not yet validated.
    • Admitting one rank later would come with its own equalities and an exllamav3 run on the 2.27 bpw pack.
    • Commit e082ebcf.
  3. Receipts. The exllamav3 comparison now replaces the captured-set comparison in the description.
    • The count. The 27 against 28 came from two exact ties: positions where exllamav3's top two fp16 logits are equal and our top token is one of them. Both are in the short text, at steps 48 and 59.
      • The count broke a tie by the lower token id. That credited step 48 and not step 59, so 27 misses (1,878).
      • The table's sort broke both ties against us, so 28 rows.
      • Precisely: 26 positions where the top tokens differ, plus those 2 ties. That is 1,879 of 1,905 with a shared top token.
      • Against exllamav3's fp32 head output there are no ties, and 29 positions differ (1,876 of 1,905).
      • The description now shows the 26 and the 2 ties in separate tables.
      • A correction to my last reply: the exllamav3 reference is its fp16 head output widened to fp32, not fp32 logits; that is why exact ties occur.
    • Revisions. The recipe pins full revisions (commit 7a8efd85):
      • model XiaomiMiMo/MiMo-V2.6-Flash-RL at 5711b268169967567844e1e560e8a3966da959b1. This is the revision the packs were converted from and the one validated. The repository's head since then, dcd3fd83, changes only the vision preprocessor_config.json.
      • ajensenwaud/MiMo-V2.6-Flash-RL-exl3-3.0bpw at 139d9a4e1eb32193890036fdddd521fe37033a30.
      • benthecarman/MiMo-V2.6-Flash-RL-exl3 at 5ae87b830eb950c705967aea6130601a8e1f74a5.
      • The recipe's download line uses the 3.0 bpw pack's revision.

Host checks. test-mimo, test-kernel-contract, test-kernel-symbols-host, test-native-options-host, test-shared-lane-regressions-host and test-tp-cli pass, and the CUDA compiles build. test-cuda-host passes except two core prompt-cache spill tests (rows_async_test, rows_close_test), which fail the same way on main 15a0b2a9 in the same container. zig fmt --check and git diff --check are clean, and lean finds 428 problems, as on main.

ajensenwaud added a commit to ajensenwaud/TensorFold that referenced this pull request Oct 28, 2026
…one GPU

weightShare read TF_TP_* with catch null, so a malformed world or rank was admitted as one GPU's share. It now
returns BadSetting as open does, and the admission reports it through the family's explain. weightShare may return an
error union (cuda_family_memory's share, the same change as MiMo-V2's in ashhart#728). Host test: unset, six malformed
settings, and a two-rank share.
@ajensenwaud

Copy link
Copy Markdown
Contributor Author

Pushed 021eb4c3: the Windows host build failed because the new admission test calls open, which reaches the TP sockets. It is now left out on Windows the way the server's speech tests are (if (comptime @import("builtin").os.tag != .windows)), folded into the sm_121 commit. The three new commits are now 48468450 (sm_121), 799c87f3 (two ranks) and 021eb4c3 (revisions); the description is updated.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants