Skip to content

model: add GLM-5-Next (GLM-5.3-Flash) - #27754

Open
danielhanchen wants to merge 38 commits into
ggml-org:masterfrom
unslothai:glm5next/upstream
Open

model: add GLM-5-Next (GLM-5.3-Flash)#27754
danielhanchen wants to merge 38 commits into
ggml-org:masterfrom
unslothai:glm5next/upstream

Conversation

@danielhanchen

@danielhanchen danielhanchen commented Aug 26, 2026

Copy link
Copy Markdown
Contributor

Adds support for GLM-5-Next (released as GLM-5.3-Flash), a 321.3B hybrid linear/sparse-attention MoE, plus its vision tower.

Running it

Two flags are currently required for correct output:

  • NVIDIA_TF32_OVERRIDE=0. ggml-cuda/common.cuh sets CUBLAS_TF32_TENSOR_OP_MATH unconditionally, so every fp32 GEMM otherwise runs at 10 mantissa bits. On a fixture this moved top-1 agreement from 0.896 to 0.9995.
  • -fa off. build_attn_mha casts the F32 latent to F16 before ggml_flash_attn_ext, which is the one place MLA cannot afford it.

KV cache type is not a correctness requirement. f16 costs +0.0005 PPL at ctx 2048 and is 0.0018 lower at 4096, both within engine-to-engine noise.

Performance

1x B200, GLM-5.3-Flash UD-IQ1_S, llama-bench -ngl 999 --flash-attn on -ctk f16 -ctv f16 -lm none -p 512 -n 32 -r 3. Before is f30bed8, after is 0069971.

test before t/s after t/s
pp512 1121.80 1118.93
tg32 62.79 63.10
tg32 @ d4096 53.52 59.50
tg32 @ d16384 41.02 57.99
tg32 @ d65536 20.66 48.99

Perplexity over one chunk is unchanged at 3.3612 +/- 0.40158.

MTP / NextN speculative decoding

--spec-type draft-mtp. llama-cli -c 32768 --temp 0 --seed 0 -n 256:

prompt MTP off n=2 n=3 n=5
short 58.6 86.5 80.2 63.7
16K 55.0 77.2

AI Usage

Used Claude and Local Models for testing, iteration and code design - manual verification of model / PR usage

danielhanchen and others added 20 commits August 26, 2026 16:48
Metadata and tensor loading only. The graph entry point throws, as qwen4exp
did at the same stage.

kda.gate_lower_bound is read as required: kimi-k3 selects the softplus branch
when it is absent, which is a different function rather than a missing clamp.

The absorbed MLA projections are 3D, so glm5next joins bailingmoe3 in the MXFP4
carve-out that would otherwise quantize them as expert tensors.
glm5next's mHC is DeepSeek-V4's hyper-connection block: same wide residual,
same 24-row mixer split, same two activations, same Sinkhorn. Only the final
collapse differs, so the graph derives from llama_model_deepseek4::graph and
reuses build_hc_pre / build_hc_post / build_hc_sinkhorn rather than restating
them, as graph_dsv4 already does in dflash.cpp.

dsv4_hc_mean becomes a static member so both archs can reach it; the body and
both deepseek4 call sites are otherwise untouched. The generated code for
deepseek4 is unchanged apart from the endbr64 landing pad the helper now needs
as a global symbol.

The four streams start as exact copies of the token embedding and collapse to
an unweighted mean after the last layer: this checkpoint has no hc_head.

KDA, DSA and the MoE land in later commits, so the two sublayers throw. The
mHC wiring around them is final.
Copy-adapts kimi-k3's KDA layer rather than kimi-linear's or bailingmoe3's: it
already matches on the recurrence ordering, the bounded-sigmoid decay gate and
its branch selection, dt_bias added per channel before the reshape, per-head A
broadcast, SiLU after the conv, f/g/beta read from the pre-convolution hidden
states, and the gated output RMSNorm with a plain weight.

Three differences from kimi-k3. The output gate is low rank, g_b(g_a(x)) as in
kimi-linear, which is what PR 1's converter emits. The q/k L2 eps is a literal
1e-6, the reference's own constant, not f_norm_rms_eps; ggml_l2_norm implements
max(sqrt(sum), eps) rather than sqrt(sum + eps), which at head_dim 128 differs
by about eps/(2*sum) and never trips the clamp, so it is close but not
bit-exact. And the cross-layer residual, latent MoE, situ activation and MLA
output gate have no counterpart here.

The conv follows the reference and convolves q|k|v as one depthwise kernel,
which keeps the conv state a single contiguous block so build_conv_state can
snapshot it. That plus build_recurrent_attn is what makes the layer safe under
recurrent-state rollback, so the arch joins llm_arch_supports_rs_rollback;
without that entry the guard in llama_context silently clamps n_rs_seq to 0.

build_delta_net_autoregressive reshaped a per-channel KDA gate onto ne1, but ne0
is the key axis everywhere else in that function, so it decayed along the value
axis. Invisible for GDN, where the gate is scalar and both spellings produce the
same [1, 1, H_v, n_seqs], and invisible to the shape checks because S_k == S_v.
Fixed rather than asserted around, since glm5next reaches that path on any
backend without the fused operator.

llama_model_deepseek4::graph now derives from llm_build_delta_net_base so
glm5next, which derives from it for the mHC residual, can reach build_delta_net.
The base is a method-only mixin over llm_graph_context with no data members and
no virtuals beyond the destructor llm_graph_context already has; deepseek4.cpp,
dflash.cpp and kimi-k3.cpp compile to byte-identical instructions across the
change.

graph_max_nodes moves the arch to kimi-k3's tier. Measured on the Tiny fixture
with the chunked fallback: 182 nodes plus 15/16 per token for each KDA layer and
46 per layer for the mHC mixers, so the 45-layer model needs 8.3k + 31.9 per
token before DSA or the MoE are counted, which overruns the n_tokens*40 budget.

test-llama-archs synthesised no MLA, hyper-connection, kpool or expert-weight
keys for glm5next, so PR 1's required get_key calls threw out of the sweep and
truncated it at 75 of 143 architectures. The fixture is complete now and the row
is skipped explicitly while the DSA and feed-forward sublayers still throw.
The routing is DeepSeek-V3 noaux_tc exactly as build_moe_ffn already implements
it: sigmoid scores, exp_probs_b added for the top-k SELECTION only, weights
gathered from the unbiased scores, normalised, then scaled by
routed_scaling_factor. n_group and topk_group are both 1, so the group-limited
stage is degenerate and build_moe_ffn's n_expert_groups > 1 guard skips it; no
group keys are written and none are needed.

The clamp is the one thing that needed a change outside this arch. glm5next
clamps the gate max-only and the up symmetrically, both BEFORE the SiLU, which
is what the branch behind the DEEPSEEK4/DFLASH arch gate already does; the else
branch clamps after the SiLU and is a different function. Adding the arch to
both gates reuses it rather than restating it. The two conditions are separate
because the dense path and the MoE path read different hparams arrays.

The leading dense layers clamp too. The reference builds them from the same
Glm5NextTextMLP as the shared expert, so swiglu_limit is not MoE-only, and the
converter already writes swiglu_clamp_shexp for every layer rather than only the
sparse ones. The shared expert is added unscaled.
nope-only MLA in the absorbed form, over every cached position. below
index_topk + index_kpool - 1 resident tokens the indexer selects all of them,
so this is exactly what the sparse path degenerates to, and it is a reference
the sparse commit can be checked against.

the attention half of the hybrid memory becomes the K-only variant: after
absorption the cache holds the kv_lora_rank latent and V is a view of K.
both are required keys for glm5next, so a model saved without them cannot be
loaded back. this is what stops test-llama-archs from round-tripping the arch.
the DSA sublayer no longer throws, so the arch can construct and run. it needs
the MLA head shape as well: with n_head_kv taken from the per-layer array it
would size the K cache row n_head times wider than the latent the graph writes.
index_topk + index_kpool - 1 is the number of positions the indexer keeps, and it
is what makes the dense attention this branch builds exactly equal to the sparse
path below that many cached tokens. an off-by-one in it is invisible to every
output comparison measured so far, on both a dense and a sparse fixture, so it is
checked against a second spelling of the same arithmetic instead.
The DSA layers of this model score pools of index_kpool consecutive positions
rather than single keys, and the pooled key cannot be rebuilt from the MLA
latents. llama_memory_hybrid therefore gains an optional third cache holding one
indexer key and one compressor gate per token, so the hybrid carries the KDA
conv+recurrent state, the MLA latents and the indexer keys at once.

Absent unless filter_idx is given, which defaults to null, so every existing
architecture gets exactly what it got before, state file layout included.

Two heads per cell, not one. GLM's compressor is not a mean pool: it is a
per-channel softmax over the kpool slots with logits gate + ape, where the gate
is a second projection of the hidden state of width indexer_head_size. Caching
it beside the key is the only way a pool survives its member tokens leaving the
batch. Architectures with indexer_kpool == 0 still get one head.

The indexer cache is handed the attention cache's slot layout rather than
finding its own, so the two agree cell for cell, and apply() asserts they do.
It also keeps its own dtype: -ctk q8_0 would otherwise quantise the gates, which
feed a softmax.

llama-kv-cache-kpool.{h,cpp} builds the pool <-> cell map host side. Pools are
defined on positions and cells are whatever find_slot handed out, so the
correspondence cannot be derived in the graph. Nothing here emits a negative
index: ggml_set_rows asserts i1 >= 0, so unpopulated entries are clamped into
range and neutralised by an additive -INFINITY instead.

Two things the map does that the qwen4exp shape it is ported from does not:

  - the top-k budget is indexer_top_k exactly, with the always-selected tail
    biased to -INFINITY so it spends none of it, and forced back in through a
    host-built base mask for the scatter. indexer_top_k is a whole number of
    pools, so the cut lands on a pool boundary; the reference's own output width
    of indexer_top_k + kpool - 1 does not, and ggml_top_k is unordered among
    equals on both CPU and CUDA.
  - one map per ubatch, shared by every indexer layer, since nothing in it
    depends on the layer. Measured on a 16 Ki cell cache with 512 tokens:
    ~4 ms once against ~4 ms x n_layers.

A unified cache with more than one sequence would let two sequences at the same
position pool each other's keys, so create_memory refuses it up front rather
than aborting mid-run.

tests/test-glm5next-memory.cpp: 74 checks, 0 failures, on both the full and the
trunk-only fixture. test-llama-archs is byte identical to the same build without
this commit at a fixed seed: 452 rows, 0 FAIL. Session state files for
qwen3next, falcon-h1, minimax-01, qwen35moe and a real Falcon-H1-0.5B are byte
identical too, across write, reload and rewrite.
Builds the pooled lightning indexer and gives the DSA layers a sparse attention
path driven by it.

Top-k runs over the POOL axis at select_k = index_topk/index_kpool, and the
selected pools are expanded to their member cells through pool_cells. That is
the reference's own two-step (modular_glm5_next.py, Glm5NextTextIndexer.forward:
topk over the pool axis, then selected_indices = pool_indices[batch_idx,
selected]), and it is not interchangeable with a single top-k of width
index_topk over member cells. The argument for the cell-level form - a pool's
members carry its score bit-exactly, so the cut must land on a pool boundary -
assumes tie groups never span pools. They do: ReLU drives most pool scores to
exactly 0.0, and ggml_top_k is explicitly unordered among equals, so the cut
falls inside an inter-pool tie group and splits a pool. Measured on TinySparse
at 512 tokens, the cell-level form leaves a partial pool on 7.51% of query rows
at layer 3 and 5.93% at layer 7; this form leaves none.

The indexer key and gate STORE is unconditional; only the SCORING is gated, on
n_ctx > index_topk + index_kpool - 1. Gating the store the same way would leave
every cell written below n_select with no indexer state, and the first ubatch to
cross n_select would pool cells that were never written.

Nothing here changes any other architecture: test-llama-archs produces a table
byte-identical to the parent's, 300 rows over 143 archs, 0 FAIL.
The tower is the GLM-OCR ViT with a clamped SwiGLU: the gate is bounded
above only, the up projection on both sides, and both before the SiLU.
ggml_swiglu_oai clamps the same way but then adds one to the up branch,
which is a gpt-oss detail this model does not share, so this adds an
FFN_SILU_CLAMP op rather than reusing it.

The clamp sits at the per-block MLP and again at the merger. Both read
hparams.ffn_op, so the graph body stays the GLM-4V one and the pair is
covered together.

It gets its own projector type rather than a flag on glm4v because the
image token limits differ (16/8000 against 8/4096, per the GLM-5.3-Flash
preprocessor) and those are hardcoded per projector, and because the
clamp must stay off for GLM-4V and GLM-OCR.

Also writes clip.vision.spatial_merge_size. No GLM4V-family mmproj has
ever carried it: Glm4VVisionModel skips the Qwen3VL parameters, which is
where it is written, so clip.cpp's hardcoded 2 has been carrying it.

Images only. glm5next spells video with its own token pair and distinct
start/end spans, and that is not handled here.
the vision tower shipped with the shared dynamic-size preprocessor, which is a
qwen-style smart_resize. the 2026-08-26 GLM-5-Next adaptation resizes
differently: both edges are aligned up by ceil rather than round, an over-budget
image is fitted by binary searching the content height for the largest aligned
canvas still within max_pixels, and the resized content is pasted into the
top-left of that canvas rather than centred and stretched to fill it. an image
already at or above min_pixels is never upscaled.

min_pixels/max_pixels stay in tokens. the reference scales them by
temporal_factor * factor**2 and compares against aligned_frames * area, and
aligned_frames equals temporal_factor for a still image, so the two cancel and
hparams.image_min_pixels / image_max_pixels (16 and 8000 tokens, 12544 and
6272000 pixels) are used directly.

glm4v and glm-ocr keep the dynamic-size preprocessor.

images only. video has its own token pair (154855, distinct from the image
token 154854) with its own start/end spans, and is out of scope here.

the resize arithmetic is covered in test-mtmd-impl against values taken from the
reference processor, including the 16- and 8000-token boundaries, extreme aspect
ratios, and inputs where the binary search and smart_resize disagree.
glm4 / chatglm-bpe tokenizer.json files set "ignore_merges": true, meaning a
pre-token that is already a vocab entry is emitted directly and the merge loop
never runs. llama.cpp implements this (llama-vocab.cpp, the get_ignore_merges()
short-circuit) but only enables it for a hardcoded list of pre-tokenizer names,
and glm4 was never added.

Without it the merges are applied - correctly - and reach a different answer,
because greedy BPE cannot always reconstruct a vocab entry from its bytes.
" 王" (Ġçİĭ, id 102322) is the case that exposed it: from Ġ ç İ ĭ the only
merges available are (Ġ,ç)=27944, (ç,İ)=76417 and (çİ,ĭ)=239209, so the lowest
rank wins first and yields Ġç İ ĭ, at which point neither (Ġç,İ) nor (İ,ĭ)
exists and it stops three tokens short. Reaching Ġçİĭ needs (Ġ,çİĭ) at 242943,
which requires never taking (Ġ,ç) at 27944.

The trigger is whitespace immediately before a CJK character, so pure Chinese
prose is unaffected and mixed Chinese-English is not:

  pure Chinese prose        620 vs 620 tokens, already identical
  mixed Chinese-English     680 -> 600 tokens, now identical to HF (-13.3%)
  wikitext-2 (289569 tok)   one divergence -> byte-identical

Found while comparing GLM-5.3-Flash perplexity against transformers, vLLM and
SGLang: the mismatch bounded how many scoring windows could be compared at long
context, and reads exactly like a model-port defect rather than a tokenizer one.
The scripted resolution used for the rebase mangled four files: it spliced a
condition into the middle of graph_max_nodes' multi-line else-if, dropped the
mtmd_image_preprocessor_glm5next declaration, dropped llama-kv-cache-kpool.cpp
from src/CMakeLists.txt (undefined llama_kpool_* and the llm_graph_input_kpool
vtable at link time), and left an "} else {" immediately followed by an
"} else if" in test-llama-archs.

These files are byte-identical between this base and the tree the glm5next
work was verified on, so each is taken from there verbatim.
@ggml-gh-bot

ggml-gh-bot Bot commented Aug 26, 2026

Copy link
Copy Markdown

Hi @danielhanchen, thanks for your contribution!

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

  • PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

  • Multiple open PRs from a new contributor: We limit new contributors (those without a previously merged PR) to 1 open PR at a time. You currently have 4 open PRs.

  • 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.

@github-actions github-actions Bot added model Model specific testing Everything test related mtmd Related to multimodal functionality (video/image/audio) conversion labels Aug 26, 2026
@eauchs

eauchs commented Aug 26, 2026

Copy link
Copy Markdown

you'll probably get the same comment from ggerganov as on #27742 — I have a version where llama_memory_hybrid and llama_kv_cache stay untouched, in #27752 if useful !

deepseek4 sets n_embd_out_impl to hc_mult*n_embd to size its MTP h input.
glm5next inherited that, but our t_embd is build_norm(build_hc_mean(...)),
which is [n_embd, n_tokens]. n_embd_out() therefore reported 4*n_embd while
the tensor held n_embd, and llama-context read n_outputs*n_embd_out floats
out of it, four times what is there.

The assert at that site sizes the destination buffer, so nothing catches the
short source. Only --embeddings and llama_get_embeddings* reach the path,
which is why plain generation never showed it.

Note for when the NextN graph starts consuming h: give MTP its own width
rather than widening n_embd_out again.
The mHC residual mixers, the lightning indexer (selection gate, learned
k-pool position table, and the three indexer projections) and the KDA
recurrence gates are about 1 GiB in total on GLM-5.3-Flash, so the size cost
is noise against a 100-240 GB quant. Quantizing them perturbs which pools
the indexer selects and how much state each KDA step retains, and those
errors compound along a sequence rather than averaging out.

Both spellings are required. The compressor tensors arrived with the
DeepSeek-V4 merge and use an underscore (indexer_compressor_ape / _gate),
while the projections use a dot (indexer.proj / .attn_k / .attn_q_b), so a
single "indexer." prefix test silently misses the compressor pair.

attn_q_a, attn_kv_a_mqa, attn_k_b and attn_v_b are deliberately not listed.
They are precision sensitive too, but the release recipe pins them to q8_0
via --tensor-type, and that is the configuration the shipped quants were
measured in.

Verified with llama-quantize --dry-run q4_k_m on the BF16: all 12 pinned
families report 0 quantized (45 mHC, 12 indexer, 34 KDA each), while
ffn_gate_exps 43/43, attn_q_a 12/12 and attn_output 46/46 still quantize.
The 101-line glm5next block was dropped from test-mtmd-impl.cpp when the
vision work was rebased, even though the commit message still claimed the
resize arithmetic was covered there. It holds the 36-case table over the
16- and 8000-token budget boundaries, including six cases annotated as ones
where a naive smart_resize disagrees, so it is the guard against sliding
back to stretch-to-fill instead of ceil-align plus zero pad.

Restored from 29c096371. test-mtmd-impl now runs 216 assertions, of which
glm5next_resize contributes 185.
@danielhanchen

danielhanchen commented Aug 30, 2026

Copy link
Copy Markdown
Contributor Author

Two changes since the last update. Both measured on 1x B200, GLM-5.3-Flash UD-IQ1_S, all layers offloaded, one GPU, --flash-attn on -ctk f16 -ctv f16.

1. Cache the pooled indexer key in the KV row. The indexer rebuilt every pool's key from the whole cache on every decode step, in all 12 indexer layers. It is now computed once when a pool closes and kept in a third head of that cell's row. llama-bench -lm none -r 3:

test before after
pp512 1121.80 1118.93
tg32 62.79 63.10
tg32 @ d4096 53.52 59.50
tg32 @ d16384 41.02 57.99
tg32 @ d65536 20.66 48.99

2.4x at 64K, and decode falloff from d0 to d65536 goes from -67% to -22%.

2. MTP / NextN speculative decoding, --spec-type draft-mtp. llama-cli -c 32768 --temp 0 --seed 0 -n 256:

prompt MTP off n=2 n=3 n=5
short 58.6 86.5 80.2 63.7
16K 55.0 77.2

+48% at n=2, +40% at n=3 after a 16K prompt. No new GGUF is needed: the NextN block already ships in the released quants (blk.45, nextn_predict_layers = 1) and is only read when --spec-type draft-mtp is passed.

Neither change affects the default path: greedy output is byte identical to before, one chunk perplexity is unchanged at 3.3612 +/- 0.40158 with and without --kv-unified, and test-llama-archs passes for all 147 architectures.

One caveat: on this 1.58bpw quant, greedy output with MTP can diverge from non-speculative decoding partway through a long generation. Both stay coherent. I have not yet separated batched versus single token numerics from an acceptance issue, so treat the equivalence as unproven.

@feni6

feni6 commented Aug 30, 2026

Copy link
Copy Markdown

Quick verification of d07e71e against the depth collapse, since the indexer-key cache rewrites the suspect path: the boundary is unchanged. Same machine and flags as the bisection above (Metal, UD-Q4_K_XL, -fa off, -c 131072), byte-identical prompts to the earlier boundary probes:

depth 2e0e57f d07e71e
96,201 OK (both codes) OK (both codes)
108,701 garbage garbage

So the caching change neither fixes nor shifts the collapse — consistent with your byte-identical greedy check; the defect predates the rewrite and survives it. Whatever produces the n_ctx-dependent boundary is invariant under the pooled-key caching.

The depth-decode win reproduces on Metal, a bit smaller than your B200 numbers at this deeper point: ~4.7 t/s decode at 96K depth on d07e71e vs ~2.5–3 t/s on 2e0e57f (M3 Ultra, Q4_K_XL, ~1.7× at 96K; your 2.4× was at 64K).

@feni6

feni6 commented Aug 30, 2026

Copy link
Copy Markdown

One more datapoint from outside this repo that fits the "pool/bias construction" suspicion: mlx-lm has an open fix for the same architecture family with the same failure shape — silent long-context collapse of the DSA indexer's top-k. There, Indexer.__call__'s per-head score reduction (mx.sum over a [1, 32, 4096, 131072] tensor) silently returns all-zeros past large shapes (MLX bug ml-explore/mlx#3784), so argpartition degenerates to selecting the trailing index_topk positions and retrieval collapses past ~128k with no crash or warning: ml-explore/mlx-lm#1454.

Different framework and kernels, obviously — but it is the same class of defect this thread is circling: a reduction/normalization feeding the indexer selection going numerically wrong only at large tensor shapes, with a context-size-dependent onset. Given our measured boundary moves with n_ctx (which sets the score/mask tensor shapes) rather than with anything else, the Metal reductions feeding the pool scores at these exact shapes (n_kv up to 524288 × n_ubatch 2048, f16 masks) might be worth the same minimal-repro treatment on the ggml side. Happy to run any test-backend-ops case at our failing shapes on the M3 Ultra.

seq_add only skipped the pooled-key rebuild when the shift itself was a
multiple of kpool. That is not sufficient: a pool straddling p0 or p1 keeps
some members and moves the rest, so it is regrouped no matter how the shift
is aligned, and its cached pooled key goes stale while still looking complete.

Both callers pass an arbitrary bound. The server's context shift uses
n_keep + n_discard and its prompt-cache reuse uses the match head, so this is
reachable in normal use: with --keep 39 and n_ctx 8192, n_discard is 4076 and
p0 is 4115, which is a multiple-of-4 shift starting mid-pool.

Also require both bounds to be pool-aligned. A negative p0 or p1 means "from
the start" / "to the end", which no pool can straddle.
@Suaroman

Copy link
Copy Markdown

CUDA SM120 before/after for 0069971 (indexer key caching): confirmed, with a remaining O(n_kv) term on the fa=0 path

Same setup as my earlier comment (2x RTX PRO 6000 Blackwell, driver 580.173.02, CUDA 13.0.2, UD-IQ3_XXS, NVIDIA_TF32_OVERRIDE=0, same llama-bench command). Before is 2e0e57f, after is a175dcd.

test fa=1 before fa=1 after fa=0 before fa=0 after
tg32 70.29 73.50 69.98 72.20
tg32 @ d8192 52.89 66.68 39.13 45.73
tg32 @ d32768 33.24 59.70 17.44 22.81
tg32 @ d65536 22.50 53.97 10.01 13.52
tg32 @ d131072 13.70 45.77 5.26 7.12

With fa=1 the depth term drops from ~0.45 to ~0.063 µs per cached token (3.3x at 131k; d0 to d131k falloff goes from -81% to -38%). That reproduces your B200 result on SM120.

With fa=0 the gain is much smaller: the slope goes from ~1.34 to ~0.97 µs per cached token, and it remains linear all the way out, so fa=0 is now 6.4x slower than fa=1 at 131k. Since -fa off is the configuration the PR currently lists as required for correct output, that residual O(n_kv) cost is worth a look separately from the indexer: it's present only on the non-FA path, so it looks like the masked attention over the full cache rather than anything in the selection.

@Suaroman

Copy link
Copy Markdown

MTP divergence data point on UD-IQ3_XXS, SM120, a175dcd. Same prompt (a 16k-token stage-gated spec for a software-rendered Rubik's cube in HTML/canvas), same build, same quant, temp 1.0. With --spec-type draft-mtp --spec-draft-n-max 2: 84 t/s over a 13.5k-token generation at 16-30k ctx, draft acceptance 0.64, mean len 2.29; output was fluent and self-reported as verified, but the render had a geometry defect (sticker quads with no inset, so faces rendered solid). With MTP off, same prompt, the render was correct on the first attempt. Single pair at temp 1.0 so not conclusive, but consistent with the divergence you noted on IQ1_S; on a 3-bit quant it shows up as a subtle code-logic error rather than incoherence. If needed I can run more pairs or the same with --temp 0 if a deterministic comparison would help. Let me know

danielhanchen added a commit to unslothai/llama.cpp that referenced this pull request Aug 31, 2026
The pin was 50 commits behind the PR head and sat before "Add MTP support",
so the nightly shipped GLM-5-Next without the NextN draft head, without the
master merge, and without the pooled-key shift fix.

Verified against b10705 by replaying the resolve step: the new commit fetches
from ggml-org, is a commit of ggml-org#27754, and merges onto the base plus the pins
listed before it with no conflict.

Unrelated, and not fixed here: ggml-org#25731 stops merging at b10705.
Upstream ggml-org#27960 touched ggml/src/ggml-rpc/ggml-rpc.cpp, which the Inkling
branch also edits, and additive_merge.py correctly refuses it. It merges on
b10698, the base of the last shipped nightly, so the next run on a newer base
will fail there until that branch is merged forward.
@danielhanchen
danielhanchen marked this pull request as ready for review August 31, 2026 10:19
@danielhanchen

Copy link
Copy Markdown
Contributor Author

I know there is another impl, but this impl has been validated and utilized by many folks and works fine.

MTP is also added and long context is checked carefully

Second reduction pass over the arch's comments: 405 comment lines on the
branch's own added lines down to 234, no code changes.

Deletes rather than reshortens. What stayed is limited to things whose absence
would let a reader make a specific mistake: reference constants and sign
conventions, the ordering and precision constraints the graph relies on, and
the shapes of ggml tensors, whose type carries none.
@feni6

feni6 commented Aug 31, 2026

Copy link
Copy Markdown

The collapse is microbatch-dependent, and -ub 128 is a full workaround at every configuration we can test.

First, at the known boundary from the bisection (-c 131072, the byte-identical 108,710-token prompt that reliably produces @@@@… at the default -ub 512):

-ub verdict
128 OK (both planted codes correct)
256 OK (both codes)
512 (default) garbage
2048 garbage

(LLAMA_GRAPH_REUSE_DISABLE=1 at default ubatch still fails, so it is not graph reuse.)

Then -ub 128 across the rest of the failure map — every previously-failing point passes, including far beyond the old boundaries:

-c depth default ub -ub 128
131072 121,695 garbage OK
262144 165,359 garbage OK
262144 248,989 (collapsed region) OK
524288 76,064 garbage OK

All on 2e0e57f, Metal, UD-Q4_K_XL, -fa off, one slot, fresh server per probe, "OK" = both planted codes exactly right and coherent reasoning.

So the failure needs (depth, n_ctx, n_ubatch) jointly — which also reframes the earlier boundary table: those boundaries are where the default n_ubatch = 512 crosses the bad region, and halving to 256 already moves the 128K boundary out of reach. The defect looks like it lives in the per-ubatch construction of the indexer inputs — the [n_kv × n_ubatch] sel/cand mask family or the pool-score path shaped by both — going wrong only when the ubatch dimension and the KV/pool dimension are both large, with n_ctx (via the pool tables) shifting where that happens.

Perf cost of the workaround on Metal is modest: ~25% slower prefill at 249K depth (43.9 vs ~55 t/s), decode unchanged.

Happy to bisect the exact failing (n_kv × n_ubatch) shape or run instrumented builds if useful.

@tobieapb

Copy link
Copy Markdown

Cross-backend datapoint on the depth-dependent repeating-token collapse, in case it helps isolate indexer logic vs backend: on the MLX path (mlx_vlm 0.6.17 glm5_next, which the thread notes has a parallel issue open in mlx-lm), we could not reproduce the collapse through 131K prompt tokens with the context set to 262144.

Setup: M3 Ultra 512 GB, macOS 26.3, pipenetwork/GLM-5.3-Flash-MLX-8bit (8-bit group-64 MLX quant of the BF16 release, MTP layer omitted), served via LM Studio's MLX engine, temperature=0, 255 generated tokens per probe, synthetic English filler + summarization instruction:

prompt tokens TTFT prefill tok/s gen tok/s output
8,256 25.3 s 326 21.2 coherent
32,641 105.2 s 310 21.5 coherent
65,546 268.1 s 244 20.9 coherent
98,149 521.8 s 188 20.3 coherent
131,029 835.1 s 157 19.6 coherent

No single-token runs at any depth; the 96K/128K outputs were the most detailed of the set. This spans the ~65–70K band reported here for -c 524288 and reaches into the 78–253K band reported for -c 262144. We have not yet probed beyond 131K or with larger declared contexts, and this quant lacks the MTP layer, in case either matters. Prefill-rate decay with depth (326→157 tok/s) is consistent with the linear degradation described in the OP.

Happy to run specific depths/configs on this hardware if useful for triangulation.

@feni6

feni6 commented Sep 1, 2026

Copy link
Copy Markdown

Cross-reference: the same boundary prompts reproduce the collapse on the independent implementation in #27752 as well (details and table posted there) — shallow control clean on both, boundaries within ~10%. Supports the shared-lineage indexer hypothesis rather than a defect unique to either PR.

@Suaroman

Suaroman commented Sep 1, 2026

Copy link
Copy Markdown

Vision path validated end to end on real weights (CUDA SM120), via a self-converted mmproj

Since no mmproj is published I converted one from the release checkpoint. model.visual.* lives entirely in model-00062-of-00062.safetensors (1.2 GB, all BF16) so only that shard plus the JSON files were needed. convert_hf_to_gguf.py --mmproj --outtype f16 on a175dcd ran clean: 348 tensors, 1.1 GB, processor_config.json picked up correctly (index file moved aside so the missing text shards don't trip the weight-map check).

Two tests: both with the UD-IQ3_XXS text model, -ngl 999 -fa on, NVIDIA_TF32_OVERRIDE=0, 2x RTX PRO 6000:

  1. llama-mtmd-cli, --temp 0, synthetic 896x448 image (blue square left, red circle right, small black "HELLO 42" top center): described the background, both shapes with correct colors and positions, and read the text exactly. Image encode 76 ms.
  2. llama-server --mmproj, OpenAI image_url data-URL content part, real 1.2 MB JPEG photograph (giraffe on savanna): accurate and specific description (coat pattern, birds perched on the neck, lone tree, overcast sky).

So both the mtmd-cli path and the server path are exercised on real weights, synthetic and photographic input.

I can share the conversion recipe if it's useful for the Unsloth repo.

@feni6

feni6 commented Sep 1, 2026

Copy link
Copy Markdown

Deep-context @ collapse: Metal mul_mm int32 batched dst offset

On Metal, GLM-5.3-Flash (this PR and #27752) emits @ from the first
sampled token once a joint (depth, n_ubatch, n_ctx) threshold is
crossed. CPU-only at the same failing depth is fine. MLX is fine.

Mechanism. Absorbed MLA builds a dense F32 KQ [n_kv, n_ubatch, 64].
mul_mm.metal writes batch im at dst + im*ne1*ne0 with int32
operands. That product wraps past 2^31. First failing head is
ceil(2^31 / (n_ubatch * n_kv)). The wrapped stores land in a contiguous
band below dst ([dst - 2^33 B, dst - (2^32 - S)*4 B), S = total KQ
elements) — measured byte-exact with a sentinel-pool probe — and the
affected region of KQ itself is left unwritten (it reads back stale
buffer content; zeros only on fresh allocations). Stores that fall below
the compute buffer's base silently overwrite other resident MTLBuffers
(verified with sentinel canary buffers). A GPU-address map
(MTLBuffer.gpuAddress) shows the persistent MLA-K and indexer-K cache
buffers allocated directly below the compute buffer (zero gap), and the
projected wrapped-store band intersects those cache buffers at every
measured GARBAGE boundary and misses them at every OK point (11/11
geometries across n_ctx 128K–512K, n_ubatch 128–2048). We describe this
as the geometrically supported fatality path — cache corruption would
also explain why a completed collapse poisons the slot until restart —
though we have not run an interventional test separating it from the
simultaneous stale-tail read.

Standalone repro (no weights, ~150 lines against ggml public API): below.
On M4 Pro, first-fail head matched that formula at n_kv ∈ {61540, 67584,
81920, 98304, 108800} for ub=512. After casting the three dst writes to
uint64_t, 256/256 head-corners pass.

End-to-end verification (synthetic needle-retrieval prompts, byte-identical across runs; "OK" = both planted codes retrieved, "GARBAGE" = endless @):

  • 96,201 @ -c 131072 -ub 512 still OK on the patched build (113.4 t/s)
  • 108,701 @ same config, previously GARBAGE, now OK, both planted codes
  • 88,417 @ -c 262144 -ub 512 (old n_ctx-slope GARBAGE) OK on patched — n_ctx was not causal
  • 177,231 @ -c 524288 -ub 128 unpatched OK — 14.94 GiB sum law is not a correctness boundary
  • 109,753 @ -c 131072 -ub 512 patched OK (independent replicate, second host)
  • 30,274 @ -c 131072 -ub 2048 (unpatched-fatal at 29,788) patched OK
  • 89,669 @ -c 262144 -ub 512 (unpatched-fatal at 88,417) patched OK
  • 72,707 @ -c 393216 -ub 512 (unpatched-fatal at 71,047) patched OK
  • 72,707 @ -c 524288 -ub 512 (unpatched-fatal at ~71,500) patched OK
  • 278,399 @ -c 524288 -ub 128 patched OK, both planted codes (first-wrap regime: 3 wrapped heads unpatched; 41.4 t/s prefill, 114 min) — was reliably GARBAGE unpatched
  • 248,989 @ -c 262144 -ub 128 patched OK, both planted codes (46.5 t/s prefill; a first attempt hit a thinking-length cap and was retried with a larger budget — full response coherent throughout)

Patch: three uint64 casts in ggml/src/ggml-metal/kernels/mul_mm.metal
(tensor-path im*N*M and the two non-tensor im*ne1*ne0 stores). Same
class, not on this model's KQ path: mul_mm_id ~716 and mul_mv Q4_0 ~292.

-ub 128 remains a valid workaround on unpatched builds because it keeps
63 * 128 * n_kv below 2^31 up to n_kv ≈ 266K.

Standalone repro: mul_mm_metal_probe.cpp
// Standalone Metal batched-mul_mm offset probe for the glm5next dense KQ shape.
//
// This constructs the same logical operation as absorbed MLA attention:
//   A [K, n_kv, 1] F16 x B [K, ub, n_head] F32
//       -> KQ [n_kv, ub, n_head] F32
//
// K is deliberately only 64, so the real 10-14 GiB output geometry is tested
// without loading the model.  Every output element in head h should equal
// K*(h+1).  Reading the corners of every head detects missing, wrapped, or
// overlapping batch writes without copying the huge result back to the host.

#include "ggml.h"
#include "ggml-backend.h"
#include "ggml-metal.h"

#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <string>
#include <vector>

static void die(const char * msg) {
    std::fprintf(stderr, "FATAL: %s\n", msg);
    std::exit(1);
}

static size_t checked_bytes(int64_t a, int64_t b, int64_t c, size_t elem) {
    if (a <= 0 || b <= 0 || c <= 0) die("dimensions must be positive");
    const uint64_t n = (uint64_t) a * (uint64_t) b * (uint64_t) c;
    if (n > SIZE_MAX / elem) die("tensor byte size overflows size_t");
    return (size_t) n * elem;
}

int main(int argc, char ** argv) {
    int64_t n_kv = 1024;
    int64_t ub = 64;
    int64_t n_head = 4;
    int64_t k = 64;

    for (int i = 1; i < argc; ++i) {
        const std::string a = argv[i];
        if (a == "--n-kv" && i + 1 < argc) n_kv = std::atoll(argv[++i]);
        else if (a == "--ub" && i + 1 < argc) ub = std::atoll(argv[++i]);
        else if (a == "--heads" && i + 1 < argc) n_head = std::atoll(argv[++i]);
        else if (a == "--k" && i + 1 < argc) k = std::atoll(argv[++i]);
        else if (a == "-h" || a == "--help") {
            std::printf("usage: %s [--n-kv N] [--ub N] [--heads N] [--k N]\n", argv[0]);
            return 0;
        } else {
            std::fprintf(stderr, "unknown or incomplete argument: %s\n", a.c_str());
            return 2;
        }
    }

    const size_t a_bytes = checked_bytes(k, n_kv, 1, sizeof(ggml_fp16_t));
    const size_t b_bytes = checked_bytes(k, ub, n_head, sizeof(float));
    const size_t out_bytes = checked_bytes(n_kv, ub, n_head, sizeof(float));
    std::printf("shape: A=[%lld,%lld,1] f16 B=[%lld,%lld,%lld] f32 "
                "KQ=[%lld,%lld,%lld] f32 (%.3f GiB)\n",
                (long long) k, (long long) n_kv,
                (long long) k, (long long) ub, (long long) n_head,
                (long long) n_kv, (long long) ub, (long long) n_head,
                out_bytes / (double) (1ULL << 30));
    std::printf("int32 last-head element offset: %llu (%s 2^31)\n",
                (unsigned long long) ((uint64_t) (n_head - 1) * (uint64_t) ub * (uint64_t) n_kv),
                (uint64_t) (n_head - 1) * (uint64_t) ub * (uint64_t) n_kv >= (1ULL << 31)
                    ? ">=" : "<");

    ggml_init_params ip = {
        /* mem_size   */ 16 * 1024 * 1024,
        /* mem_buffer */ nullptr,
        /* no_alloc   */ true,
    };
    ggml_context * ctx = ggml_init(ip);
    if (!ctx) die("ggml_init failed");

    ggml_tensor * a = ggml_new_tensor_3d(ctx, GGML_TYPE_F16, k, n_kv, 1);
    ggml_tensor * b = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, k, ub, n_head);
    ggml_set_name(a, "k_cache");
    ggml_set_name(b, "q_absorbed");
    ggml_tensor * out = ggml_mul_mat(ctx, a, b);
    ggml_set_name(out, "kq");

    if (ggml_nbytes(a) != a_bytes || ggml_nbytes(b) != b_bytes || ggml_nbytes(out) != out_bytes) {
        die("unexpected ggml tensor byte size");
    }

    ggml_cgraph * gf = ggml_new_graph(ctx);
    ggml_build_forward_expand(gf, out);

    ggml_backend_t metal = ggml_backend_metal_init();
    if (!metal) die("Metal backend unavailable");
    ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, metal);
    if (!buf) die("Metal tensor allocation failed");

    // Clear first so an omitted write is deterministic, then upload only the small inputs.
    ggml_backend_buffer_clear(buf, 0);
    std::vector<ggml_fp16_t> host_a((size_t) k * (size_t) n_kv, ggml_fp32_to_fp16(1.0f));
    std::vector<float> host_b((size_t) k * (size_t) ub * (size_t) n_head);
    for (int64_t h = 0; h < n_head; ++h) {
        const float value = (float) (h + 1);
        const size_t begin = (size_t) h * (size_t) ub * (size_t) k;
        const size_t end = begin + (size_t) ub * (size_t) k;
        std::fill(host_b.begin() + begin, host_b.begin() + end, value);
    }
    ggml_backend_tensor_set(a, host_a.data(), 0, a_bytes);
    ggml_backend_tensor_set(b, host_b.data(), 0, b_bytes);

    const enum ggml_status st = ggml_backend_graph_compute(metal, gf);
    if (st != GGML_STATUS_SUCCESS) {
        std::fprintf(stderr, "FAIL: Metal compute returned status %d\n", (int) st);
        ggml_backend_buffer_free(buf);
        ggml_backend_free(metal);
        ggml_free(ctx);
        return 1;
    }

    int failures = 0;
    const int64_t ms[] = {0, n_kv - 1};
    const int64_t ns[] = {0, ub - 1};
    for (int64_t h = 0; h < n_head; ++h) {
        const float expected = (float) k * (float) (h + 1);
        for (int64_t n : ns) {
            for (int64_t m : ms) {
                const uint64_t elem = (uint64_t) m
                    + (uint64_t) n_kv * (uint64_t) n
                    + (uint64_t) n_kv * (uint64_t) ub * (uint64_t) h;
                float got = NAN;
                ggml_backend_tensor_get(out, &got, (size_t) elem * sizeof(float), sizeof(got));
                if (!std::isfinite(got) || std::fabs(got - expected) > 1e-3f) {
                    std::printf("FAIL head=%lld n=%lld m=%lld expected=%.1f got=%g\n",
                                (long long) h, (long long) n, (long long) m, expected, got);
                    failures++;
                }
            }
        }
    }

    std::printf("checked %lld head corners: %s\n",
                (long long) n_head * 4, failures ? "FAIL" : "OK");

    ggml_backend_buffer_free(buf);
    ggml_backend_free(metal);
    ggml_free(ctx);
    return failures ? 1 : 0;
}
Fix: three uint64 casts in mul_mm.metal
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);

@feni6

feni6 commented Sep 1, 2026

Copy link
Copy Markdown

Upstream fix PR for the Metal mul_mm int32 dst-offset wrap diagnosed above is now open: #28210 — with independent third-machine confirmation (M3 Max: probe repro, negative control, and an Apple Silicon non-regression sweep) by @eauchs in #27752 (comment). -ub 128 remains the workaround here until it lands.

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

Labels

conversion model Model specific mtmd Related to multimodal functionality (video/image/audio) testing Everything test related

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants