model : add GLM-5.3-Flash (glm5next)

[eauchs]

Overview

Adds glm5next (HF model_type: glm5_next, Glm5NextForConditionalGeneration) —
converter plus text graph. Text-only for now.

Almost nothing here is new: the model is an assembly of subsystems already in tree.

• KDA linear attention (34 of 45 layers) maps onto the ssm_* tensors of
  kimi-linear (conv1d_q/k/v, f_a/f_b, g_a/g_b, beta, a, dt, norm).
• Hyper-connections (mHC) are DeepSeek-V4's, unchanged. Glm5NextTextHyperConnection
  inherits from DeepseekV4HyperConnection with a bare pass, so the Sinkhorn
  formulation carries over as-is, and the checkpoint uses the exact same tensor names
  (hc_{attn,ffn}_{fn,base,scale}). The only delta is the final stream collapse, an
  unweighted mean — which is what dsv4_hc_mean already does, so there are no
  hc_head_* tensors in the checkpoint.
• The k-pool compressor reuses INDEXER_COMPRESSOR_{APE,WGATE} from deepseek4.
• MLA / MoE / NextN come from GlmMoeDsaModel (GLM-5.2).

The DSA indexer is wired now, with k-pool compression. It reuses the existing
llama_memory_hybrid_idx, which gains an idx_row_size constructor parameter and a
set_input_kpool() input; llama_memory_hybrid and llama_kv_cache are untouched.

deepseek4.cpp, glm-dsa.cpp and kimi-linear.cpp are untouched.

Shape: 45 layers (34 KDA + 11 MLA), hidden 4096, MLA with qk_rope_head_dim = 0
(NoPE), 288 experts top-8 sigmoid/noaux_tc plus 1 shared, 3 leading dense layers,
hc_mult 4 with 20 Sinkhorn iterations, SwiGLU clamped at 10.

Additional information

What is not implemented

• Not validated numerically against HF. The fixture checks backend consistency
  and GGUF round-trip, not k-pool fidelity to modular_glm5_next.py.
• Multi-sequence on a unified cache (-kvu): the cross-sequence cell mixing is now
  fixed — set_input_kpool skips cells that do not belong to the stream's sequence, which
  brings whole-pool selection back to the single-stream value (108 -> 236 at -s 1234).
  What remains cannot be read with this metric: count_row assumes cell index == position,
  and on a unified cache the second sequence occupies cells 32..63 at positions 0..31, so it
  is discarded by construction. Attention masking stays correct either way (the KQ mask is
  per token), so there is no cross-sequence leak. Worth knowing when reading reports:
  leaving --parallel unset selects auto slots, which set n_parallel = 4 and
  kv_unified = true, so a plain llama-server invocation is on this path by default.
• Pooling by absolute position, not in cache-array order as HF does: diverges
  after seq_rm or a context shift. This is the container's existing behaviour rather than
  something this PR introduces — mainline set_input_qsa blocks cells the same way
  (b = p/r over cells.pos_get(j)), and set_input_kpool follows it for consistency.
  Instrumented over a seq_rm + seq_add context shift: 144 calls, 0 cells dropped — the
  per-forward recompute realigns it, so the divergence is narrower than the reference's
  cache-array ordering would suggest.
• Footprint: the pool bias is n_kv/index_kpool * n_tps * n_stream * 4 B, so
  ~0.5 GiB at 1M context with n_ubatch = 512 (it was n_kv-sized, four times that,
  before the pooled selection landed). The indexer cache adds
  11 * n_ctx * 256 * 2 B (~5.6 GiB at 1M).
• The MTP draft head is not numerically validated against a reference. HF's
  modular_glm5_next.py ignores layer 45 entirely, so the graph follows GLM-family
  convention (post-norm h, concat order, shared head, plain residuals). @Nokodoko has
  since measured it on GLM-5.3-Flash IQ4_XS (2x RTX 6000 Blackwell): draft acceptance
  51-67% on a mixed prose/code/JSON bench, and the nextn tensors are Q8_0/F32 in that
  quant rather than crushed by quantization.
• index_share_for_mtp_iteration is not implemented. The reference shares the trunk
  index for the MTP step, which a separate MTP context cannot see, so the draft head
  attends densely. Documented at the call site. @Nokodoko measured no acceptance drop
  between a ~200-token and a 5.7k-token context, so this does not look like the dominant
  gap, but it is a real divergence.
• Greedy speculative decoding is not byte-reproducible on this arch. Traced by
  @Nokodoko and not an accept-loop or MTP-graph defect: every accepted token is the
  argmax of its own verify-batch logits. The forward is batch-decomposition sensitive -
  the same 79-token prefix chunked differently moves a top-2 logit gap by ~1 logit, and by
  ~2.7 logits inside a real verify batch, which is enough to flip moderately confident
  greedy decisions. Suspects are KDA chunked-scan accumulation and 288-expert top-k
  routing flips; the DSA indexer is excluded at short context. Any drafting mode is
  affected, so spec-decode losslessness here holds only up to batch-shape numerics.
• Repeating-token collapse past a depth that scales inversely with the configured
  n_ctx. Resolved — cause is in shared ggml code, not this arch. @feni6
  root-caused the Metal deep-context @ collapse on
  #27754 and
  the fix applies here verbatim: mul_mm.metal computes the batched dst offset
  im*ne1*ne0 in int32, which wraps past 2^31 for the dense F32 KQ [n_kv, ub, 64] at
  deep context, and the misdirected stores land on GPU-VA-adjacent K-cache buffers; the
  fix is three uint64_t casts in ggml/src/ggml-metal/kernels/mul_mm.metal. His
  discovery, his patch — confirmed here on M3 Max (third machine class, after his M4 Pro
  and M3 Ultra) with his standalone probe: at n_kv = 108,710, ub = 512, 64 heads (13.27
  GiB output) the first faulty head is 39, exactly ceil(2^31 / (512 × 108,710)), 100 of
  256 checked corners failing, all read at 0; with the casts, 256/256 — including his
  five geometries (61,540 / 67,584 / 81,920 / 98,304 / 108,800) — and the negative
  control (casts removed, rebuilt) fails again at head 39. Non-regression, which nobody
  had measured on Apple Silicon and which matters because mul_mm is a hot kernel every
  model borrows: NMSE unchanged at 8.52e-08 / 2.75e-14 / 4.84e-14, test_mtp draft +
  reload OK on three backends, seven neighbouring archs at FAIL=0. Until it lands
  upstream, -ub 128 is a usable Metal workaround (keeps 63*128*n_kv < 2^31 up to
  n_kv ≈ 266K). Distinct from the UD-IQ1_M garbage above, which reproduces on a
  single sequence.
• No vision.
• Not run against the real weights on my side — the dev machine is a 128 GB M3 Max,
  so my own numbers below are synthetic weights plus shape checks against the remote
  checkpoint. Two people have since exercised the real thing: @tjluyao converted the
  official checkpoint (305.8 GiB, 72 files) — 288-expert stacking and the FP8
  weight_scale_inv dequant both healthy, n_tensors = 1412, total_size = 641.6G at
  BF16, peak RSS ~25 GB so ~32 GB is the practical floor; and @0-F0xtr0t ran the branch
  on CUDA Blackwell (RTX 5090, sm_120), where test-llama-archs -a glm5next gives
  NMSE 8.03e-08. @yakimoto then ran three real unsloth quants on an M1 Ultra: UD-IQ4_XS
  and UD-IQ3_XXS generate coherently on Metal and on CPU, while UD-IQ1_M returns
  repeated-token garbage on both backends — so the 1-bit tier looks numerically dead on
  this arch, and the earlier UD-IQ1_M datapoint is not reproduced. Pick IQ3 or above.

Verification

test-llama-archs -a glm5next:

|        glm5next|Apple M3 Max|   MoE|  OK (8.55e-08)|       OK|
|        glm5next|  Accelerate|   MoE|  OK (2.38e-14)|       OK|
|        glm5next|Apple M3 Max|   MoE|  OK (4.82e-14)|       OK|
|        glm5next|        Meta|   MoE|SKIP           |     SKIP|

The Meta SKIP matches every other recurrent/hybrid arch (kimi-linear, glm-dsa,
bailingmoe3).

DSA path is live. Random-weight fixture, --temp 0, same seed, only index_topk
changes:

topk=4     md5=d6203820
topk=8     md5=e067ce80    <- differs: different selection widths
topk=2048  md5=dc8845e5
topk=4096  md5=dc8845e5    <- identical: both exceed n_kv, width clamps to n_kv

Output changes only when the effective width changes. Between 2048 and 4096 the
hyperparameter changes but the width does not, and the output is bit-identical.

llama-eval-callback shows all six indexer tensors consumed in the executed graph,
with TOP_K(indexer_score_cells-3{256,8}) = {7,8}, matching
width = min(n_kv, topk + kpool - 1) = min(256, 4 + 4 - 1) = 7.

Note: the earlier all-zero fixture could not prove this (117 of its 120 tensors were
zero, so the DSA layer contributed nothing); these numbers use a randomised one.

Pooled DSA selection. The indexer now cuts on the pool axis and expands each
selected pool into its members, instead of running top_k over cells. A reference-free
check counts partially selected pools (credit to #27754 for describing the failure mode
and this metric):

                       whole pools   partial pools   broken rows
  before  Metal            200            129             93     FAIL
  before  Accelerate       179            151             93     FAIL
  after   all backends     236              0              0     OK

236 is the theoretical maximum, so the pool budget is fully spent and no pool is picked
twice. Before the fix the backends disagreed on the partial count while agreeing on the
broken-row count: ggml_top_k does not order ties, so each backend split the same pools
differently. They now agree exactly.

MTP (NextN) draft graph. test_mtp builds the draft head and reloads it:

test_mtp: glm5next, 16 tokens
  Apple M3 Max    draft OK, reload OK
  Accelerate      draft OK, reload OK
  CPU             draft OK, reload OK

The graph is genuinely built, not silently bypassed: patching a GGML_ABORT into
graph_mtp makes the test abort inside graph_reserve. Draft logits are bit-identical
between the synthesized fixture and a real GGUF round-trip loaded with load_mtp.

Fixed along the way: attention.recurrent_layers was read with n_layer_all while the
saver writes per-layer arrays with n_layer(), so any GGUF carrying a NextN block failed
to load ("wrong array length; expected 5, got 4"). Inert for current files, since no
in-tree converter emits that key.

test-save-load-state, which #27755 now runs over every generated arch, passes on
glm5next — all five subtests, no skip:

=== Test 1: baseline ===
63 104 78 98 110 19 25 104 20 14 1 62 42 102 12 7

=== Test 2: sequence removal isolation ===
PASS

=== Test 3: state load ===
63 104 78 98 110 19 25 104 20 14 1 62 42 102 12 7
PASS

=== Test 4: seq copy (host) ===
63 104 78 98 110 19 25 104 20 14 1 62 42 102 12 7
PASS

=== Test 5: seq copy (device) ===
63 104 78 98 110 19 25 104 20 14 1 62 42 102 12 7
PASS

All tests passed.

The restored continuations are identical to the baseline in all three restore paths,
including seq copy on device, so the indexer cache survives save/load and cross-sequence
copy.

Indexer workspace, chunked over tokens — contributed by @matteoscalabrini. The indexer
scores [n_pool, n_head, n_tokens] and materialises that tensor twice, once by the
mul_mat and once by the permute+cont that brings the head axis into ne[0]. That is
2*n_pool*nh*n_tokens*4 B per device, since the DSA layers spread across a layer
split, and it capped ubatch — which with -ot exps=CPU is also what amortises the
expert stream over PCIe, so the cap cost prefill as well as memory. Scoring a token
depends only on its own query row and on pooled, which is shared across the batch, so
the token loop splits into chunks with identical results and ggml-alloc reuses one buffer
across them. The chunk is derived from a fixed scratch target, so short contexts come out
unchunked on the previous code path. Measured on 5x RTX 3090, UD-Q4_K_XL, f16 KV,
experts on CPU:

ub 512  unchunked   8788 MiB/device scratch    75 tok/s prefill @10k
ub 4096 chunked      ~2 GiB/device scratch    272 tok/s prefill @10k

Greedy output is byte-identical to the unchunked graph at matched ubatch, with
--parallel 1, with --parallel 2, and with a ragged final chunk; needle retrieval
passes at 28,788 and 243,314 tokens. The same cont(permute(...)) is present in #27754
and #27773, so the workspace cost is not specific to this PR — the chunking is.

No regression: glm-dsa, kimi-linear, deepseek4, qwen3next and deepseek32 all still pass.
Build is clean (0 errors, 0 warnings, 173 TUs rebuilt, Metal + Accelerate).

End to end on a 4-layer test model (3 KDA + 1 MLA/DSA): HF -> converter -> GGUF ->
llama-cli loads and generates.

Tensor shapes were checked against the real checkpoint over HTTP range requests:
hc_attn_fn [24, 16384], index_kpool_compress_gate [128, 4096],
index_kpool_compress_ape [4, 128].

Shared files touched

file	why
src/llama-context.cpp	one line: glm5next joins the existing n_tokens*40 graph_max_nodes list (non-fused Sinkhorn: 20 iters x 2 sites x 45 layers)
src/llama-graph.{h,cpp}	glm5next added to deepseek4's SwiGLU clamp semantics (clamp the gate before SiLU) in build_ffn and build_moe_ffn; llm_graph_input_mem_hybrid_idx + build_inp_mem_hybrid_idx(); build_attn_mask_top_k() lifted verbatim out of the DSA build_attn overload so both share it; new build_attn overload where top_k == nullptr means a dense mask
src/llama-memory-hybrid-idx.{h,cpp}	idx_row_size constructor parameter (0 keeps the indexer_head_size default), and set_input_kpool(), which fills the pool cells/bias/tail inputs so the top-k picks whole pools
src/llama-model-saver.cpp	guard so dsv4_hc_mult > 0 does not force a DSV4-only key
src/llama-arch.cpp	llm_arch_is_hybrid, llm_arch_supports_sm_tensor -> false
src/llama-model.{h,cpp}	hybrid memory filters, rope type NONE, print_info, LLM_TYPE_312B_A17B

16 files changed, 2326 insertions(+), 46 deletions(-).

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES — the implementation (src/models/glm5next.cpp,
  conversion/glm5next.py) and the test entry were written with Claude Code and GLM 5.3-flash on ds4. Scope,
  design decisions and verification are mine; I ran the build, the arch tests and the
  end-to-end conversion myself.

[ggml-gh-bot]

Hi @eauchs, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

• Large PR: Large changes require prior discussion (e.g. an issue or RFC) and maintainers may not be able to review this PR as-is. Consider splitting it into smaller, focused PRs.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[eauchs]

and no perplexity / top-1 numbers: the checkpoint is 328 GB and my machine
has 128 GB. Looking for someone who can convert the real weights and run a
full check.

[eauchs]

convert_hf_to_gguf.py runs end to end on a 4-layer synthetic model, but the
288-expert stacking, the FP8 weight_scale_inv dequant and the NextN block
have never been exercised on the real checkpoint.

[tjluyao]

Converted the real zai-org/GLM-5.3-Flash checkpoint with this PR (305.8 GiB, 72 files, snapshot 3f1971b7), since the description notes the real weights hadn't been exercised yet. One blocker, then good news.

Blocker: Failed to detect model architecture on the unmodified checkpoint

File "conversion/base.py", line 1182, in __init__
    self.hf_arch = get_model_architecture(self.hparams, self.model_type)
File "conversion/base.py", line 2721, in get_model_architecture
    raise ValueError("Failed to detect model architecture")

The cause isn't in this PR's code — it's an interaction between load_hparams and the text_config merge in TextModel.__init__:

1. config.json has architectures: ["Glm5NextForConditionalGeneration"] at the top level and no architectures key inside text_config.
2. ModelBase.load_hparams goes through AutoConfig.from_pretrained(...).to_dict(), which materialises text_config["architectures"] = None — HF config objects always carry that attribute.
3. TextModel.__init__ then does self.hparams = {**self.hparams, **self.hparams["text_config"]} to move text_config to root. Because the key now exists with value None, it overwrites the valid top-level value.
4. The next get_model_architecture call sees architectures: None with text_config["architectures"] also None, so arch stays None and raises.

Demonstrated directly:

RAW config.json text_config:  'architectures' KEY PRESENT: False
After load_hparams():         'architectures' KEY PRESENT: True  -> value: None

merge {**hparams, **hparams["text_config"]}:
  top-level architectures BEFORE: ['Glm5NextForConditionalGeneration']
  top-level architectures AFTER : None

Instrumenting get_model_architecture shows it called twice — the first call (pre-merge) resolves Glm5NextForConditionalGeneration correctly; the second (post-merge) gets model_type: glm5_next_text, architectures: None and raises.

Workaround: inject text_config.architectures = ["Glm5NextForCausalLM"] into a copy of config.json. Everything downstream then works.

Possible fix — don't let None values clobber during the merge:

self.hparams = {**self.hparams, **{k: v for k, v in self.hparams["text_config"].items() if v is not None}}

Note this isn't GLM-specific in mechanism: any wrapper config whose text_config omits architectures in the JSON should hit it. It just surfaces here because the whitelist a few lines above (StepVLForConditionalGeneration, Sarashina2VisionForCausalLM, ...) doesn't cover Glm5NextForConditionalGeneration, and adding it there wouldn't help anyway since arch is already None by that point on the second call.

With that patched, the conversion works

--vocab-only exports cleanly (tokenizer and chat template both fine), and the full BF16 conversion is running now at ~180/642 GB with no errors. Both paths flagged as unexercised look healthy:

• 288-expert stacking — blk.N.ffn_down_exps.weight ... shape = {2048, 4096, 288}
• FP8 weight_scale_inv dequant — expert tensors arrive as torch.float32, no errors

The tensor census matches the published architecture: 34 ssm_* layers (KDA), 43 MoE layers (45 minus the leading dense), hc_* on all 45, n_tensors = 1412, total_size = 641.6G at BF16.

Memory requirement, re: the 128 GB note in the description

Peak RSS is ~25 GB, driven by a single stacked expert tensor: {2048, 4096, 288} = 2.42B params, which is 9.0 GiB as float32 plus a 4.5 GiB BF16 copy. It was OOM-killed twice on a host with ~17 GiB free before I moved it to a larger box. So 128 GB is comfortable; ~32 GB looks like the practical floor. --use-temp-file did not lower the peak, since the driver is one tensor rather than write buffering.

Happy to report perplexity numbers once the quantization finishes.

[eauchs]

Your fix is pushed — commit 9327de5, exactly your version (filtering out the Nones at merge).
Your correction on memory is right, and I've updated my description: peak RSS is ~25 GB, driven by a single stacked expert tensor, so the practical floor is ~32 GB — not 128.
Two sentences were wrong and I've fixed them thanks to you: "Never run against the real weights" and the note about 128 GB.
And the honest point, the one that matters: my DSA selection is per-cell, which is wrong. The reference picks whole index_topk / index_kpool pools and then expands them. #27754 does it correctly and explains why Jaccard doesn't catch the error (the ReLU puts a lot of pools at exactly 0.0, and ggml_top_k doesn't order ties). If you're about to measure perplexity, you should know this before burning hours of compute — your numbers will be worth more on #27754.

[0-F0xtr0t]

Ran this branch on NVIDIA CUDA (Blackwell, sm_120 / RTX 5090) — adding a data point since the verification table above is Apple-only.

Build: CUDA 12.8, -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=120, clean.

test-llama-archs -a glm5next on CUDA:

|Model arch.|Device                  |Config|  NMSE vs. CPU|Roundtrip|
|-----------|------------------------|------|--------------|---------|
|   glm5next|NVIDIA GeForce RTX 5090 |   MoE| OK (8.03e-08)|       OK|
|   glm5next|CPU (Ryzen 9950X3D)     |   MoE| OK (0.00e+00)|       OK|
|   glm5next|Meta                    |   MoE|          SKIP|     SKIP|

NMSE ~8e-08 vs CPU (well under the 1e-4 bar), roundtrip OK; a second run read 8.79e-08. So the KDA / DSA / mHC path is numerically consistent on CUDA, not only Metal/Accelerate. (Meta SKIP matches the other hybrid archs, as you noted.)

Real weights: loaded unsloth/GLM-5.3-Flash-GGUF UD-IQ1_M (-ngl 99 --cpu-moe) and it generates coherent text with thinking mode intact.

One load snag + a suggested fallback. That GGUF writes the k-pool size under glm5next.attention.indexer.kpool (=4) rather than glm5next.attention.indexer.block_size, so it trips GGML_ASSERT(hparams.indexer_block_size > 0) at src/models/glm5next.cpp:81. --override-kv glm5next.attention.indexer.block_size=int:4 is a workaround, but accepting indexer.kpool as a fallback lets those GGUFs load unmodified:

--- a/src/llama-arch.h
+++ b/src/llama-arch.h
@@ enum llm_kv {
     LLM_KV_ATTENTION_INDEXER_BLOCK_SIZE,
+    LLM_KV_ATTENTION_INDEXER_KPOOL,
     LLM_KV_ATTENTION_INDEXER_LOCAL_BLOCKS,

--- a/src/llama-arch.cpp
+++ b/src/llama-arch.cpp
@@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
     { LLM_KV_ATTENTION_INDEXER_BLOCK_SIZE,           "%s.attention.indexer.block_size"           },
+    { LLM_KV_ATTENTION_INDEXER_KPOOL,                "%s.attention.indexer.kpool"                },
     { LLM_KV_ATTENTION_INDEXER_LOCAL_BLOCKS,         "%s.attention.indexer.local_blocks"         },

--- a/src/models/glm5next.cpp
+++ b/src/models/glm5next.cpp
@@ void llama_model_glm5next::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_ATTENTION_INDEXER_BLOCK_SIZE,  hparams.indexer_block_size, false);
+    if (hparams.indexer_block_size == 0) {
+        // some converters (e.g. unsloth) write the k-pool size under `indexer.kpool`
+        // rather than `indexer.block_size`; accept it so those GGUFs load without an override
+        ml.get_key(LLM_KV_ATTENTION_INDEXER_KPOOL, hparams.indexer_block_size, false);
+    }

With the patch, UD-IQ1_M loads and generates with no --override-kv. Happy to open it as a PR against the branch if that's easier to take.

[eauchs]

Thanks to both of you — this PR now has two independent validations on real hardware.

@tjluyao converted the official checkpoint (305 GiB): stacked experts and FP8 dequant both sound.
@0-F0xtr0t ran it on CUDA Blackwell (RTX 5090) at NMSE 8.03e-08, with real weights loading and generating coherent text.

That's exactly what was missing, and I've updated the description accordingly — it still claimed the code had never been run against the real weights.

@0-F0xtr0t your patch is in (llama-arch.h, llama-arch.cpp, glm5next.cpp): indexer.kpool is now accepted as a fallback for indexer.block_size, so GGUFs written with the old key load without --override-kv. No need to open a separate PR, but shout if I've mangled anything.

[eauchs]

Selection now selects whole pools and expands them—no more top-k on individual cells. Pushed in 0deb55d.

Partial pools dropped from 129 (Metal) / 151 (Accelerate) down to 0 across all three backends. The backends were diverging on partial counts while agreeing on the 93 broken rows, which was the classic signature of a tie-break issue.

Credit to @danielhanchen: the description in #27754 nailed the diagnosis and provided a reference-free metric.

[Nokodoko]

Ran a validation pass against this branch (6ca4915) on GLM-5.3-Flash IQ4_XS
(unsloth UD-IQ4_XS, 2×RTX 6000 Blackwell). Three findings that might help:

1. The nextn head tensors are not being crushed by quantization in this particular
   GGUF — blk.45.nextn.{eh_proj,enorm,hnorm,shared_head_norm} are all Q8_0/F32 in the
   unsloth IQ4_XS build, so "not numerically validated" isn't explained by a naively
   over-quantized head, at least not for this quant.
2. index_share_for_mtp_iteration (true in the GLM-5.3 HF config, index_topk=2048)
   is not implemented — graph_mtp()'s comment already documents this ("the reference
   shares the trunk index for the MTP step, which this separate context cannot see, so the
   draft head attends densely"). Measured effect was smaller than expected (a 5.7k-token-
   context test showed no acceptance drop vs a ~200-token test), so this may not be the
   dominant accuracy gap, but it's a known, real divergence from the reference worth fixing
   eventually.
3. More importantly: measured draft acceptance on a mixed prose/code/JSON bench via
   /v1/chat/completions with --spec-type draft-mtp,ngram-map-k4v was 51–67% (well
   above the 7% this PR's own description reports) — but greedy (temp=0) output was NOT
   byte-identical to --spec-type ngram-map-k4v alone on 2 of 3 fixed test prompts,
   despite ngram-map-k4v alone being provably lossless vs no speculation at all on the same
   3 prompts, same build. That strongly suggests a real (if narrow — divergences look like
   near-tied-logit flips, not gross corruption) verifier/accept-reject bug specific to the
   MTP path, separate from the acceptance-rate question. Happy to help narrow it down with
   a logit dump at the divergence point if useful — didn't have time to instrument that in
   this pass.

[Nokodoko]

Follow-up on the greedy non-losslessness I reported with --spec-type draft-mtp on
GLM-5.3-Flash: root-caused, and it is not a bug in this PR's accept loop or MTP
graph. Instrumenting common_sampler_sample_and_accept_n shows every accepted token is
the argmax of its own verify-batch logits, and the embeddings_nextn trunk-graph change
is numerically inert (bit-identical logits with/without it). The cause is that the
glm5next forward is extremely batch-decomposition-sensitive: the same 79-token prefix,
chunked differently (cache + suffix batch of 1/2/3/5/8 vs one prefill), moves a
top-2 logit gap across a ±1-logit range deterministically, and inside a real verify
batch the same position flipped by ~2.7 logits vs its prefill value — far beyond normal
batched-matmul noise, enough to flip moderately-confident greedy decisions (suspects:
KDA chunked-scan accumulation and 288-expert top-k routing flips; DSA indexer excluded
at short context). Consequence: any drafting speculative mode breaks greedy
byte-reproducibility on this arch (my earlier "ngram-map is lossless" control turned out
to be vacuous — ngram never actually drafted on those prompts). Might be worth
documenting spec-decode losslessness as "up to batch-shape numerics", which glm5next
makes user-visible.

[eauchs]

Rebased on master — new head 24652c57a.

The branch is now rebased on current master. New head is 24652c57a (the SHA you know, 8a8d0bcc4, is the old tip). The state that matters: MERGEABLE, 8 commits, +2378/−46 across 16 files, 0 commits behind master. Your arm B should be running on 24652c57a, not the old SHA.

Two things the rebase brought in:

1. ggml_swiglu_clamp. Master introduced a fused op that replaces the manual clamp + clamp + swiglu_split sequence. I checked in ops.cpp: it computes exactly the same thing — min(gate, limit), clamp(up, -limit, limit), silu(gate) * up. So the FFN path changed shape without changing its result — and that's the path you measure on.

2. The CUDA ceiling reshape is still uncommitted, measured numerically inert on Metal and Accelerate, waiting on your arm B. No urgency there.

Sanity after the rebase:

• test-save-load-state: 111/111
• Non-regression: 7 archs, FAIL=0 — deepseek4, qwen4exp, glm-dsa, kimi-linear, qwen3next, deepseek32, minimax-01
• NMSE glm5next at -s 1234: 8.52e-08 / 2.75e-14 / 4.84e-14 — identical to pre-rebase
• Build: 0 errors, 0 warnings

[eauchs]

@tjluyao a correction on something I told you, not a nudge.

The advice from the 27th — measure perplexity elsewhere — is void now. It was based on a real defect: my DSA selection was per-cell, so any number you produced would not have matched the reference. That is fixed the same day in 0deb55d: partial pools dropped from 129 (Metal) / 151 (Accelerate) to 0 on all three backends, with 236 whole pools — the theoretical maximum, so the pool budget is fully spent and no pool is picked twice.

What the branch has gained since, briefly: head 24652c57a, 8 commits, MERGEABLE, 0 behind master. The indexer chunking from @matteoscalabrini is in (3.6× prefill), a real unified-cache bug is found and fixed (1f817ef), and the CUDA ceiling at 262,140 tokens is diagnosed.

The reason I'm writing: nobody has measured perplexity on this branch. It is the first gap in the description's "what this doesn't cover" list, and you are the only person who has converted the official checkpoint — 305.8 GiB, 72 files, n_tensors = 1412 — so a perplexity number from you would be the first one grounded in the real weights.

No pressure on timing: four days have passed, and you may well have deleted the ~640 GB by now. If you have, or if you'd rather not rerun, a plain "no" is a perfectly good answer and I'll leave it there.

[feni6]

@eauchs Data toward your "one bug, not two" read from #27754: this branch reproduces the depth collapse too, with byte-identical prompts — and its boundary sits a little shallower than the other implementation's.

Setup: M3 Ultra 512 GB, Metal, Unsloth UD-Q4_K_XL, -fa off, one slot, default ubatch, this branch at head 24652c57a. The prompts are the exact synthetic-records probes from the bisection posted on #27754 (same seeds, same token counts), fresh server per probe.

-c	prompt tokens	#27754 @ 2e0e57f	this PR @ 24652c57a
131072	~8K (control)	OK	OK (both planted codes correct; 292 t/s prefill, 15 t/s decode)
131072	96,201	OK	@@@@…
131072	108,710	@@@@…	@@@@…
131072	121,695	@@@@…	@@@@…
262144	165,359	@@@@…	@@@@…
524288	76,064	@@@@…	@@@@…

The shallow control matters: the 27754-converted GGUF loads and answers correctly here (your indexer.kpool fallback does its job), so the deep failures are a genuine depth collapse rather than a conversion mismatch. Caveat all the same: this quant was converted by the other branch's script; a native conversion could shift details.

So: two independent implementations, same failure shape, boundaries within ~10% of each other — which supports the shared-lineage hypothesis over anything specific to either PR's kernels. One more lead from the other thread that may transfer: over there the collapse turned out to be microbatch-dependent (-ub 128 moved the boundary from ~88K out past 248K at -c 262144, though it still failed by ~480K at -c 524288). Whether -ub moves this branch's boundary is untested — happy to run that sweep here if useful, or any specific (n_ctx, ub, depth) points on this hardware.

[matteoscalabrini]

@eauchs Arm B run against the reshape, on 24652c57a — it lifts the ceiling.

arm	n_kv	n_pool	unpatched (Sat)	with reshape (today)
A: prompt 261,861	261,888	65,472	survives, 142.5 t/s	survives, 142.6 t/s
B: prompt 261,963	262,144	65,536	CUDA error: invalid argument, ~1 s	survives, 8 s, coherent

Arm B reuses arm A's warm prefix, so it hits n_kv 262,144 within ~100 fresh tokens — same 1-second reproduction that crashed on Saturday, now completing cleanly with the server healthy after. Numerically inert on CUDA too: greedy output at 2k is hash-identical between the patched and unpatched builds (0c5e1504… both), and the full-depth prefill rate is unchanged (142.6 vs 142.5). 5×RTX 3090, same setup as my earlier reports.

Two incidental datapoints from the same runs: the rebase is speed-neutral at this config (full-depth average identical pre/post), and the depth collapse @feni6 reproduced on this branch does not occur on CUDA with -fa on at -ub 4096 — both arms decode coherent text at 261k+. Backend- or fa-path-dependent, consistent with the fragile-reduction hypothesis.

So from this side: reshape in the series, guard as its own PR, exactly as you proposed. Happy to re-run arm B on whatever lands.

[eauchs]

@matteoscalabrini Arm B validates the reshape — it is in the series as of a few minutes ago.

• Commit 5390dd741 on top of 24652c57a is the reshape, unmodified from the patch you measured (the snippet proposed above, nothing added). Your measurement is in the commit message: n_pool = 65,536 goes from CUDA error: invalid argument to a coherent decode, greedy output hash-identical between patched and unpatched (0c5e1504… both), full-depth prefill 142.6 vs 142.5 t/s. That closes the ceiling the way the diagnosis predicted it would close: the failing dimension was n_pool, and arm B dying in ~1 s off the warm prefix and now completing cleanly is the confirmation that needed no deep batch to reproduce.
• The second commit, 4bc437ab8, declares ssm_a as NOSCAN — the gate multiplies by it instead of running a scan, as qwen3next already declares for the same reason; both spellings write blk.N.ssm_a, so no GGUF changes. Series head is now 4bc437ab8.

Of your two incidental datapoints, the second is the more useful one: the depth collapse @feni6 reproduced on this branch (Metal, -fa off, boundary just under 100K, posted above) does not occur on CUDA with -fa on at -ub 4096 — both arms coherent at 261k+. That is a backend/fa-path split on the same failure shape, exactly the discrimination the fragile-reduction hypothesis needs, and it hands me the cheapest experiment to run on the Metal side: same prompts, -fa on. I have asked @feni6 for exactly that in a separate reply — if it heals there, the fault is in the non-fused attention path and the k-pool is out of cause, and the -ub sweep is unnecessary.

On the split: reshape in the series (done), guard as its own PR (still open on the ggml side) — unchanged. And your offer to re-run arm B on whatever lands is taken: the guard PR is where that offer is most valuable, since it is the ggml-wide fix and your box is the only one that can reproduce the boundary in ~1 s off a warm KV.

[eauchs]

@feni6 This is the datapoint the "one bug, not two" read needed — same failure shape on two independent implementations, boundaries within ~10% of each other, and byte-identical prompts carried over from the #27754 bisection so the comparison is apples-to-apples. The shallow control is worth as much as the deep failures: the 27754-converted GGUF answering correctly at ~8K (both planted codes, 292 t/s prefill) rules out conversion mismatch as the explanation for the deep rows, and incidentally confirms the indexer.kpool fallback does its job on a foreign GGUF.

One request before anything else, and it is cheaper than the sweep: same Metal, same prompt, -fa on. matteoscalabrini just measured that the collapse does not occur on CUDA with -fa on at -ub 4096 — both arms decode coherent text at 261k+ (comment above). If your 96,201-token probe heals with -fa on, the fault is in the non-fused attention path and your k-pool is out of cause — and the -ub sweep is replaced, no need to run it. If it does not heal, -fa is exonerated on Metal and the sweep is back on the table, at which point your offer to run it stands.

Your caveat is noted and I agree with it: the quant came from the other branch's script, so a native conversion could shift the boundary a little — though it would take a lot of shifting to explain a ~10% agreement between two independent implementations.

[eauchs]

Rebased on master, new head a9904cd9a.

The rebase was 25 commits apart, and the only conflict was in tests/test-llama-archs.cpp, which came from #28147 (log_level became int verbosity, original_logger became log_old). test_dsa_kpool and test_mtp are ported to its motif, behavior identical. The two commits you know by SHA were rewritten by the rebase and kept their subjects: the reshape is now 90fe0ba4d, the NOSCAN declaration 518e38800, so please retest against the new head rather than the old SHAs.

The point that matters: #28159 broke the MTP context, and this branch was the only place it could be seen. The hoist moved the n_layer_nextn read into load_hparams, so n_layer() excludes the NextN block by the time the base fills the per-layer arrays. Measured here at n_head_kv[3] = 1 against [4] = 0 (n_ff[3] = 192, [4] = 0, n_layer_all = 5, nextn = 1): the MTP context reduces to that single layer and sized its KV cache from n_head_kv == 0, dying on failed to allocate buffer for kv cache. Fixed by mirroring the last trunk layer into the NextN entries, which is what a scalar key would have broadcast anyway; the local read of the key is now a duplicate of the base and is dropped. Reported upstream on #28159.

Verified state on the new head:

• build: 0 errors, 0 warnings
• glm5next NMSE at -s 1234: 8.52e-08 / 2.75e-14 / 4.84e-14, unchanged
• test_mtp: draft OK, reload OK x3, repaired
• k-pools, 4 passes: 236/0/0, 216/0/0, 236/36/96, 236/0/0
• 7 archs non-regression: FAIL=0
• save-load-state: 111/111
• conversion e2e: 141 tensors, loads and generates with no error

One non-regression footnote so nobody reads it wrong: qwen4exp now shows SKIP on Meta. It gained a PLE (Per-Layer Embeddings) via #27941, and a PLE conv history is a row of the recurrent cache, which linear layers alone have, so the fixture now carries a recurrent cache, and Meta skips every recurrent arch, exactly as it already skips glm5next, kimi-linear and glm-dsa. That is master's doing, not a regression on this side.

[eauchs]

@CISC — closing the promise from #28173: "I'll drop a9904cd once this lands." It landed, and the mirror is gone.

#28173 and #28183 have both merged, the branch is rebased on the merge, and the mirror commit (a9904cd, the trunk-into-NextN hoist workaround) is dropped. New head is c9ddd6821 — 10 commits, MERGEABLE, +2394/−46.

Verified after the drop, on the new head:

• glm5next NMSE at -s 1234: 8.52e-08 / 2.75e-14 / 4.84e-14 — unchanged
• test_mtp: draft OK, reload OK, on three backends
• k-pools, all four passes: identical to pre-drop (236/0/0, 216/0/0, 236/36/96, 236/0/0)
• 7 archs non-regression: FAIL=0
• test-save-load-state: 111/111

Your fix suffices on its own — measured both with and without the mirror on top, and with neither applied the MTP context still dies, sizing its KV cache from n_head_kv == 0. So the mirror was exactly that: a stopgap against a master that had #28159 and not yet this. Nothing is owed to it anymore.

@matteoscalabrini — the gridDim guard is written, and I need you for the half I cannot produce.

It lives on cuda/softmax-griddim-guard (pushed to my fork), +23/−5 in ggml/src/ggml-cuda/softmax.cu, exactly the ggml-wide fix we agreed should stand alone from the series. The mechanism: when ne02 exceeds 65,535, the grid becomes (ne01*ne02, 1, ne03) and the kernel recovers the indices with i02 = blockIdx.x / ne01, i01 = blockIdx.x - i02*ne01. rowx needed no change, because gridDim.x * gridDim.y is ne01*ne02 in either layout — same flat row. The fallback shape follows binbcast.cu:317 and cpy.cu:236.

What I have checked: the index arithmetic, simulated host-side over 34,417,610 blocks and seven geometries, including (128, 65536, 1) — zero divergence between the two layouts.

What I cannot check, and want to be plain about: the compile. No CUDA on my Mac, and my workflows still await maintainer approval, so CI will not compile it for me either. The branch has never seen a compiler. That is why the PR stays unopened until you are back — I'd rather not open code that has never compiled on a PR whose entire argument is that it is clean.

The ask: a compile, and arm B — the one that died in about a second off the warm prefix. Since the series now folds n_pool off ne[2] (074880cbf), arm B has to run on the stack without that commit for the guarded launch to be the one in play. With your measurement in hand, the PR opens with it at the top of the description; without it, there is nothing to open.

[feni6]

Heads-up: the Metal deep-context @ collapse root-caused on #27754 applies to this branch as well — the defect is in shared ggml code, not in either PR's model implementation.

Mechanism (full write-up with standalone repro, GPU-address maps, and an 11/11 end-to-end verification battery: #27754 comment): mul_mm.metal computes the batched dst offset im*ne1*ne0 in int32; glm5next's dense F32 KQ [n_kv, n_ubatch, 64] wraps it past 2^31 at deep context, and the misdirected stores land on the GPU-VA-adjacent K-cache buffers. We measured this branch failing the same class earlier (96,201 tokens @ -c 131072 -ub 512, boundary within ~10% of #27754's — table posted upthread), which is expected since ggml/src/ggml-metal/kernels/mul_mm.metal is identical in both lineages.

The fix is three uint64 casts; it applies to this branch verbatim:

mul_mm.metal: three uint64 casts

diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal
index ee848eed6..df3be0038 100644
--- a/ggml/src/ggml-metal/kernels/mul_mm.metal
+++ b/ggml/src/ggml-metal/kernels/mul_mm.metal
@@ -134,7 +134,8 @@ kernel void kernel_mul_mm(
 
     // Store result tile to output matrix (with batch offset)
     // cT.store handles bounds checking via tD's extents (M, N)
-    device float * dstBatch = (device float *)dst + im * N * M;
+    // int32 im*N*M wraps for KQ [n_kv, ub, 64] past 2^31 elements (glm5next collapse)
+    device float * dstBatch = (device float *)dst + (uint64_t)im * (uint64_t)N * (uint64_t)M;
 
     auto tD = tensor(dstBatch, dextents<int32_t, 2>(M, N), array<int, 2>({1, M}));
     cT.store(tD.slice(ra, rb));
@@ -318,7 +319,7 @@ kernel void kernel_mul_mm(
         // if no bounds checks on the output are needed, we can directly write to device memory
         device float * C = (device float *) dst +
             (r0 + 32*(sgitg &  1)) + \
-            (r1 + 16*(sgitg >> 1)) * args.ne0 + im*args.ne1*args.ne0;
+            (uint64_t)(r1 + 16*(sgitg >> 1)) * (uint64_t)args.ne0 + (uint64_t)im*(uint64_t)args.ne1*(uint64_t)args.ne0;
 
         for (short i = 0; i < 8; i++) {
             simdgroup_store(mc[i], C + 8*(i%4) + 8*args.ne0*(i/4), args.ne0, 0, false);
@@ -337,7 +338,7 @@ kernel void kernel_mul_mm(
 
         if (sgitg == 0) {
             for (int j = tiitg; j < nr1; j += NR1) {
-                device float  * D  = (device float  *) dst + r0 + (r1 + j)*args.ne0 + im*args.ne1*args.ne0;
+                device float  * D  = (device float  *) dst + r0 + (uint64_t)(r1 + j)*(uint64_t)args.ne0 + (uint64_t)im*(uint64_t)args.ne1*(uint64_t)args.ne0;
                 device float4 * D4 = (device float4 *) D;
 
                 threadgroup float  * C  = temp_str + (j*NR0);

Until it's applied, -ub 128 is a usable workaround on Metal (keeps 63*128*n_kv < 2^31 up to n_kv ≈ 266K). Happy to re-run our boundary battery on this branch's head if that's useful for verification.

[eauchs]

@feni6 — your probe compiles against this branch's head (c9ddd6821) and runs here on an M3 Max — third machine class, after your M4 Pro and your M3 Ultra, so these are numbers neither of us had before.

Reproduced on this hardware:

• ./probe --n-kv 108710 --ub 512 --heads 64 — 13.27 GiB of KQ output. First faulty head: 39, exactly ceil(2^31 / (512 × 108,710)). 100 of the 256 checked corners fail, all read at 0.
• With your three uint64_t casts (mul_mm.metal:137, 321, 340): 256/256. Same on the five geometries you listed — 61,540 / 67,584 / 81,920 / 98,304 / 108,800 — 256/256 everywhere.
• Negative control, done: casts removed, rebuild, and the probe fails again at head 39. The test knows how to fail, so the 256/256 is a real pass, not a tautology.

And the piece you hadn't published, because you had no Apple Silicon: the non-regression. mul_mm is a hot kernel every model borrows, so with your casts in: glm5next NMSE at -s 1234 unchanged at 8.52e-08 / 2.75e-14 / 4.84e-14, test_mtp draft + reload OK on three backends, and the seven neighbouring archs at FAIL=0. Nobody had verified that on Apple Silicon until now.

The find and the fix are yours — the write-up, the repro, the diff, all yours. What I did is confirm on a machine you hadn't measured, nothing more. I am not opening the ggml PR: it's your discovery, the PR is yours to open, and the branch here stays untouched — nothing pushed, nothing committed. Meanwhile -ub 128 stands as the Metal workaround you gave until your fix lands upstream. The description's entry on the depth collapse gets replaced, as of this post, with a pointer to your diagnosis and this confirmation.

[feni6]

@eauchs — thank you, this is exactly the confirmation the fix needed, and more. The third machine class matters (all our Apple Silicon data was M4 Pro + M3 Ultra), the negative control makes your 256/256 meaningful in a way a bare green never is, and the non-regression sweep was the one piece we could not produce here — NMSE/test_mtp/neighbouring-archs on Apple Silicon was the open box on the checklist. Head 39 at ceil(2^31/(512×108,710)) on a third machine closes the arithmetic beyond argument.

And thank you for the deferral on authorship — noted, and taken up: the upstream fix PR is now open as #28210, with your M3 Max confirmation (probe, negative control, and the non-regression run) credited in the evidence section. The three casts are byte-identical to the ones you tested at mul_mm.metal:137/321/340.

Agreed on the interim guidance: -ub 128 remains the workaround for this branch and #27754 until #28210 lands. If it would help this PR, once the fix is upstream the description can drop the workaround note entirely — the collapse is a shared-ggml defect, not a property of either model implementation, and both PRs reproduce it for the same reason.

[feni6]

One final note, relayed per the account owner's instruction in their own words: "Don't post anything else on GitHub except for announcing that you're not going to post anymore or make any PRs."

@eauchs — this is that announcement, so you aren't left waiting on us: there will be no further posts or PRs from this account. #28210 stands as submitted, for anyone who wants to carry the fix forward.