Ornith 1.5 + Qwen 3.8 27B hybrid runtime + DFlash speculative decoding on Metal - #844
Open
audreyt wants to merge 272 commits into
Open
Ornith 1.5 + Qwen 3.8 27B hybrid runtime + DFlash speculative decoding on Metal#844audreyt wants to merge 272 commits into
audreyt wants to merge 272 commits into
Conversation
# Conflicts: # gguf-tools/deepseek4-quantize.c # tests/ds4_test.c
# Conflicts: # ds4_server.c
Q2_IMATRIX_FILE now resolves to the aligned imatrix GGUF (cyberneurova-...-chat-v2-imatrix-aligned.gguf) on audreyt/CyberNeurova-DeepSeek-V4-Flash-abliterated-GGUF. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The refreshed m5 prefill numbers reflect the default auto MPP routes (Q8_0 prefill, attention-output low projection, routed-MoE) enabled on M5. Make that explicit in the column header. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Raise the default Metal prefill chunk to 4096 and reuse the range-capable layer-major prefill graph for chunked ranges. Enable the guarded Q8_0 attn_q_b MPP route for <=2048-token prompt batches, dynamic Q8_0 tile width, the routed-MoE fast layout from layer 0, and the RB16 indexed decode path. M5 Max post-patch ds4-bench profile with 64 generated tokens: prompt 443/459/522/486/465 t/s and generation 38.6/38.2/37.6/34.0/33.6 t/s at 0.5k/1k/2k/4k/8k. Tests: make all ds4_test; make test; git diff --check.
Detect macOS Low Power Mode and widen the Q8_0 prefill MPP route only under that condition, while preserving the guarded default for normal-power runs and explicit Q8_0 filters. Low-power M5 Max baseline vs patched auto with 128 generated tokens: 0.5k: prefill 133.46 -> 196.89 t/s, gen 13.53 -> 15.08 t/s 1k: prefill 118.65 -> 188.91 t/s, gen 12.23 -> 14.93 t/s 2k: prefill 130.90 -> 220.33 t/s, gen 11.02 -> 14.65 t/s 4k: prefill 118.09 -> 212.81 t/s, gen 13.25 -> 14.00 t/s 8k: prefill 185.52 -> 206.49 t/s, gen 12.94 -> 13.84 t/s Tests: make all ds4_test; make test; DS4_METAL_MPP_LOW_POWER_DISABLE=1 ./ds4_test --metal-mpp-equivalence; git diff --check.
Carries forward the pending "MPP -> Metal Tensor" naming refactor and adds: - --dump-logits FILE CLI flag and run_logits_dump() so prefill-time logits can be captured for A/B drift comparison. - bench/compare_logit_drift.py + bench/compare_bench.py + run helper. - Macro plumbing in ds4_metal.m's library compile step for five env-gated drift flags (DS4_METAL_HC_STABLE default-on, DS4_METAL_NORM_RSQRT_DISABLE default-on, DS4_METAL_KV_RAW_F32 default-off, DS4_METAL_ROPE_EXP2_LOG2 default-off, DS4_METAL_TENSOR_MATMUL_DISABLE default-off). - Logs the active flag set on first device init so test runs are self-documenting. Per-kernel changes that consume each macro land in follow-up commits so they can be reverted independently if a drift measurement regresses. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The HC=4 and scalar Sinkhorn split paths use 1/(1+exp(-z)) directly, which overflows when z is sufficiently negative (exp(-z) explodes). M5 Max's faster ALU is more likely than M3/M4 to push HC mixer inputs into that regime upstream, so the latent fragility may surface as logprob drift on M5 only. Replaces 1/(1+exp(-z)) with the identity 0.5*tanh(0.5*z) + 0.5 and 2/(1+exp(-z)) with 1 + tanh(0.5*z). Bounded across the full float range. The iter-0 vs iter-1+ epsilon application difference is left intact -- it is mirrored identically in the scalar reference path and appears to be an intentional Sinkhorn warm-up. Gated by DS4_METAL_HC_STABLE so the historical form can be A/B'd. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…ult on) kernel_rms_norm_fuse_impl uses 1.0f/sqrt(mean+eps); the fused kernel_dsv4_qkv_rms_norm_f32_4 was using rsqrt(...) for the same value. Apple Silicon's hardware rsqrt has implementation-defined precision and can differ from 1.0f/sqrt by ~1 ULP. Across the 43 layers of DeepSeek V4 Flash that per-layer ULP drift compounds visibly, and the rounding gap between rsqrt and div+sqrt isn't guaranteed to match between M3/M4 and M5 hardware families. Switch the fused QKV norm to 1.0f/sqrt(...) so both norm kernels share a single formula. Gated by DS4_METAL_NORM_RSQRT_DISABLE so the rsqrt path can be A/B'd. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
kernel_dsv4_kv_fp8_store_f32 deliberately writes the raw cache row as (float)((half)q) so its precision matches the half-typed FlashAttention KV buffer the indexer references. With DS4_METAL_KV_RAW_F32 set, the half cast is skipped and the FP8-dequantized FP32 value is written verbatim. This is diagnostic only: enabling it makes the indexer see higher- precision values than FlashAttention, which is a deliberate mismatch that reveals how much drift the FP16 quantization contributes but is not safe to ship. Default off. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Metal's pow(freq_base, k) is not IEEE-754 strict and the rounding can differ between GPU families. With DS4_METAL_ROPE_EXP2_LOG2 set, the RoPE angle is computed as exp2(k * log2(freq_base)) instead, using two primitives with tighter precision specifications. The change touches both the NeoX and default RoPE branches of kernel_dsv4_rope_tail_f32. Default off -- this is a diagnostic to quantify how much RoPE pow precision contributes to logprob drift on M5 Max relative to M3/M4. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
When the macro un-defines DS4_METAL_HAS_TENSOR at library compile time the cooperative-tensor _mpp kernel templates are no longer in the library, but g_metal4_tensor_api_enabled was still truthy so the host dispatch layer kept attempting to fetch them. The result was a flood of "Metal kernel kernel_mul_mm_*_mpp_* function not found" warnings on the legacy fallback path. Flip g_metal4_tensor_api_enabled = 0 inside the same branch so the host code's ds4_gpu_use_mpp_*() and ds4_gpu_*_mpp_tensor() guards see the disabled state and skip _mpp lookups entirely. Measured on M5 Max with the short reasoning prompt: drift between -mt off and DS4_METAL_TENSOR_MATMUL_DISABLE=1 -mt auto is now exactly zero (rms=0, max_abs=0, max_rank_delta=0), confirming that the M5 Max logprob drift is sourced entirely in the Metal 4 cooperative-tensor matmul codepath and not in HC, norm, RoPE, or KV. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Bisecting the M5 Max logprob drift on -mt auto:
- -mt off baseline: reference
- -mt auto (all routes): rms=0.150, max_abs=0.750, top20=0.263
- -mt auto + DS4_METAL_MPP_Q8_0_DISABLE=1: rms=0, max_abs=0 (exact)
- -mt auto + DS4_METAL_MPP_F16_DISABLE=1: still rms=0.150 (no help)
- -mt auto + DS4_METAL_MPP_ATTN_OUT_DISABLE=1: still rms=0.150
- -mt auto + DS4_METAL_MPP_MOE_{GATE,UP,DOWN}_DISABLE=1: still rms=0.150
The Metal 4 cooperative-tensor Q8_0 matmul (kernel_mul_mm_q8_0_f32_mpp
and direct_rhs variants in dense.metal) is the *sole* drift source on
M5 Max vs the legacy simdgroup_multiply_accumulate path. The other
Tensor routes (F16 compressor, attention-output low projection, routed
MoE gate/up/down) are bit-clean against -mt off.
Flip ds4_gpu_mpp_q8_0_default_target() to return 0 when the device
name contains "M5". Other Tensor routes continue to default on, so the
Q8_0 carve-out preserves the bulk of the Metal Tensor speedup (F16
compressor at layers 0-19, MoE at layers 20+, attn-out at layers 32-42).
Users who care more about prefill throughput than bit-equivalence can
opt back in with DS4_METAL_MPP_Q8_0_ENABLE=1.
Verified on M5 Max with default flags only: -mt auto now produces
exactly the -mt off logits (rms=0, max_abs=0, max_rank_delta=0,
same_top1=yes, top5_overlap=5/5, top20_overlap=20/20).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
MTLCompileOptions.fastMathEnabled defaults to YES and Apple's headers explicitly note this "may violate the IEEE 754 standard". With safe math forced via MTLMathModeSafe (macOS 15+) or fastMathEnabled=NO (deprecated fallback), drift between -mt off and -mt auto on M5 Max shrinks ~4x (rms 0.150 -> 0.037, max_abs 0.75 -> 0.19) -- showing that fast-math optimizations applied differently across the two hardware paths were amplifying the underlying matmul2d divergence. Default OFF: enabling safe math also moves -mt off away from the fast-math production reference (rms=0.63 vs original fast-math baseline) so it isn't a drop-in fix. Useful as a diagnostic to localize remaining drift sources and as an option for users who prefer strict IEEE-754 semantics over fast-math speed. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The previous commit (75f0930) added the M5 carve-out by editing ds4_gpu_mpp_q8_0_default_target(), but that helper was also being reused as the default-target for ds4_gpu_use_mpp_f16_compressor_matmul (line 1363) and for the verbose memory-report banner that prints mpp_f16 (line 2102). That coupled F16 compressor default-on/off to the Q8 carve-out, which is wrong: the per-route bisection showed F16 is bit-clean on M5; only Q8 needed to flip default-off. Introduce a dedicated ds4_gpu_mpp_f16_default_target() that always returns 1 and use it at the two F16 call sites. The Q8 helper keeps its M5 carve-out unchanged. Verified on M5 Max with default flags: -mt auto still produces zero drift vs -mt off (rms=0, max_abs=0, max_rank_delta=0), and the F16 compressor Tensor route is now back to default-on on M5 as intended. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
test_metal_q8_0_mpp_matmul_case() built the reference output by calling ds4_gpu_matmul_q8_0_tensor() after ds4_gpu_set_quality(false). The set_quality(false) call enables MPP routing, and the dispatcher at ds4_metal.m:6277 then routes to ds4_gpu_matmul_q8_0_mpp_tensor() when the MPP can_use gate passes. So on M5 with Metal 4 tensor API enabled, the "reference" was actually the MPP output, and the test compared the MPP kernel to itself -- the max_abs/rms numbers were always near zero and any divergence in the MPP kernel itself would not have been caught. Force ds4_gpu_set_quality(true) around the reference call so the dispatcher takes the legacy simdgroup_multiply_accumulate path, then restore set_quality(false) before invoking ds4_gpu_matmul_q8_0_mpp_tensor() directly for the candidate. The reference and candidate now exercise the two different code paths the test was originally meant to compare. Verified on M5 Max: ./ds4_test --metal-kernels still passes, meaning the M5 cooperative-tensor Q8 matmul agrees with the legacy path within the 0.10 max-abs kernel target on the test shapes. The systemic drift in -mt auto comes from many small matmul deltas compounding through 43 layers, not from any single kernel exceeding the per-call threshold. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…bers
Two corrections triggered by another reviewer's audit:
1. The auto-suite description claimed "auto enables Q8_0 prefill ...";
on M5 that is no longer true now that 75f0930 defaults Q8_0 Tensor
off on M5. Reword the section so it lists F16 compressor, attn-out,
and MoE as the auto-enabled routes, then call out the M5 carve-out
for Q8_0 explicitly with the env-var opt-in.
2. Refresh worst-case suite numbers measured on the current branch
(codex/metal4-m5-drift-patches after the F16-coupling fix 78fa48f
and the test-self-reference fix 580e896) on M5 Max:
worst_rms = 0.169 (was documented ~= 0.170)
worst_top20_max_abs = 0.306 (was documented ~= 0.342)
worst_max_abs = 0.922
min_top5_overlap = 5/5
min_top20_overlap = 20/20 (was 19/20)
worst_rank_delta = 1
Three short fixtures (short_italian_fact, short_code_completion,
short_reasoning_plain) are now bit-exact (rms=0); the residual
drift is concentrated on the two long-context fixtures and comes
from the F16 compressor, attention-output, and routed-MoE Tensor
routes still being default-on, compounding small per-matmul
deltas through 43 layers.
The Q8_0 isolation paragraph also picks up the M5 default-off note
so the env-var docs stay consistent with the runtime behavior.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…ync-pr-15-metal4-m5 # Conflicts: # README.md # ds4_cli.c # ds4_metal.m # ds4_server.c # tests/ds4_test.c
Codex/sync pr 15 metal4 m5
# Conflicts: # README.md # ds4_server.c
# Conflicts: # ds4_server.c
Brings the 10 ornith15 commits (through 39b15f0, fast-forwarded from Min.local) onto main alongside main's upstream DwarfStar sync, the server idle-timeout work and the directional-steering decode loop. Conflict resolutions, all unions rather than side picks: - .gitignore: keep ornith15's /*.gguf (superset of /ds4flash.gguf); no *.gguf path is tracked. - Makefile: clean target gets ornith15's test binaries plus main's $(CUDA_CONFIG) entries; keep the bench-qwen38/test-ornith* targets. - download_model.sh: both sides only added help text and trailing hints (headroom128 DSpark vs qwen/ornith DFlash2); keep both. - ds4.h: both sides added a different eval entry point at the same spot; declare both ds4_session_eval_no_mtp and ds4_session_eval_argmax. - ds4_server.c decode loop: keep main's effective_decode_sampling() refactor (a superset of ornith15's inline think-mode/tool clamping, covered by test_thinking_*), and thread ornith15's greedy_next hand-off through it. The speculative gate keeps ornith15's dflash_ready alternative and batched-mode inference_mu pairing plus main's !dynamic_steering exclusion; the argmax fast path additionally requires !dynamic_steering so steered tokens still go through server_eval_token_no_mtp. - ds4_server.c send_models: keep ornith15's ornith9/ornith/qwen alias lists, but select on the cached s->model_id / s->model_is_glm_dsa instead of re-reading the engine, since --idle-timeout can close the engine while HTTP handlers still serve /v1/models. Verified on the merge result: all five targets link warning-free; ds4_test --server, ds4_agent_test, ds4-eval --self-test-extractors, test_layer_pack, test_engine_mgpu_placement, test_dflash_shape, test_q64a_quant, test_gpu_args, test_layer_stash, test_gpu_args_cli.sh, q4k/mxfp4 dot tests and tests/test_qwen38_bench.sh pass. Also verified by hand that ds4-server binds and serves over IPv6 (::1) and IPv4. Known inherited failure, NOT introduced by this merge: 39b15f0 ("F16-shadow tiled decode attention") breaks Metal decode for every Qwen-family model tested. ed9249f passes and 39b15f0 fails on all three gates below, and the merge result reproduces each 39b15f0 failure with an identical signature: make test-metal-session-batch DS4_TEST_MODEL=Qwen3.8-27B q4_64a FAIL: logits mismatch expected_top=13 actual_top=440 differing=248320 max_abs=13.564 (cheapest reproducer, ~13s) tests/test_ornith15_bench.sh exact-match vs llama-cli tests/test_ornith9_bench.sh exact-match vs llama-completion Greedy generation on Qwen3.8-27B shows the damage directly: ed9249f "The capital of France is Paris.\nThe capital of ..." 39b15f0 "The capital of France is a popular and desirable ..." Note when re-running these gates: the local test model is qwen38-q4_64a-full.gguf (Qwen3.8-27B-Q4_K_M.gguf on Min), arch qwen35, 851 tensors. A synthetic 579-tensor qwen38.gguf lacking context and attention metadata makes test_qwen38_bench.sh fail at "inspect missing train context" and makes test_metal_session_batch pass vacuously.
Brings 889af6b onto main, which repairs the two defects in 39b15f0 that 4944860 inherited and documented as red. Auto-merged with no conflicts; ds4_metal.m and metal/qwen_gdn.metal are now byte-identical to ornith15, and .gitignore keeps main's /*.gguf plus the three gate binaries. 39b15f0's shadow decode kernel never ran: its `constant args &` is declared first (implicit buffer index 0) but the host bound args last at index 8, so every buffer shifted, args.head_dim was garbage, the hd != 256 guard returned from every thread, and the host still reported success -- leaving the caller to skip the legacy attention and emit whatever heads_out already held. A second latent defect dropped the kv_base offset on the k_new/v_new side of the cache-persist block, copying KV head 0 into every head. Both are fixed, and the path is now opt-in (DS4_QWEN_DECODE_SHADOW=1) because F16 history cannot satisfy the bit-exact batched-vs-sequential contract in test_metal_session_batch. Full gate ladder re-run on this merge, all green: make -B ds4 ds4-server ds4-bench ds4-eval ds4-agent no warnings tests/test_ornith15_bench.sh exact match vs llama-cli tests/test_ornith9_bench.sh exact match vs llama-completion tests/test_qwen38_bench.sh Qwen3.8-27B q4_64a make test-metal-session-batch exact_logits=1 ds4_test --server, ds4_agent_test, ds4-eval --self-test-extractors, test_layer_pack, test_engine_mgpu_placement, test_dflash_shape, test_q64a_quant, test_gpu_args, test_layer_stash, test_gpu_args_cli.sh, q4k/mxfp4 dot Both ornith exact-match gates also pass with DS4_QWEN_DECODE_SHADOW=1, so the opt-in fast path is correct on main too. e885392's IPv6 listen change re-verified end to end on the merge result: ds4-server --host ::1 shows TCP [::1]:8793 (LISTEN), answers /health and /v1/models over IPv6, and /v1/models returns the ornith alias list, which also exercises the send_models resolution from 4944860. Local test model note: qwen38-q4_64a-full.gguf (Qwen3.8-27B-Q4_K_M.gguf on Min), arch qwen35, 851 tensors. The synthetic 579-tensor qwen38.gguf fails test_qwen38_bench.sh at "inspect missing train context" and makes test_metal_session_batch pass vacuously.
The generic mul_mv_id family has kernels for Q8_0/Q2_K/Q4_K/Q8_K/ IQ2_XXS/MXFP4 only, so qwen_gpu_moe_ffn skipped its fused call whenever any of gate/up/down was Q5_K or Q6_K and fell through to the per-expert path: three matmuls plus swiglu, scale and add for each selected expert, 24 dispatches per layer instead of 2. The GLM routed-MoE family already covers exactly the missing types (Q2_K/Q4_K/Q5_K gate+up pair, Q2_K/Q4_K/Q5_K/Q6_K down), and its semantics match this shape: unclamped SwiGLU (qwen35moe has swiglu_clamp_exp = 0), the route weight folded into mid (expert_weight_scale = 1), and a down kernel that overwrites out with the full routed sum. Try it after the generic call and before the per-expert fallback; it returns 0 for anything it cannot prove safe, so that fallback stays reachable. moe_experts already has room for the DS4_N_EXPERT_USED * n_ff mid buffer, and the qwen hybrid path has no TP expert split, so the expert-range rebasing is the identity here. This matters for mixed-precision files. The repo's own Ornith-1.5-35B-Q4_K_M has Q6_K ffn_down_exps on 20 of its 40 MoE layers, so half the model was already on the per-expert path; with this change all 40 report fused (DS4_QWEN_MOE_PATH_DEBUG=1). APEX-style Ornith builds are Q5_K/Q6_K on every routed tensor, so today every MoE layer takes that path. Decode throughput on Ornith-1.5-35B-Q4_K_M at 1k ctx is unchanged (15.02 -> 15.02 t/s): at this shape the routed MoE is not the bottleneck, so this removes dispatch overhead rather than buying time. The measured lever there is the shadow decode attention (DS4_QWEN_DECODE_SHADOW=1: 15.01 -> 52.43 t/s), still opt-in because F16 history cannot satisfy test_metal_session_batch's bit-exact batched-vs-sequential contract. Adds DS4_QWEN_MOE_PATH_DEBUG, which reports the path chosen per layer once per process. With per-tensor mixed precision the choice varies layer to layer, so "why is this model slow" needs an answer that does not require a debugger. Gates, all green: build of all five targets warning-free; test_ornith15_bench.sh exact match vs llama-cli (this is the direct correctness proof, since the 20 Q6_K-down layers now take the new kernels); test_ornith9_bench.sh exact match vs llama-completion; test_qwen38_bench.sh; test_metal_session_batch exact_logits=1; ds4_test --server, ds4_agent_test, ds4-eval --self-test-extractors, test_layer_pack, test_engine_mgpu_placement, test_dflash_shape, test_q64a_quant, test_gpu_args, test_layer_stash, test_gpu_args_cli.sh, q4k/mxfp4 dot.
APEX (github.com/localai-org/apex-quant) is a per-tensor recipe, not a new quant type: each tensor is classified by role and given a precision from a layer-wise gradient, so one file mixes several standard GGML types and the type for a role varies by layer index. On mudler/Ornith-1.5-35B-A3B-APEX-GGUF Balanced that is Q5_K routed experts in layers 5-34, Q6_K in 0-4 and 35-39, Q8_0 shared experts on every layer, Q6_K attention and output, and a Q5_K token_embd. Before this, the file did not run: --inspect rejected token_embd, and with that bypassed the decode path died at the first Q5_K layer with "Metal quant matmul received unsupported type/dim (13, in=2048)" -> "Metal MoE failed at layer 5". Three gaps: - tensor_type_is_glm_dense_quant, which weights_validate_qwen_layout uses for token_embd, accepted Q6_K but not Q5_K. - embed_token_any had no Q5_K/Q6_K case, and dflash2_embed_row only knew F32/Q4_K. Added row dequantizers that mirror ds4_vec_dot_q5_K_f32 and ds4_vec_dot_q6_K_f32 element for element, so a gathered row agrees with what the matvec paths compute from the same bytes. - ds4_gpu_embed_token_quant_tensor reported an unsupported quant as "range is outside the mapped model", once per token, on a path whose documented contract is to return 0 and let the host gather the row. That is now a silent fallback and the range error means what it says. The spam also broke the harness comparison: it strips "ds4:.*", which left one word per line and looked like a text mismatch. Combined with the GLM routed-MoE bridge in the previous commit (which is what lets the 30 Q5_K layers run at all), APEX Balanced now matches llama-cli exactly under tests/test_ornith15_bench.sh, with and without DS4_QWEN_DECODE_SHADOW=1. Throughput on APEX Balanced, M5 Max, 32 generated tokens: ctx 1024 prefill 1998 t/s decode 15.30 -> 57.51 t/s (3.8x) ctx 4096 prefill 1682 -> 1993 t/s decode 4.59 -> 32.34 t/s (7.0x) where the decode pair is DS4_QWEN_DECODE_SHADOW off -> on. The shadow stays opt-in for the reason 889af6b gives, and its value grows with context because it replaces an O(ctx) serial scan per token. The routed MoE splits 30 fused / 10 per-expert (DS4_QWEN_MOE_PATH_DEBUG=1): the GLM pair kernels cover Q2_K/Q4_K/Q5_K, so the ten Q6_K-pair layers still take the per-expert path. Not worth a hand-written Q6_K pair kernel yet: converting 20 layers from per-expert to fused on Ornith-1.5-35B-Q4_K_M moved decode 15.02 -> 15.02 t/s, so at this shape the routed MoE is not the bottleneck. Adds DS4_QWEN_MOE_DISABLE_GLM_FUSED to A/B that claim. Only Balanced is covered. Compact and I-Mini put routed experts on Q3_K/IQ2_S and drop the shared expert below Q8_0, which needs kernels that do not exist here yet. Gates, each run alone (two 35B models back to back produces the phantom failure the regression-gate skill warns about): all five targets build warning-free; test_ornith15_bench.sh exact match vs llama-cli on both Ornith-1.5-35B-Q4_K_M and APEX-Balanced, and on APEX-Balanced again with the shadow enabled; test_ornith9_bench.sh; test_qwen38_bench.sh; test_metal_session_batch exact_logits=1; ds4_test --server, ds4_agent_test, ds4-eval --self-test-extractors, test_layer_pack, test_engine_mgpu_placement, test_dflash_shape, test_q64a_quant, test_gpu_args, test_layer_stash, test_gpu_args_cli.sh, q4k/mxfp4 dot.
APEX-MTP-style files bundle the model's MTP/NextN draft head inside the
target GGUF as the trailing block: blk.<n>.nextn.{eh_proj,enorm,hnorm,
shared_head_norm} plus a full MoE layer at blk.<n>.*, with block_count
counting the head. ds4's own Ornith-1.5-35B-Q4_K_M.gguf already has
exactly this shape, and until now the head was dead weight: every qwen
draft path assumes either a separate sidecar file or the dense
mtp.0.* namespace.
qwen_mtp_bind_from now falls back to the bundled layout when the legacy
names are absent, mapping nextn tensors into the existing
qwen_mtp_weights_t (post_attention_norm takes the ffn_norm slot). Two
structural differences are handled in both draft steps:
- eh_proj: the fused projection runs over [enormed; hnormed] (one
matmul over a 2*n_embd concat) instead of separate e/h projections
summed afterwards. The Metal pool gains a concat buffer.
- MoE block: when the head carries routed experts instead of a dense
FFN, the CPU step routes through qwen_compute_routed_experts_cpu and
the Metal step through qwen_gpu_moe_ffn by presenting the bound
tensors as a regular ds4_layer_weights -- identical router, top-k,
weighting and shared-expert semantics by construction.
Engagement stays inside qwen_generate_hybrid (the driver whose batched
verifier rewinds recurrent state correctly); DS4_QWEN_NEXTN_DRAFT=0 is
the kill switch. One hard gate: drafting is enabled only when every MoE
layer's (gate, down) type pair is verifiable by the batched path,
probed up front via the new ds4_gpu_qwen_routed_batch_pair_supported.
The batched verifier rejects unsupported types only after earlier layers
have advanced KV/GDN state, so a mid-round fallback corrupts one token
(observed: first generated token changed); probing first keeps the fail
closed. Consequence: on today's 35B files -- local Q4_K_M (Q6_K down)
and APEX Balanced (Q5_K/Q6_K routed) -- drafting self-disables with an
explanatory line, because the generic mul_mm_id family has no Q5_K/
Q6_K kernels. It engages automatically on files whose pairs are covered
(Q8_0/Q2_K/Q4_K/IQ2_XXS/MXFP4), or once those two kernels land.
Also downgrades the per-round "unsupported routed batch MoE quant types"
abort to warn-once; it fired every verify round and its stderr interleaving
masqueraded as a text mismatch under the harness cleaner.
Gates, serialized where 35B models are involved: build warning-free;
test_ornith15_bench.sh exact match vs llama-cli on both
Ornith-1.5-35B-Q4_K_M and APEX-Balanced; test_ornith9_bench.sh;
test_qwen38_bench.sh unchanged earlier this session;
test_metal_session_batch exact_logits=1; ds4_test --server,
ds4_agent_test, ds4-eval extractors, test_layer_pack,
test_engine_mgpu_placement, test_dflash_shape, test_q64a_quant,
test_gpu_args, test_layer_stash, test_gpu_args_cli.sh, q4k/mxfp4 dot.
Investigated making the F16-shadow decode the default. Two independent
defects in the shadow path surface under tests/test_metal_session_batch's
mixed prefill+decode phase, both producing the same 6.2e-4 logit spread
between a batched decode row and its sequential control (argmax equal):
1. The slot table's incremental converted counter drifts under
interleaved multi-session execution.
2. Forcing a full F32->F16 reconversion per call removes the counter and
STILL fails intermittently -- a state race between the shadow path and
the pooled command-buffer/slot machinery when sessions interleave with
a prefill chunk. Single-session use is unaffected: every exact-match
harness passes with DS4_QWEN_DECODE_SHADOW=1 there.
So the default stays OFF with the kill switch now doubling as the
opt-in; the gate comment records both findings and the two candidate
fixes (per-session shadow isolation, or session-graph-integrated draft
attention like GLM's mtp step). An interim "default on" experiment was
reverted after its initial pass turned out to be a stale test binary.
Also:
- --inspect now reports bundled nextn draft heads honestly:
"layers: 40 (+1 bundled nextn draft head, executed separately)"
instead of claiming 41 model layers.
- The dense-quant rejection message no longer lists a type set that has
since widened ("expected q8_0, q4_K, or q4_0" -> "not a supported
dense quant"); both validators say the same true thing.
Gates: forced full rebuild warning-free; ds4_test --server;
test_metal_session_batch exact_logits=1; test_ornith15_bench.sh exact
match vs llama-cli on Ornith-1.5-35B-Q4_K_M and APEX-Balanced.
Root cause of the batched-vs-sequential drift, found by bounding runs: the shared slot table had 24 slots while three interleaved sessions need thirty (3 caches x 10 full-attention layers), so the prefill path hit its table-full reset mid-phase -- and that reset released every slot's shadow MTLBuffer (last ARC reference) while dispatches referencing them were still encoded into a shared, not-yet-committed command buffer. An intermittent use-after-free: 6.2e-4 logit spread, argmax unchanged, pass/fail flipping with scheduling. The incremental converted-counter drift suspected first was a symptom of the same churn. Fix is two properties, not a rewrite: - reset() no longer releases shadow buffers. They are scratch reused by whoever claims the slot next; reads are bounded by the claiming call's pos, so stale bytes are unreachable. Only process teardown frees them. - capacity 24 -> 64, covering 6 sessions' full-attn layers before any fallback matters. With those, DS4_QWEN_DECODE_SHADOW defaults to ON (=0 restores the legacy F32 scan bit-exact with llama.cpp). Evidence, all post forced rebuild of every binary including the test: - tests/test_metal_session_batch exact_logits=1 five consecutive runs under =1 (previously 0/5), then three more with the default unset. - Exact match vs llama-cli/llama-completion unchanged on Ornith-1.5-35B Q4_K_M, APEX-Balanced and Ornith-1.5-9B. - Unit battery green; test_qwen38_bench.sh green. Throughput on APEX-Balanced, M5 Max, now by default with no env: ctx 1024 decode 55.47 t/s (was 15.3 legacy) ctx 4096 decode 30.04 t/s (was 4.6)
…eduling Small speculative verify rounds (2-9 tokens) cannot take the mul_mm_id path, whose full-tile threshold is 32 rows. They fell back to mul_mv_id, which had no Q5_K/Q6_K kernels, so the up-front capability probe kept the bundled nextn drafter disarmed on every APEX/Ornith 35B file. Add kernel_mul_mv_q5_K_f32_impl and kernel_mul_mv_q6_K_f32_impl by adapting the bit-exact GLM Q5/Q6 down kernels to the generic mul_mv_id argument layout; instantiate the two host-visible variants and register them in the routed pipeline/nr0 selectors. This arms batched verification for Q4_K/Q6_K local35 and Q5_K/Q6_K APEX-MTP. The first scalar attempt used nb00 (block bytes) as row stride and corrupted output; the final kernels advance rows by nb01 exactly like the proven GLM loops. The high-precision APEX-MTP head is nevertheless not profitable on this model/prompt: 0.45-0.50 accepted tokens/round, 32 t/s with auto drafting versus 76 t/s serial. This agrees with mlxfast's own published Qwen3.8 calibration: the unmodified native-MTP depth-2 path is 0.994x serial, and its contract explicitly allows non-drafting rounds. Default K is therefore 0 unless DS4_QWEN_MTP_K is set or DS4_QWEN_MTP_AUTO=1 opts into the current EMA scheduler. That makes the default exactly serial (1.000 ratio, 75.86 t/s on APEX-MTP Balanced) while preserving all drafting machinery for head/scheduler work. Gates: forced builds clean; test_metal_session_batch 3/3 exact_logits=1; test_ornith15_bench exact-match vs llama-cli on local35 and APEX-MTP; APEX-MTP drafting path also exact before scheduler default change.
Speculative decode died on any prompt past 4096 tokens: ds4: unsupported tensor type for dense matvec The nextn head kept a fixed 4096-row private KV (QWEN_MTP_KV_CAP) while decode passes absolute positions, so qwen_mtp_draft_one_metal() refused every draft beyond that row, every caller fell through to the CPU head, and matvec_any() had no BF16 case for the bf16 sidecar weights, so the run aborted mid-generation. Qwen3.8 reports swa=0: the head is a full-attention layer, so its KV has to span the whole session. Size it from the pooled target context (g_qwen_pool.max_ctx, else the effective pool ctx) with a 4096/512 fallback ladder for machines that cannot allocate it, and teach the CPU dense matvec BF16 so the fallback degrades instead of dying. Measured on Qwen3.8-27B Q4_64A + Qwen3.8-27B-MTP-bf16, M5 Max, greedy: 9.5k-token prompt now drafts at 0.70 accepted/round and decodes 14.75 t/s against a 9.7 t/s no-draft baseline; it aborted before. MLX BF16 parity logits stay bit-exact (0/248320 differ).
…d MTP priming - metal/qwen_gdn.metal: kernel_qwen_gqa_attn_decode_shadow_grp dispatches one threadgroup per KV head (16 simdgroups), removing the 16x ngrp DRAM read amplification that caused the >=8k decode cliff (7-8s GPU/token at 10k -> 0.9-1.0s; greedy text byte-identical). DS4_QWEN_SHADOW_GRP=0 reverts. - ds4.c: qwen_mtp_choose_k cap 7 + DS4_QWEN_MTP_H override; qwen_mtp_update_ema stop-token suppression at all call sites; qwen_mtp_priming_batched (256-row window) replaces sequential head priming (DS4_QWEN_MTP_PRIMING_SEQ=1 reverts); DS4_QWEN_POOL_DEBUG, DS4_QWEN_FALLBACK_TALLY instrumentation; early note_requested_ctx. - ds4_metal.m: gated grouped-kernel dispatch; DS4_QWEN_PROFILE_ALWAYS prints per-command-buffer GPU time with aggregates. - next-step.md: Min run results addendum.
Ornith-1.5-35B (GQA group 8) and 9B (group 4) previously fell back to the per-q-head shadow kernel; the grouped path now covers any ngrp<=16 with identical per-head scalar order. Verified: ornith 35B/9B exact-match regressions green, grouped-serial bitwise-equals legacy-serial at ngrp=8, Qwen3.8 golden unchanged. Ornith results appended to next-step.md: 35B serial ~88-90 t/s decode, nextn K=1 +13% (99.8), batched priming 2.1-2.8x on with-head prefill; 9B DFlash break-even-loss documented.
…the grouped kernel dead Qwen3.8-27B decode was collapsing with context on M5 Max: 25.6 t/s at 523 tokens against 10.0 t/s at 10295, while pure MLX stays flat near 24-26. Two defects, one on top of the other. 1. The grouped shadow decode kernel added in e5ce86d never ran on this model. It was gated on `n_head / n_head_kv == 16`, but Qwen3.8's full-attention layers are 24 query heads over 4 kv heads, so ngrp is 6 and every decode silently took the legacy per-query-head kernel. The 3.3x "grouped kernel win" recorded in next-step.md was measured on a machine that loses GPU sustain under load; on Max, toggling DS4_QWEN_SHADOW_GRP moves nothing (9.82 vs 9.79 t/s) and the generated text is byte-identical, which is the tell. The gate now admits any ngrp in 2..16 and the host dispatches ngrp*32 threads instead of a hardcoded 512, so exactly ngrp simdgroups are live and the staging barrier stays uniform. 2. Grouping alone buys nothing here (9.75 vs 9.69 t/s) -- the decode is occupancy-bound, not DRAM-amplification-bound. One threadgroup per kv head is 4 threadgroups per layer per token. Split-K along position (flash-decoding style) partitions [0, pos) into QWEN_DECODE_SHADOW_BT blocks, runs one threadgroup per (kv_head, split), and merges the partials with the same online-softmax combine the single-pass kernel performs inline. Chunk boundaries are BT multiples so each chunk replays the single-pass block order exactly; the current token is folded in the merge so the overall order stays "cache blocks, then k_new/v_new". Measured on M5 Max (median of 3, greedy, decode-only t/s): | prompt | before | after | MLX bar | |--------|--------|-------|---------| | 523 | 24.07 | 26.03 | 26.5 | | 2554 | 18.40 | 25.52 | 29.7 | | 10295 | 9.69 | 25.63 | 24.8 | Decode is now flat in context and beats pure MLX at 10 k. nsplit sweep at 10295: 2/4/8/16/32/41/81 -> 14.6/18.4/21.3/22.9/24.1/24.8/25.9 t/s, so the auto rule takes one BT block per split, capped at 256. Correctness: generated text is byte-identical to the single-pass path at 523, 2554 and 10295 across every nsplit tested, and a --dump-logits pair at 2554 with DS4_QWEN_SHADOW_SPLITK=0 vs default is bit-identical (md5 272d8aa6070c725fe974c882fa5e27af both ways). MLX-exact parity is unchanged at 0/248320 diffs, max_abs 0.0. Full gate ladder green, including exact_logits metal-session-batch and both Ornith exact-match harnesses. DS4_QWEN_SHADOW_SPLITK=0 restores the single-pass grouped dispatch; 2..256 forces a split count for sweeps. DS4_QWEN_SHADOW_DEBUG=1 prints, once per process, which geometry actually ran -- the missing instrument that let a dead kernel look like a 3.3x win. Also in this commit: - MTP verify-width profile counters. DS4_QWEN_MTP_PROFILE=1 now prints mean verify milliseconds per verified width, so the draft-depth price h can be fitted inside a single run instead of across thermal windows. Fitted on Max at 2554: verify(w) = 30.6 + 8.74*w ms against a 37.4 ms plain step. - Top-2 confidence clamp at depths 0 and 1 (DS4_QWEN_MTP_TOP2=1, default off), using the target's top-1/top-2 logit gap from a real k=2 indexer reduction plus two 4-byte reads per position -- no full-logits copy. Measured neutral at 523 (+1.8%, identical text) and negative at 2554 (24.6 vs 26.6 t/s), so it stays opt-in. - The top-2 index buffer lives in g_qwen_pool so teardown frees it; a function-static would dangle after Metal cleanup. Default h stays 0.18: on Max post-split-K it is the best of {0.18, 0.35, 0.59, 0.90} across 523 and 2554 (49.6/30.9 t/s vs 51.4/26.6, 52.6/24.9, 25.9/25.1). Min's 0.10 fit would be worse.
… correction Records that the grouped shadow kernel was gated on ngrp==16 and never ran on Qwen3.8 (24/4 heads), so the previously documented 3.3x long-context win was Min's sustain collapse rather than the kernel; states the split-K numbers and the byte/bit-identity evidence behind defaulting it on; corrects the AUTO measurements that were silently fixed-depth because DS4_QWEN_MTP_K overrides DS4_QWEN_MTP_AUTO; and files the EMA latch that makes the adaptive policy unable to resume drafting as the top open item.
Min reached the same dead-gate conclusion independently and in parallel on 2026-08-24, from the Ornith side rather than the Qwen side: the grouped shadow decode kernel was gated on n_head/n_head_kv == 16, which matches no model that ships here (Qwen3.8 full-attn 24/4 => 6, Ornith-35B 16/2 => 8, Ornith-9B 16/4 => 4), so every one of them silently used the legacy per-query-head kernel. Conflict resolution, all three files: - ds4_metal.m / metal/qwen_gdn.metal: take Min's runtime bound of ngrp [1..16] over the [2..16] arrived at here. Wider is right, because split-K is gated behind `grp` and an ngrp==1 (MHA) model still wants the split even though it has nothing to group. Min's host gate also folds in head_dim == 256, which this side left to the kernel. - Drop the `sgid >= ngrp` guard from both the grouped and the split kernels. Min omitted it deliberately and is right: the host dispatches exactly ngrp*32 threads so it can never fire, and if it ever did it would be a non-uniform return in front of a threadgroup_barrier, which is a worse failure than the out-of-range write it was meant to prevent. The empty-split return in the split kernel keys off sp, which is uniform across the threadgroup, so that one stays. - Keep this side's split-K dispatch structure, which already carried the identical ngrp*32 threadgroup sizing. - next-step.md: this side's rewrite supersedes the old body; Min's Ornith section is folded in whole, since it is measurement this side does not have. Split-K now defaults on for the Ornith models too, so it was measured rather than assumed. At 523 tokens, -n 128, two reps, byte-identical output: | model | ngrp | split-K | SPLITK=0 | gain | |-------|------|---------|----------|--------| | 9B | 4 | 84.86 | 76.37 | +11.3% | | 35B | 8 | 110.61 | 89.76 | +22.4% | Also fixes the DS4_QWEN_SHADOW_DEBUG instrument added on this side. It was warn-once, and warn-once is misleading here: the first shadow decode of a run happens during priming at a tiny pos where nsplit is always 1, so a 35B run that spends every steady-state token in the split path reported "grp nsplit=1". It now reports each distinct (path, nsplit) pair, bounded to 8 lines so it still cannot shred the exact-match harnesses. Full gate ladder green after the merge, including exact_logits metal-session-batch and both Ornith exact-match harnesses.
The adaptive draft-depth policy was an absorbing Markov chain. Once
qwen_mtp_choose_k returned 0, every subsequent round had drafted == 0;
qwen_mtp_update_ema matches none of its update branches in that case, so the
EMA froze at whatever value had stopped the drafting and no amount of
newly-predictable text could ever restart it. Reproduced before the fix:
DS4_QWEN_MTP_AUTO=1 DS4_QWEN_MTP_H=0.35 on the 2554-token prose fixture
accepted 13 tokens across a whole 256-token run (0.07/round) against 113
(0.63/round) at h=0.18, and h=0.90 on the 523-token fixture accepted zero
for the entire run despite the text being highly predictable.
A zero-draft round genuinely carries no evidence, so the only honest escape
is to spend a draft occasionally and re-observe. choose_k now returns depth 1
after `interval` consecutive evidence-free rounds, and both entry points
carry a qwen_mtp_probe { quiet, interval, probing } owned by the caller
alongside its ema[8].
The interval backs off, and -- the non-obvious part -- a probe backs off even
when its draft is accepted. Resetting to the eager interval on any accepted
draft looks right and is wrong: at 10295 tokens probes land ~0.3 tokens/round
while speculation is still a net loss there, so treating acceptance as "this
works" kept re-arming the probe and cost 26%. Accepted is not profitable.
Only a round the cost model itself priced and chose resets the interval; a
probe always doubles it, capped at QWEN_MTP_PROBE_MAX 256. If drafting really
is worthwhile the EMA rises, the cost model starts choosing depth on its own,
and that resets the interval.
Measured on M5 Max, 2 reps. Inert at the shipped default h=0.18 -- identical
output sha AND identical accept counts on all three fixtures, so the shipped
configuration is provably unaffected:
h=0.18 @523 50.88 t/s acc 210 (1.59/rnd) sha 59a3312301db (golden)
h=0.18 @2554 22.94 t/s acc 113 (0.63/rnd) sha 7bcf7cd460ab (pre-fix)
h=0.18 @10295 19.84 t/s acc 58 (0.63/rnd) sha 4589b20f7f04 (pre-fix)
It acts only where the policy was stuck:
h=0.90 @523 latched 23.5 acc 0 -> 36.0 acc 114 +52%, output golden
h=0.35 @2554 latched 24.7 acc 13 -> 23.3 acc 71 -6%, 5.5x acceptance
h=0.90 @10295 latched 24.5 acc 0 -> 21.4 acc 3 -12.5%, backoff-bounded
The 10295 row is the honest cost: where speculation cannot pay, the probe
still samples occasionally. Backoff is what holds that at -12.5% instead of
the -26% the accept-resets-eagerly variant cost.
DS4_QWEN_MTP_PROBE=0 restores the absorbing behaviour bit-for-bit (verified:
h=0.35 @2554 returns to acc 13 / sha b64f80639a09, h=0.90 @523 to acc 0 /
golden). N sets the eager interval, default 4.
Default h stays 0.18: with the latch gone it is still the best of the four
tested values, so the cliff in the h table was the latch rather than the
price, and the real marginal price at depth 1 is ~0.30.
Full gate ladder green, including exact_logits metal-session-batch and both
Ornith exact-match harnesses.
Caveat recorded in next-step.md: no fixture yet exercises a genuine
mid-generation regime change. h=0.90 @523 is a proxy that latches at round 0
because ema[0] inits to 0.85 < 0.90 on uniformly predictable text. The prose
fixtures stay hard throughout, so they can only ever show the probe's cost
and never its benefit -- which is exactly why the first measurement pass
looked purely negative.
A qwen35moe-style draft GGUF carrying blk.<n_layer>.nextn.* tensors (a full MoE draft block) bound nowhere: --mtp called it unsupported and --dflash died in dflash2_bind. qwen_nextn_sidecar_install() now handles that shape on both flags. It probes qwen_mtp_bind_from + qwen_mtp_is_valid before arming g_qwen_mtp_sidecar_model, so anything else falls through untouched: --mtp keeps its legacy-MTP/DSpark paths, and --dflash keeps dflash2 plus the dense-Qwen family gate. On --mtp the nextn attempt sits behind those detections; on --dflash it is tried before dflash2_bind. qwen_mtp_bind() prefers the armed sidecar over the DS4_QWEN_MTP_HEAD env fallback, so drafting engages through the hybrid generate paths while the legacy session verifier stays off. The sidecar model is engine-owned: ds4_engine_close() disarms the pointer before the maps are unmapped.
5c9980e made standalone qwen35moe nextn heads attachable via --mtp and --dflash, but called any structurally valid head "loaded". That is a correctness footgun: the regular Q4_K_M and APEX Balanced Ornith trunks are both qwen35moe 40L/2048 and accept identical nextn tensor shapes, yet the same standalone NVFP4 head behaves completely differently: Q4_K_M + sidecar K=1: accepted 1 / 126 rounds, 52.02 t/s, output differs (serial 107.88-109.19 t/s) APEX + sidecar K=1: accepted 64 / 63 rounds, 54.39-56.87 t/s, output byte-identical to serial Shape and arch cannot distinguish the pair. Classify from the tensor types that produced those measurements, never filenames: trunk Q4_K_M: gate/up routed experts all Q4_K (q4=80) trunk APEX: gate/up routed experts Q5_K/Q6_K (q5=60 q6=20) known head: NVFP4 output readout + Q5_K embedding + Q4_K/Q6_K nextn MoE Pair policy: APEX Q5/Q6 + known NVFP4 nextn ALLOW_MEASURED_GOOD uniform Q4_K + known NVFP4 nextn REFUSE_MEASURED_BAD (not forceable) any other full-MoE nextn pair refuse unless DS4_QWEN_MTP_FORCE=1 legacy dense MTP sidecars unchanged / classifier not applicable Every full-nextn attempt now prints both identities, type counts, path, and a machine-readable verdict. "Loaded" only follows ALLOW. Known-bad refusal prints the measured acceptance/throughput evidence; DS4_QWEN_MTP_FORCE does not override a pair already measured bad. Unknown artifacts remain available for research through the explicit FORCE flag. Cover both attach mechanisms: - qwen_nextn_sidecar_install now receives the already-loaded target trunk and returns -1 for recognized-but-refused pairs, so both --mtp and --dflash abort with the pairing verdict rather than falling through to misleading DFlash/unsupported errors. - DS4_QWEN_MTP_HEAD runs the same check lazily against the target passed to qwen_mtp_bind. A rejected explicit head leaves MTP off and does not silently fall back to the target's bundled head. Behavioral matrix against the real artifacts: APEX + known head, --mtp rc 0, ALLOW_MEASURED_GOOD APEX + known head, --dflash rc 0, ALLOW_MEASURED_GOOD Q4 + known head, --mtp rc 1, REFUSE_MEASURED_BAD Q4 + known head, --dflash rc 1, REFUSE_MEASURED_BAD Q4 + known head + FORCE rc 1, still REFUSE_MEASURED_BAD APEX + unknown full-nextn rc 1, REFUSE_UNKNOWN same + FORCE rc 0, ALLOW_FORCED_UNKNOWN Q4 + known head via env rc 0 serial, MTP off, no bundled fallback dense Qwen3.8 BF16 env head rc 0, legacy path unchanged Full gate ladder green, including exact_logits metal-session-batch and both Ornith exact-match harnesses.
The install probe only checked that blk.<DS4_N_LAYER>.nextn.* tensors bind, which a full qwen35moe file also satisfies (it bundles its head), so --mtp/--dflash happily armed a 24 GB model as the draft head. GPU placement then tried to absorb the exec layers: hundreds of 'Metal F32 tensor matmul range is outside the mapped model' warnings and corrupted output instead of a rejection. qwen_nextn_head_only() scans the tensor table and requires every blk.<k>/ tensor to sit at k == DS4_N_LAYER; anything else falls through to the existing unsupported/dflash paths after printing a specific stderr line. Found by the validation pass with Ornith-1.5-35B-A3B-APEX-MTP-Balanced. Re-verified after the guard: APEX rejected verbatim, the NVFP4 bare head still loads and drafts (drafted=23), legacy-MTP rejection text unchanged, and MLX parity stays 0/248320 with a byte-identical logits dump.
…it-K, probe regime fixture Records the post-latch-fix re-measurement round on Max: - Sidecar validation found a real hole in 5c9980e: the bind probe was satisfied by a FULL model's bundled nextn tensors, so a 24 GB APEX file armed as a 'standalone head' and drafted corrupted output through ~180 Metal range warnings. Documented alongside 46dd3e7's qwen_nextn_head_only() guard and the post-fix re-validation. - Top-2 clamp: deterministic counters show it drops accepted 113->102 at p2k and diverges output bytes under greedy; stays opt-in. - h re-fit: accepts flat across {0.12,0.18,0.26} — no h separates; stays open. Default h=0.18 unchanged. - Ornith split-K @10k: byte-exact on both models, +354.8% (9B) and +78.9% (35B); gain grows with context; default stays on. - 2554 fixed-cost localization blocked by newly-documented Qwen-path window drift (identical binary swung prefill 205->78 between windows while Ornith holds +-0.6%); Ornith-35B reference profile recorded (steady decode CB median 8.51 ms). - Regime-change fixture demonstrates the probe's structural benefit: PROBE=0 latches 0/0 frozen at h=0.90, probe restarts drafting (8 accepted).
…ard work Resolution, both hunks inside qwen_nextn_sidecar_install(): - Take a48cf0c's signature and return contract wholesale (trunk + sidecar params; -1 = recognized-but-refused pair aborts engine open on both --mtp and --dflash; DS4_QWEN_MTP_HEAD lazy check unchanged). - Keep 46dd3e7's qwen_nextn_head_only() structural check AHEAD of the pairing classifier, returning 0 rather than -1. The two answer different questions: a file carrying exec-layer tensors blk.0.. blk.<N>-1 is not a standalone head at all — GPU placement would absorb those layers and corrupt drafting — so it must refuse unconditionally and fall through as not-this-format even under DS4_QWEN_MTP_FORCE. FORCE remains a48cf0c's deliberate escape hatch for unknown *head* pairs; this merge guarantees it can never reopen the full-model corruption path. Post-merge verification (local://merge-verify.md): build green; CPU spot-gates green (ds4_test --server, dflash_shape 8/8, engine_mgpu_placement 101/101). Behavioral matrix as a48cf0c predicts: APEX-full-model via --mtp hits the structural line then the three-format rejection, rc 1, no arming; APEX trunk + NVFP4 head ALLOW_MEASURED_GOOD, drafted=107 accepted=67 coherent; Q4_K_M trunk + same head REFUSE_MEASURED_BAD rc 1 with the measured evidence. MLX parity: /tmp/l3.json is md5-identical to both pre-change dumps (9a0af92fc056a2f0ac05429bb7639269) — bit-exact preserved; the 2409@5e-10 reading from one comparator script is the previously diagnosed comparator-side artifact, not model drift.
|
Hey, any chance we could get MTP support added? That would be awesome — MTP seems way more agent-friendly than DFlash. |
Unaligned Q4/Q6 prefill chunks never hit n128 NAX. Split like Q8_0: aligned prefix on NAX, remainder on tiled mm (DS4_METAL_Q4_NAX_SPLIT=0 restores the old %32 gate). 523-tok longcopy prefill 177 -> 364 t/s, decode text byte-identical to serial. MTP K=3 verify is still ~2.1x serial (~62 t/s), not ranked mlx.fast ~94 t/s. Width-4 GPU is 50 ms vs serial 33 ms. Env-gated probes (xsums QMV, pair SwiGLU, nsg/nxpsg/r1, GDN snap/rows4) stay off; next-step.md has the ABBA log.
Fused gate+up SwiGLU is on for Q4_64A n_tok 2/4/6/8 (PAIR=0 disables). GDN conv/core can write scratch; SNAP=0 restores prefix on reject and swaps on full accept. Longcopy/prose sha-match SNAP=1. Dump still paid so SNAP stays on. ~58 t/s K=3 vs mlx.fast ~94; FFN occupancy is the hole.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Ornith 1.5 + Qwen 3.8 27B hybrid runtime + DFlash speculative decoding on Metal
6 commits against
main@84cc882, 34 files, +15314/−269:6 commits against
main@84cc882, 34 files, +15314/−269:feat(qwen): Qwen3.8 27B family + DFlash2 speculative (squashed, qwen-only)— theQwen 3.8 27B runtime and the DFlash2 block-diffusion drafter, squashed from a
separate branch.
feat(qwen): Ornith 1.5 hybrid runtime + classic DFlash on Metal— the Ornith 1.5(Qwen 3.5) delta on top.
fix(review): harden DFlash binding and restore PR scope— the review fixes keptexplicit: metadata safety, reference-process cleanup, truthful help/defaults, and
removal of 2.4k lines of fork-main server/KV/default-model scope drift.
fix(audit): close remaining Qwen/DFlash correctness gaps— whole-PR self-audit:recoverable binding/rollback, backend and distributed refusal, false-green harness
fixes, Qwen diagnostics/alias correctness, custom-quant regression coverage, and
distillation/requantizer input safety.
perf(qwen): batch DFlash target-layer capture at prefill width— packedcaptured-layer storage at the pool batch width, eliminating the classic-DFlash prefill
cliff and the five-target-layer DFlash2 readback overflow; includes a behavioral Metal
test and full capture/verify failure propagation.
fix(audit): fail closed after recurrent DFlash/target errors— after the batched capture change, failed Metal hybrid/target forwards invalidate checkpoint and draft state so the next request resyncs instead of mixing GDN states.What
(dense) and 35B A3B (MoE) on Metal: batched prefill/decode, per-session KV and GDN
state, checkpoint resume.
--dflash): binds classic (non-DFlash2) draftGGUFs, per-layer persistent draft KV across rounds, causal + sliding-window masks,
Qwen rotate-half RoPE, batched target output-head argmax, one-shot and session
speculative loops. Metal only; other backends error explicitly.
draft), fused gate+up+SwiGLU for 1..8 rows, batched output head, batched Ornith 35B
routed experts. Row-offset correctness verified against target-only output.
training/dflash_distill_mlx.py(MLX) —target-generated corpus, offline target-hidden caching, paper block objective,
full/projection scopes. Published as
audreyt/Ornith-1.5-9B-DFlash-GGUF
(Q4_K_M, 766 MB);
./download_model.sh ornith9-dflash.probes K=1 → K=3 → K=7, keeps speculating while acceptance EMA ≥ 80%, permanently
falls back to target-only on low yield.
DS4_DFLASH_ADAPTIVE=0disables.make cpubuilding (the Metal-only callsites on the Qwen hybrid andDFlash paths are now
DS4_NO_GPU-guarded), tracksds4_dflash2.incas a prerequisiteof the objects that include it, and gitignores root
*.ggufartifacts.Measured
Apple M5 Max (128 GiB), Ornith 1.5 9B Q4_K_M target + the 766 MB distilled classic
draft,
--raw --temp 0.0 -n 128, median of three runs on an idle machine:Speculative output is byte-identical to target-only in both workloads, and with
DS4_DFLASH_ADAPTIVE=0as well.The low-yield row is the cost of finding out: the scheduler spends one probe
(85.1 ms/token against a 17.7 ms target step), measures 0% acceptance and falls
back to target-only permanently. That cost is fixed, not per token — the same
prompt at
-n 512is 39.38 t/s target-only versus 36.49 t/s with the drafterattached (−7.3%), and it keeps shrinking with longer generations. Acceptance on
the base-Qwen3.5 draft was 6/27 on literary text and 10/13 on simple text before
distillation; distillation lifted mean accepted length 2.50 → 2.77
pre-quantization.
Verified on this tree
All on this tree at Apple M5 Max, 128 GiB:
make -B ds4 ds4-server ds4-bench ds4-eval ds4-agentandmake cpu: both link,both zero warnings under
-Wall -Wextra, same as main.tests/test_ornith15_bench.sh: Ornith 1.5 35B A3B output exact-match againstllama-clireference — pass.tests/test_ornith9_bench.sh: Ornith 1.5 9B output exact-match againstllama-completion— pass.tests/test_qwen38_bench.sh: pass against a valid Qwen model; GGUF architecture,layer/context/attention metadata, CPU CSV and native Metal CSV are mandatory. This is
still a smoke test, not output correctness:
llama.cpphas noqwen3_5_textarchitecture, so the Qwen 3.8 side has no exact-match reference on this host.
make test-metal-session-batch DS4_TEST_MODEL=qwen38-escha-w2-q4_k_m.gguf: PASS, sessions=2steps=6 mixed_suffix=8 exact_logits=1.
./tests/test_dflash_shape: 8/8 metadata-bound checks; real cloned GGUFs withhead_dim > 128,n_kv > n_head, non-divisible GQA, or a missing required tensor arerejected through the recoverable bind path.
./tests/test_q64a_quant: Q4_64A and Q2_64A codes are derived from the serialized BF16scale/bias; fixtures distinguish 67 and 18 codes respectively from the old F32 logic.
./ds4_test --long-context --think-tool-recovery --metal-short-prefill --metal-kernels --metal-tensor-equivalence --streaming-decode-prefill-correctness --server: all OK.high-yield probe; distributed
--dflashis rejected before model load.ds4-server+ Ornith 9B + classic DFlash/v1/chat/completionssmoke: 200 JSONresponse; adaptive fallback exercised. A forced 30-second
llama-clitimeout is killedand reaped (
still_live=no)../ds4_test --logprob-vectors --local-golden-vectors --metal-ssd-streaming-cache-pressurefail on this host with the flash-0731fixtures — and fail identically on main @84cc882, same assertions and counts.
--tool-call-qualitypasses on both. Somake testis not clean here, butthis PR does not change it.
Resolved: classic DFlash prefill batching
Commit
f3331e8removes the fixed 8-token capture path. Target-layer activations are nowpacked as
[token][actual captured layer][embedding]in a lazily allocatedbatch_cap × MAX_TARGET × embeddingMetal stash, and each captured layer is interleavedwith one GPU dispatch instead of one blit per token. Both classic capture loops use the
pool's normal prefill batch width.
Apple M5 Max, Ornith 1.5 9B,
speed-bench/promessi_sposi.txt; current values are mediansof three runs:
--dflash--dflashThat is a 10.8×/12.9× improvement for draft-attached prefill, now within normal run
variance of target-only rather than ~12× slower. Decode remains byte-identical (109/109
accepted on the counting probe); the interactive session path is byte-identical too.
The packed actual-layer-count layout also fixes a separate safety defect: DFlash2 uses
five target layers, but the old Metal readback always copied eight slots into host buffers
allocated for five. A guarded-malloc DFlash2 run (
targets=5, 14/14 accepted) and the new37-row × 5-slot × width-13 Metal test both pass.
Scope notes for review
Q4_64A= 36,Q2_64A= 37) withMetal
get_rows/mul_mv_extkernels, a reference implementation ingguf-tools/quants.c, andgguf-tools/gguf_requant_q4_64a.pyto produce such files.They are selected by GGUF content only and are not on any Ornith or DFlash path. Happy
to split them into their own PR if you would rather review them separately.
ds4_cuda.cuandds4_rocm_compat.cuadd unsupported non-Metal entry points so theQwen symbols resolve; they do not implement the hybrid runtime. I have no CUDA or ROCm
toolchain here, so the 22 entry points this PR adds were audited for definition and
matching arity in both backend files, and the new stubs compile standalone under
clang++
-Wall -Wextra. This is not a substitute formake cuda/make rocm.ds4-benchgains--dflash/--dflash-n-max; snapshots are disabled while DFlash state is live so frontier replayresets the draft cache.
were slower than the existing cache-coalesced row kernel on Apple silicon and were
removed. The win comes from projection fusion, not tile reuse.
idle-timeout/queue/health counters, directional-steering policy, KV eviction work,
MTP/NVFP4/deepspec converters, or fork-only agent/contributor notes.