feat(escha): serve Qwen3.6-35B-A3B-Escha-W2 natively (2-bit eschamoe + int8), lossless AR ~58 tok/s - #242
feat(escha): serve Qwen3.6-35B-A3B-Escha-W2 natively (2-bit eschamoe + int8), lossless AR ~58 tok/s#242davidtai wants to merge 7 commits into
Conversation
…TP spec-decode Escha-W2 is a 2-bit Qwen3.6-A3B (qwen3_5_moe hybrid: 30 GDN + 10 full-attn, 256 experts / top-8) whose MoE experts are the eschamoe trellis format and whose non-expert weights are per-out-channel symmetric int8. It now loads and serves through the ordinary mtplx runtime — mtplx.runtime.load(path, mtp=True) -> MTPLXRuntime.forward_ar -> batched_decode speculative decode — exactly like any other MTP model. What's here: - mtplx/eschamoe.py: on-the-fly 2-bit expert decode (never caches dense weights); escha_qmv Metal matvec ported from dusterbloom/higgs (cited in NOTICE). - mtplx/int8_linear.py: fused int8 non-expert matvec (no int8->bf16 dequant at load). This is the big serving win: decode 25.1 -> 59.6 tok/s, quality-preserved. - mtplx/escha_load.py: builds the standard qwen3_5_moe trunk and swaps in the eschamoe experts + int8 Linears, returning an ordinary qwen3_5_moe model. - mtplx/qwen3_5_mtp_patch.py: Escha's MTP head ships no routed experts, so it borrows the trunk's eschamoe experts and shifts its (w-1)-convention norms +1. Also fixes the object level the qwen3_5 MTP injector attaches to. mlx-lm's outer qwen3_5_moe.Model exposes .language_model (no .model/.lm_head); its .language_model (the TextModel) exposes .model + .lm_head (no .language_model). validate_mtp_support inspects _text_model(model).mtp, so the MTP surface must live on the TextModel and be re-exposed on the outer wrapper via delegation — the same dual-level pattern as the generic inject_mtp_support. The previous code set .mtp on the outer model and mixed self.model (TextModel-level) with self.language_model.lm_head (outer-level) on one object, so validate failed and forward hit AttributeError: 'no attribute model'. This repairs the base qwen3_5_mtp path too (its object level was never exercised by CI). - mtplx/runtime.py: dispatch Escha to the qwen3_5 MTP injector; load via escha_load; skip the standard-A3B dense-affine MoE/MLP serving opts (they don't apply to eschamoe/int8 and batched_decode doesn't depend on them). Tests: - tests/test_qwen3_5_mtp_object_level.py: hermetic CPU regression for the injector object level (fails on the old level-mixed code, passes here). - tests/test_eschamoe_decode.py: bit-exact 2-bit decode against a vendor-golden fixture. Evidence (real M5 Max, docs/escha_w2_serving_evidence.md): loads in 5.5s at 12.5 GiB, mtp_enabled=True; a programming prompt returns a correct is_palindrome function; MTP speculative decode is bit-for-bit lossless vs greedy AR (identical output SHA). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…refill split bf16 is the checkpoint's native dtype; the fp32 round-trip was a wasteful cast with no quality change (identical greedy output). int8 uses the fused matvec for decode (M<=32, no dequant) and a transient dequant->GEMM for prefill (M>32) where a dense GEMM tiles the weight far better than a per-row matvec. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
… 2x activation bytes for nothing) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…ILE) Decode at batch=1 is host-encode-bound: the eschamoe on-the-fly 2-bit MoE path (_forward_ondevice, S<=256) issues many tiny Metal launches per token (2 gathers, 3 t128 Hadamard kernels, 2 escha_qmv matvec kernels, silu, casts) and the CPU can't hand them to the GPU fast enough. The decode shape is fixed (S=top_k=8, fixed H/I), so wrap the pure compute in mx.compile: MLX traces the launch chain once and re-issues it as a compact captured graph every token, cutting the per-token Python + kernel-encode overhead. - Extract _forward_ondevice's body verbatim into _escha_decode_compute (minus the trailing reshape, kept in the caller so `lead` never enters the graph). Weights are passed as ARGUMENTS, not closed over, so ONE compiled graph is reused across all 40 layers and every token (one trace per distinct S). - mx.fast.metal_kernel composes with mx.compile, so escha_qmv's custom primitive is embedded in the graph rather than re-encoded per call. - ESCHA_COMPILE=0 restores the exact eager path (A/B + debug escape hatch). Default on. Prefill (S>256) and the int8 non-expert path are untouched. CPU-verified (no GPU): mx.compile is bit-exact (max_abs_diff 0.0) on the real t128/gather/silu/cast chain via a dense stand-in for escha_qmv, and the real compiled fn traces to the correct shape/dtype. Numerical + perf validation on GPU is left to the coordinator. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
At batch=1 the eschamoe decode is host-encode-bound: the runtime's synchronous per-token AR loop (generate_ar) samples on the HOST every step, which forces an mx.eval sync and leaves the token-feedback loop unable to overlap host graph-encode with GPU compute (~46 tok/s vs mlx-lm generate_step's async-pipelined ~56 tok/s at 1024 ctx). Add MTPLX_AR_ASYNC_PIPELINE (default OFF, fail-closed): for the greedy, guard-free serving path it engages a generate_step-style async pipeline (_run_greedy_ar_async) that keeps the feedback loop on-device (argmax) and mx.async_eval's step N+1's forward before reading token N, so the host graph-encode of the next step overlaps the current step's GPU compute. Bit-identical to the synchronous greedy path (async_eval reschedules, never changes values); gated off for any host-feedback feature (constraint grammar, loop/thinking guards, repetition trim, sampled/penalty decoding, ar_return_hidden). The trailing stats/output assembly is factored into a shared _finalize() closure reused verbatim by both paths, so the synchronous path stays byte-identical. CPU-verified (no GPU): synthetic + cyclic-stub checks confirm the async token stream matches the synchronous argmax stream (including mid-stream stop and max_tokens=1), and that the branch engages only when eligible. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…uped prefill Grouped 16-row-tile prefill (444 tok/s @16k) beats the per-row on-device path (276) via decoded-weight reuse, despite its per-layer host grouping sync — measured, so keep it. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
youssofal
left a comment
There was a problem hiding this comment.
Good work. The reverse-engineering writeup (PTX + goldens, cross-validated, 0/3.1M mismatches) is thorough, the vendor-golden fixture tests are the right approach, and the attribution in NOTICE/int8_linear is how borrowing should look. Nothing here is a refusal; this is sequencing plus a set of asks.
Verified on our side: your hermetic injector test fails on current main exactly as you claimed. .mtp lands on the outer wrapper while validate_mtp_support inspects the TextModel, so it is a real pre-existing break for qwen3_5_mtp-type checkpoints (our shipping 35B is unaffected via a different load path, which is why nobody noticed). We'd like to take the injector fix and its two tests as a standalone cherry-pick so it rides the imminent release. Can you split that commit out, or bless us doing the cherry-pick with attribution?
The eschamoe serving lane itself we'd land in the release after. The diff is 1.7k lines into two of the hottest files (generation.py, runtime.py), which have moved ~100 commits since your base, so it needs a rebase plus our own A/B regression gate at native sampling settings before it ships.
Asks for that round:
- Rebase onto the new main once it's tagged.
is_escha_checkpointruns at the head of every_load_base_model. Checkmodel_typebefore touching the filesystem so non-A3B loads don't pay the index read, and lose the bareexcepton the index parse.- Promote the
bench/validate_int8.pycases intotests/(the strided-slice case earned permanent battery coverage).bench/validate_qmv.pyis referenced but not in the diff. - The evidence ladder (49.4 to 56.5 to 58.5) is greedy-measured. Headline numbers need to be at the model's shipped sampling settings, with greedy as a labeled extra. Related:
MTPLX_AR_ASYNC_PIPELINEcurrently helps only greedy requests. Plain temperature/top-p/top-k sampling can stay on-device the same way (penalties genuinely can't), and that follow-up is what makes the +14% real for product traffic. ESCHA_MTP_SHARE_LAYER=20cites an acceptance sweep; put the sweep numbers in the evidence doc.- Your own doc recommends AR until the verify-kernel near-tie flip is fixed. Make the code enforce that (default the escha path to AR) rather than trusting readers.
…ry-pick from PR #242) Cherry-picked from davidtai's PR #242 (commit 039b7a2), injector fix + its hermetic test only, as announced in the review; the eschamoe serving lane lands separately after its rebase. mlx-lm's outer qwen3_5_moe.Model exposes .language_model (no .model or .lm_head); its .language_model (the TextModel) exposes .model + .lm_head (no .language_model). validate_mtp_support inspects _text_model(model).mtp, so the MTP surface must live on the TextModel and be re-exposed on the outer wrapper via delegation, the same dual-level pattern as the generic inject_mtp_support. The previous code set .mtp on the outer model and mixed TextModel-level self.model with outer-level self.language_model.lm_head on one object, so validate failed and forward hit AttributeError: 'no attribute model'. The object level was never exercised by CI; our shipping 35B takes a different load path, which is why nobody noticed. tests/test_qwen3_5_mtp_object_level.py (his hermetic CPU regression, taken verbatim) fails on the pre-fix tree exactly as the PR claimed (verified at tip before applying) and passes here. Escha-specific branches from the original commit (is_escha_qwen3_5_mtp, expert borrowing, norm +1 shift) are deliberately not carried; they depend on escha_load and belong to the lane.
… to Unreleased 2.5.4 shipped with release notes only on GitHub and no CHANGELOG section; backfilled condensed from those notes (warm-turn cache reuse, cache observability #229/#230-half, paged-attention thresholds #228, vision restore cap, adaptive-depth constants, --no-auth #235, timings #237). Unreleased (2.6.0) Fixed gains the session-bank restore aliasing (#247), the streaming tool-arg duplication (#249), the solo-penalty 500 on the composite scheduler, and the qwen3_5_mtp injector object level (PR #242 cherry-pick, credited).
…LFM2, temperature-0 exactness (#212, #235, #239, #242, #245, #247, #249) Speculative decoding is no longer a single-user feature. The new --scheduler-mode mtp_batch serves independent requests through fixed-width MTP cohorts: each row owns its state and sampling, drafts verify in one batched target forward, and rows join and leave mid-flight. Two cohort widths (3-wide and 8-wide) install side by side and each cohort seals at the narrowest width that fits. Against the previous production ar_batch route, concurrent agent workloads decode at 1.6-2.25x per lane on an M5 Max (Qwen3.6-35B-A3B, shipped settings). --mtp-batch-numerics selects throughput/balanced/b1-exact profiles with documented trade-offs and an install-time self-check; per-request stats report each row's own truth. The session bank composes with the cohorts, so agent fleets with shared system prompts keep warm TTFT under concurrency. By David Tai (@davidtai). /v1/embeddings (OpenAI shape, Matryoshka `dimensions`) and /v1/rerank (Cohere/Jina shape) are served by the same daemon, opt-in per model, with lazy loading, an LRU resident cap, idle release, and capability-separated /v1/models so chat pickers are never offered an embedder. Checkpoints shipping their own Python code require --retrieval-trust-remote-code. By @Cyb3rb1ade (PR #212). LiquidAI LFM2 / LFM2.5 serve natively with a bit-exact ShortConv decode fast-path and a verified think/tool grammar. By David Tai (@davidtai). Temperature-0 speculative output is token-identical to plain decoding again: the MTP lane's cold prefill fed the whole prompt in one window while plain decoding splits body + final token; the one-ulp cache difference flipped greedy argmax at a near-tie. All cold-prefill paths now partition identically. Fixes: session-bank prefix-restore corruption (#247), streaming tool-call argument duplication into delta.content (#249), solo penalty requests 500ing on the composite scheduler, `mtplx serve --no-auth` not parsing (#235 follow-through), qwen3_5_mtp validation (#242, @davidtai), artifacts launching at depth ceilings instead of measured depths, --reasoning-parser being silently overridden, the mlx-lm ArraysCache Metal buffer-object leak on long decodes (vendored in-tree), missing MTP heads now degrading to AR serving instead of refusing, and AR-batch hardening (fail-closed cache removal, no completed-stream starvation, MLX-ABI-keyed shader cache). QA: full pytest battery and 567 Swift tests green; Speed-V2 greedy exactness gate 3/3 depths (0/3 on 2.5.4); four-arm perf sweep vs 2.5.2/2.5.3/2.5.4 wheels (same harness, fans-verified, die-temp gated, interleaved) — decode flat-to-faster than 2.5.4, cold TTFT improved; a serial-lane sampling regression introduced mid-cycle was found and fixed during this QA (receipts in the release records).
|
The qwen3_5_mtp validation fix from this PR (MTP surface at the TextModel level) shipped in 2.6.0 with your regression test, credited in the notes. The Escha W2 serving lane itself is still under review here. |
Summary
Native serving for
EschaLabs/Qwen3.6-35B-A3B-Escha-W2— a 2-bit Qwen3.6-A3B checkpoint(
qwen3_5_moehybrid: 30 GDN + 10 full-attn, 256 experts / top-8). It loads and serves through theordinary runtime (
mtplx.runtime.load(path, mtp=True)), decoding the 2-biteschamoeexperts andint8 non-experts on the fly — no transcode, no dense-weight materialization, weights never cached.
Recommended path is autoregressive; MTP spec-decode is experimental (see end).
Improvements (before → after)
mx.compiledecode chain (ESCHA_COMPILE)MTPLX_AR_ASYNC_PIPELINE)qwen3_5_mtpMTP injector object levelAttributeError)Decode @1024, lossless throughout: 49.4 (sync) → 56.5 (async) → 58.5 tok/s (async + compile).
Benchmark — prefill / decode / memory (1024 → 128k)
Example generation
Prompt: Write
is_palindrome(s)(ignore case/spaces/punctuation), with a docstring and two asserts.The base-path fix
inject_qwen3_5_mtp_supportattached the MTP surface at the wrong object level, sovalidate_mtp_supportfailed and forward raisedAttributeError. Fixed to the dual-level pattern(
TextModelcarries.mtp; the outer model delegates), matchinginject_mtp_support. Covered by anew hermetic CPU test. This repairs the base (non-escha)
qwen3_5_mtppath too.Files
mtplx/eschamoe.py+eschamoe_gather.npz(2-bit decode kernel + compiled decode path) ·mtplx/int8_linear.py(fused int8 matvec / GEMM-prefill) ·mtplx/escha_load.py(loader) ·mtplx/qwen3_5_mtp_patch.py(Escha head + object-level fix) ·mtplx/runtime.py(dispatch/load hook) ·mtplx/generation.py(opt-in async AR pipeline, default off) ·tests/(object-level + bit-exactdecode) ·
NOTICE(attribution).MTP status
The native MTP head binds and drafts; short-context spec-decode is bit-exact. Long-context
spec-decode is experimental — its multi-token verify uses a different
scaled_dot_product_attentionkernel than single-token decode, which can flip a near-tie greedy argmax. Serve AR for now.
Attribution
escha_qmvand the int8 tiling are ported/adapted fromdusterbloom/higgs; see
NOTICE.🤖 Generated with Claude Code