Repository navigation
feat(ep): expert parallelism stage 0 + stage 1 (EP=2) for V4's MoE - #89
Conversation
Splits the 256 routed experts across ranks so each card reads only its slice, with attention left data-parallel. Design and arithmetic: memory/mission/wave44-BV-expert-parallelism.md; full log for this wave: memory/mission/wave60-CK-expert-parallel.md. NO GPU RAN. Nothing here is a multi-GPU measurement. Stage 0 - QTIP CUDA wrappers validated device KIND, not ORDINAL: activations on cuda:0 with expert weights on cuda:1 passed the check and launched on the wrong stream. Fixed in all 10 multi-operand wrappers via a new qtip/device_guard.rs whose core predicate is a pure function over DeviceLocation, so the ordinal comparison is testable without a GPU. - Neither candle nor cudarc calls cuCtxEnablePeerAccess, so every cross-GPU copy is staged through host RAM. Adds cuda_peer.rs (enable_peer_access -> PeerAccessReport) and, crucially, a CALL SITE in DeviceMapSetting::into_mapper that WARNs when any pair is not peered. cudarc symbols verified against the real 0.19.4 source, not memory. - Offline tid2eid expert->device distribution (arc-engine) plus the runtime placement it feeds, and a per-rank balancedness counter (tokens AND distinct experts), off unless ARC_EP_BALANCE=1. Stage 1 - ExpertPlacement (contiguous / explicit / LPT-balanced) and ExpertParallelPlan: global->local expert remap plus a mask that zeroes unowned slots, from device-resident tables built once at load. - ExpertSubset slices the expert axis at load, and REFUSES BY NAME on formats whose leading dim cannot be sliced rather than putting every expert on every rank. - QuantMethod::select_experts, implemented for QtipLayer and Qtip2bLayer (V4's serving rung), so a UQFF artifact can be narrowed after deserialize. Until that runs, MoEExperts::forward refuses. - ep_size now does something: read from the model config, overridable with ARC_EP_SIZE, and it errors when it disagrees with the communicator's world size instead of silently degrading. Tests (41 new, mutation-proved) - Collective correctness: ring_algebra asserts the all-reduce RESULT equals the true sum, and that the construction guard agrees with the algebra at every world size 2..16 -- the disagreement that let the original single-exchange ring bug ship. - EP=2/EP=4 combined output reproduces EP=1, including under a permuted placement. Proved discriminating by three applied-and-reverted mutations: dropping the remap, dropping the mask, dropping the combine's masked weights. Corrects wave44-BV 3.2: its per-N imbalance spread (+4.8/4.6/4.2% at B=128) does not follow from its own definition -- above E/N coverage the N cancels and the bound is a single +5.05%. Slightly worse than claimed; the conclusion (skip EPLB in stage 1) is unchanged. Upstream mistralrs-* files were not mass-reformatted (fork policy, fab114f): incidental rustfmt churn in deepseek4.rs and distributed/layers.rs was reverted, leaving only semantic hunks. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01RRYaA9dQfd7LECusk3dQrk
Code Metrics Report━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Language Files Lines Code Comments Blanks ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ C Header 5 305 210 52 43 CSS 2 1181 1036 34 111 CUDA 72 24645 17702 4215 2728 Dockerfile 1 39 22 8 9 JavaScript 16 3546 2676 482 388 Jinja2 7 694 656 5 33 JSON 74 4600 4597 0 3 Makefile 1 6 5 0 1 Metal Shading Lan| 33 12224 9431 1142 1651 PowerShell 1 300 227 30 43 Python 143 14830 12217 797 1816 Shell 24 5924 4071 1457 396 Plain Text 4 3801 0 2479 1322 TOML 33 1485 1292 43 150 YAML 3 25 23 2 0 ───────────────────────────────────────────────────────────────────────────────── HTML 4 2687 2604 43 40 |- CSS 2 543 479 37 27 |- JavaScript 1 1233 1215 12 6 (Total) 4463 4298 92 73 ───────────────────────────────────────────────────────────────────────────────── Jupyter Notebooks 4 122 83 23 16 |- Markdown 1 60 30 22 8 |- Python 1 122 113 1 8 (Total) 304 226 46 32 ───────────────────────────────────────────────────────────────────────────────── Markdown 194 38699 0 29715 8984 |- BASH 71 1632 1192 324 116 |- C 3 17 17 0 0 |- CUDA 2 84 56 16 12 |- JSON 18 708 708 0 0 |- PowerShell 1 1 1 0 0 |- Python 23 1008 787 113 108 |- Rust 65 2048 1713 77 258 |- TOML 6 207 164 0 43 |- YAML 4 38 33 5 0 (Total) 44442 4671 30250 9521 ───────────────────────────────────────────────────────────────────────────────── Rust 663 313742 271458 14550 27734 |- Markdown 478 24544 471 21171 2902 (Total) 338286 271929 35721 30636 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1284 461100 335289 76855 48956 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
|
BLOCKED — triage pass 2026-08-17. Disposition: BLOCKED, not closed. Stage 0 is real and independently valuable (the QTIP device-ordinal guard is a genuine latent-corruption fix; What blocks it: this PR claims EP=2 works. Arc has never executed a forward pass on more than one GPU — not once. 41 green CPU tests establish that the algebra is right (ring all-reduce equals the true sum at every world size 2..16; EP=2/EP=4 reproduce EP=1 under a permuted placement). They establish nothing about NCCL, peer access, or the actual expert-weight slicing on device. Merging this to master would put "expert parallelism works" in the tree on the strength of a simulation. Exact unblocking condition — all three, on a 2-GPU box:
This PR ships no script. That is part of the blocker: add Per D15 I do not call |
Jish set the architectural decision: "we're making our own systems here... name everything that deserves a name and no subsystem should be left without an absolute parent system name." Adds memory/mission/TAXONOMY.md — the full tree, enumerated from the code rather than drawn as an idealised diagram. Eight parent systems under Arc: ArcServe, ArcInfer (ArcSched/ArcKV/ArcAttention/ArcSpec/ArcMoE/ArcGraph/ ArcSample/ArcBoost), ArcModels, ArcQuant (QTIP/TurboQuant/ArcBake), ArcKernels (ArcTarget), ArcFormat (ArcOverlay), ArcLab, ArcGate. New models land in ArcModels; new GPUs land in ArcTarget, so future support always has a home. QTIP and TurboQuant keep their paper names per Jish; the attribution model is the existing kv_sharing/NOTICE breakdown (ADAPTED/VERBATIM/OURS). Documentation only — no code symbol, crate, or env var is renamed. All .rs changes are module-doc comments. Code renames are proposed as a sequenced follow-up in the taxonomy's last section. Correctness fixes across the sweep: - GSM8K 87.0% (n=100, provisional) superseded by 96.3% (1270/1319, full set, 0-shot, 0 degenerate / 0 truncated / 0 errors). Reference 90.8 is 8-shot. - b=1 decode 14.58 tok/s superseded by 18.27 aggregate / 17.99 per-user p50; end-to-end serving 111.69 tok/s at B=256, $12.06/Mtok now recorded. - TurboQuant's "4.27x KV compression, measured end-to-end" RETRACTED: it is format arithmetic and was never produced by a forward pass. TurboQuant is off by default, has a kernel at head_dim 128 only, none at 512, and has never been measured. - codebook.rs claimed Beta(d/2,d/2); the density reproducing the shipped tables is (1-x^2)^((d-3)/2), i.e. Beta((d-1)/2,(d-1)/2). generate.rs already had this right and pins it to 1e-9; the two modules now agree. - V4 cannot use FlashAttention at any generation: sinks on all 43 layers and the fused flash-sinks gate excludes head_dim 512. Stated wherever it was implied otherwise. - Marked what does not exist rather than letting it read as shipped: the fused 512 kernel (no PR at all), expert parallelism (#89), KV segment allocator (#90), LDLQ, EoRA, the xKV pool. - Historical numbers are kept and labelled superseded, not deleted; the cudnn -62% A/B ratio still stands and is preserved as such. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Unblocks the Typos CI gate; the only failing check on this PR.
…l-ep2 # Conflicts: # mistralrs-core/src/moe/experts.rs
|
Two updates from the PR sweep. The Typos gate is fixed ( The body says To be explicit about what that does and does not buy: it proves the cuda-gated arms compile. EP=2 remains entirely unmeasured — no forward pass has run on more than one GPU, and it cannot be tested on the current single-H200 box, so the honest framing in the header stays exactly as written. Merging this is merging inert, default-off code ( One thing I checked while resolving and want on the record rather than in a review queue: |
…-reduce `dummy_ops::SumAllReduce::sum_all_reduce` is `Ok(xs.clone())` — a silent no-op. Under expert parallelism each rank computes only the partial sum of the experts it owns, so a no-op combine leaves every rank serving a partial sum as if it were the whole: fluent, wrong, and caught by nothing downstream. The production build ships no `nccl` feature, so `Comm` is `Dummy`. The only thing between a mis-built binary launched with 2 ranks and that outcome is `build_expert_parallel_plan`'s `comm.world_size() != ep_size` bail — and that bail fires only because `DummyComm` DISCARDS the rank/world_size it was constructed with and hardcodes `0`/`1`. That discard reads like a bug. The "fix" — honour the constructor arguments — is precisely the change that re-arms the no-op, silently, because every other test still passes and the model still speaks. So it gets a tripwire. Mutation-checked: rewriting `DummyComm` to store and return its constructor arguments fails `dummy_comm_pins_world_size_to_one_...` with the reason, at world sizes 2, 8 and 64. The second test asserts the consequence as arithmetic — a no-op combine returns a finite, plausible, wrong number rather than an error — and deliberately survives that mutation, because it is about what the guard protects, not about the guard. Verified separately while pricing EP, and worth recording here: `MISTRALRS_NO_NCCL=1` is not an escape hatch — it drops `Comm` to `Dummy`, so EP bails rather than running. The `ring` backend is the only one giving a real all-reduce without tripping `use_nccl()`, but it is TCP-socket based and would understate an NVLink EP win. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…l-ep2 # Conflicts: # mistralrs-quant/src/lib.rs
…ns, 7 rules The headline, recorded so the next session inherits it instead of re-deriving it: **when instruction-bound, bits-per-weight CANCELS OUT of the throughput equation** (`achieved = 23,409/X` at B=256, bpw-invariant). The 2-bit format does not buy speed — it buys the model existing on one card: 74 GB vs MXFP4's 151 and FP8's 284, neither of which fits an H200. SGLang's V4 numbers run on two GH200s. The capacity wedge is already won by the format. It is a DIFFERENT claim from speed, where we are 15–25× behind. They must never share a table — conflating them produced the "16,600 → 1,477" framing Jish rejected. Recorded, with provenance: - The trellis ladder is COMPILED, not estimated (`nvcc -cubin`, CUDA 12.4.131): shipped 15.125 inst/wt → K8/V4/L12 + row-scale hoist 4.375 = 3.46× fewer instructions, still 3.1–3.9× short of the 1.13–1.41 budget. Both LUTs are exactly 32,768 B (no `cudaFuncSetAttribute`); occupancy 62.5%, register-limited. - V4 KV over-retention 30–57×; launches 9,131 → 8,650/token; frozen Gumbel 1 → 64 distinct tokens. - 13 retractions, four of them main's, including the ceiling claim, the MTP acceptance figure and its explanation, and an inverted re-bake cost. - 7 new kernel rules, each paid for — chiefly "a number you did not compile is an estimate" (hand-counting undercounts SASS by 2.05×) and "score inst/weight against its OWN bpw budget". - 7 sequencing gates, each a gate and not a preference. - TAXONOMY: expert parallelism IS on master (PR #89, `610c4506b`) — the gap is a missing CLI flag; `mark_unreachable` is itself dark; `CrossPrefixMeter` is not dead; head_dim 512 is not the blocker (448 is); Arc has never run a forward pass on more than one GPU. Also mirrors the four previously-unversioned mission docs into the repo. They lived only in the agent memory directory, which has no backup — and the repo's TAXONOMY copy had already drifted a revision ahead of the live one, which is exactly the failure the new sync note makes visible. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
Splits V4's 256 routed experts across ranks so each card reads only its slice, with attention left data-parallel. Design and arithmetic:
memory/mission/wave44-BV-expert-parallelism.md(PR #70). Full log:memory/mission/wave60-CK-expert-parallel.md.🔴 NO GPU RAN. $0.00 spent. Nothing here is a multi-GPU measurement — Arc has still never run a forward pass on more than one GPU. The per-user tok/s number is not in this PR; §6 of the wave log is the exact script and its cost.
Stage 0 — three ways stage 1 could have failed for reasons unrelated to EP
cuda:0with expert weights oncuda:1passedmatches!(.., Device::Cuda(_))and the kernel launched on the wrong device's stream — silent corruption, not an error. Unreachable whilemapper.mapco-locates everything; exactly the invariant EP breaks. Fixed in all 10 multi-operand wrappers viaqtip/device_guard.rs, whose core predicate is a pure function overDeviceLocationso the ordinal comparison is testable without a GPU.cuCtxEnablePeerAccess, so every cross-GPU copy is staged through host RAM — a cost the existing pipeline-parallel map already pays. Addscuda_peer.rsand, crucially, a call site inDeviceMapSetting::into_mapperthat logs at WARN when any pair is not peered. cudarc symbols verified against the real 0.19.4 registry source, not from memory.tid2eiddistribution, which wave44-BV flagged as "free to check and nobody has": V4's first 3 layers route by a fixed token-id table, so their expert load is a closed-form property of the tokenizer. Plus a per-rank balancedness counter recording both tokens (compute) and distinct experts (bytes), off unlessARC_EP_BALANCE=1— it needs a D2H read, which is incompatible with CUDA-graph capture.Stage 1 — EP=2
ExpertPlacement(contiguous / explicit / LPT-balanced) andExpertParallelPlanremap global expert ids to rank-local ones and zero the weights of unowned slots, from device-resident tables built once at load. Rankr's partial isΣ_{j owned by r} w·E_g(x); every slot is owned by exactly one rank, so the combine is a plain sum.ExpertSubsetslices the expert axis at load and refuses by name on formats whose leading dimension cannot be sliced, rather than quietly putting all 256 experts on every rank — "EP appeared to work and bought nothing" is a capacity bug that looks exactly like a healthy run. A UQFF artifact holds every expert, soQuantMethod::select_experts(implemented forQtipLayerandQtip2bLayer, V4's serving rung) narrows the stacks after deserialize, andforwardrefuses until it has.ep_sizenow does something. It was deserialized atarc-engine/src/deepseek_v4.rs:48and never read — the twelfth "wired but never invoked" case in this repo. It is now read from the model config, overridable withARC_EP_SIZE, and it errors when it disagrees with the communicator's world size instead of silently degrading.The tests are the teeth
41 new tests, all green, each proved discriminating.
Collective correctness —
ring_algebraasserts the all-reduce result equals the true sum, not merely that it ran, and asserts the construction guard agrees with the algebra at every world size 2..16. That disagreement is precisely how the original single-exchange ring bug shipped: it returnedself + left neighbourat N≥4 with no error. It lives outside theringfeature so CI actually runs it.EP=2 / EP=4 reproduce EP=1, including under a permuted placement where local 0 ≠ global 0. Fixture discrimination: every expert computes a different function, so routing to the wrong one cannot hide. Three mutations were applied, run, and reverted:
localize_masks_...,ep2_reproduces_ep1_output,ep2_..._permuted_placementmasks_partition_the_routing_slots_across_ranks,ep4_reproduces_ep1_outputep4survived the first mutation — with one expert per rank the local index is always 0, so the remap is right by coincidence. Recorded rather than hidden.A correction to wave44-BV §3.2
The note's per-N imbalance spread (+4.8/4.6/4.2% at B=128) does not follow from its own definition. Above
E/Ncoverage theNcancels:(E/N)/(distinct/N) = E/distinct. The real bound is a single +5.05% at B=128 (+28.1% at B=64, +0.23% at B=256) — slightly worse than claimed, with the conclusion unchanged: skip EPLB in stage 1, instrument for it. The test asserts the corrected values.What is not done
--features 'cuda flash-attn',CommisDummyandsum_all_reduceisOk(xs.clone()). Theep_size != world_sizecheck is the only thing standing between a mis-built binary and a fluent, wrong model.normal.rs:374-380skips auto device mapping under NCCL, andattention/mod.rs:232silently forcesnaive_sdpa— turning NCCL on turns flash-attention off.cargo check --features cudahas never run on this branch (cudarc's build.rs needs nvcc).Verification
cargo check --workspace;cargo test -p mistralrs-core -p mistralrs-quant -p arc-engine(383 / 262 / 149, 0 failures); the scoped clippy lane clean. Upstreammistralrs-*files were not mass-reformatted (fork policy,fab114fe3) — incidental rustfmt churn indeepseek4.rs(341 lines) anddistributed/layers.rs(854 lines) was reverted, leaving only semantic hunks.🤖 Generated with Claude Code