[gopinath87607]

Add GLM-5.3-Flash (glm5-next) support

GLM-5.3-Flash is a 312B-parameter / 17B-active hybrid MoE. Its trunk is 45 layers: 11 absorbed-NoPE-MLA
full-attention layers (every fourth) and 34 KDA linear-attention layers, all wrapped in hyper-connections. In the
quantisation this was developed against (UD-IQ4_XS, 5 shards) the dense weights stream about 7.5 GiB per token, so
it runs on one card but not quickly.

This adds the architecture end to end — geometry detection, the GGUF tensor map, five new kernel files, the
host-side layer order, expert-pool binding — plus everything the model carries that the first port did not run: the
SwiGLU clamp both references apply, the DSA k-pool indexer, the model's own MTP draft block, a VRAM tier for the
routed experts, and a chunked prefill that reads each projection once per group of tokens.

24 commits on glm5-next-support (22 of them its own, plus a merge), started from v0.1.40.3 and now merged with
today's main (v0.1.40.4-126-gfb58e0d). It merges cleanly — checked with git merge-tree, which reports a
tree and no conflict list. The merge had two conflicts, both the same shape and neither a real disagreement: each
side added a member to FileExpertSource at the same spot, so both are kept (copy_blobs from upstream, slices
from here), and the branch's zero-blob tolerance for glm5-next's dense-lead layers auto-merged intact.

What a glm5-next block has that nothing else in the engine does

Each of these is here rather than bent into an existing family, because each one produces plausible output when
wrong rather than crashing:

• Hyper-connections (mHC) own every residual. The sublayer's output is mixed back into four streams through a
  Sinkhorn-normalised 4x4 matrix computed per token from a 24-wide projection with its own base/scale. Ported from
  the reference's hc_pre/hc_post, not re-derived.
• KDA is a delta-rule recurrence with a per-head state, a depthwise causal convolution over three separate
  q/k/v streams that nevertheless share one state (the reference concatenates their filters before convolving), an
  L2 normalisation whose epsilon is a floor on the norm rather than a term inside the sum, and a decay gate
  bounded to (-5, 0).
• Absorbed NoPE MLA. K and V are the same tensor: the cache holds the 512-wide kv_lora latent, one row per
  token, and one token's value is that latent. There is one KV head, no RoPE (rope.dimension_count is 0), the
  query is absorbed into latent space, and the softmax scale is 1/sqrt(256) — the per-head width before
  absorption — not 1/sqrt(512).

The router is not qwen4exp's

probs  = sigmoid(logits)          <- no softmax anywhere in the model
select = probs + bias             <- the bias steers SELECTION only
ids    = top_k(select, k)         <- ties by ascending index
w      = probs[ids]; w /= sum(w); w *= scale

The weights are the un-biased probabilities, normalised by their sum rather than their max, and scale
(2.5) is applied after. ffn_gate_inp.weight is F32 here, where qwen4exp's is BF16 — reading it as BF16 gives 288
plausible logits and a different top-8.

Running it

The routed experts are served from a pack built by tools/iq_pack.py (they live in native_experts.txt), and the
engine reads the architecture out of that pack. Add --mtp to the pack command for the draft block (below); a pack
built without it still loads.

strata --serve \
  --pack            <pack dir>          # tools/iq_pack.py
  --native          <model>-00001-of-00005.gguf   # the dense tensors, served natively
  --mmap-experts
  --max-context     16384 \
  --prefill         4096 \
  --glm-gpu-experts 0                   # optional: the VRAM expert tier

The switches this adds

New engine flags

• --glm-gpu-experts N — keep N MiB of the routed experts in VRAM and compute those on the card. 0 means
  every free byte less STRATA_GLM_GPU_RESERVE_MIB; -1 (the default) is off. Inert on another architecture —
  the first family never reads the flag. Refused nowhere, because it has nothing to refuse.
• --dsa — attend the 2,048 cache cells the DSA k-pool indexer selects, instead of the whole cache. Refused by
  name on another architecture, which has its own indexer behind --native-qsa-indexer.
• --mtp-block — load the model's own draft block (the nextn.* tensors past the trunk) and run it. Needs a
  pack built with iq_pack.py --mtp. Refused by name on another architecture.
• --mtp-probe — measure only: read the prompt, run the head, run ONE draft step, print MTP <draft id> and
  exit. Implies --serve.
• tools/iq_pack.py --mtp (pack side) — write the draft block's rows, blk.45.* and nextn.*.

New environment variables

• STRATA_GLM_GPU_RESERVE_MIB — the VRAM the tier leaves alone when --glm-gpu-experts 0. 512 MiB by
  default (it was 2,048): sampling all four cards once a second through a whole run, every card ended with
  2,473-2,580 MiB free, of which 2,048 was this reserve and the rest the rounding of the budget into whole slots,
  and the entire chunk path — plan, activations, scratch — moved the card by 2 MiB peak to peak. The 1.5 GiB
  it hands back pays for the second half of a chunk's slots, which is what lets a wave's copies run under the
  previous wave's kernel.
• STRATA_GLM_GPU_CHECK=1 — recompute every VRAM-tier hit on the CPU, print the per-layer mean and rms, and
  fail the request when a hit differs from the CPU's by more than that row's own rms. It covers both paths:
  the decode path's hits entry by entry, and run_chunk's whole plan, sampled with a stride. =2 also prints the
  wrong-slot figure as it goes, so the bar can be re-derived rather than taken on faith.
• STRATA_GLM_GPU_CHECK_N — how many entries a layer the chunk path's check recomputes, 16 by default. It has
  to sample: a full walk of a 512-token chunk is 103 GMAC a layer, about six minutes at the pool's measured
  0.28 TMAC/s, which is the very cost the card path exists to avoid.
• STRATA_GLM_GPU_CHECK_FROM=<layer> — start the per-layer check at that layer instead of 0. The check is
  what fails a request when a slot is wrong, and it costs the CPU time it takes to recompute the sample; a run
  that has proved the layers below can skip them, and it says so once in the log.
• STRATA_GLM_GPU_SLOT_CHECK=1 — walk the slot table's own invariant (slot_[layer * n_expert + e] against
  owner_[s]) at the top of the decode's run_hits and at the end of every chunk, and fail the request naming the
  layer, the slot and both experts. A few hundred comparisons a layer a call. It is what caught the stale table
  behind "layer 6 holds the wrong expert", which no per-layer value check can see one token before it bites.
• STRATA_GLM_GPU_CHUNK_DIRECT=0 — put the staging ring back in front of the source for a chunk. Default is
  direct (the source's own bytes DMA'd into the slot, the host copying nothing); =0 is the A/B arm that measures
  what the copy costs, and it is worth 1.26×.
• STRATA_GLM_GPU_CHUNK_THREADS=<n> — threads that assemble one wave's blobs. Only meaningful on the staged
  path; the direct path has nothing to assemble.
• STRATA_GLM_GPU_PIN_GIB=<GiB> — off unless set. Reserve that much RAM as one registered (pinned)
  anonymous mapping a MoE layer, so a chunk DMAs straight out of it instead of through the driver's staging buffer
  for a pageable file mapping. The probe says pinned is 2.2× faster than pageable and blocks the caller for none
  of the transfer — and as an auto-sized default it measured a 2.4× regression (245.7 s / 21.6 tok/s against
  103.3 s / 51.4 tok/s for the same binary with the knob at 0), because every registered byte is one the kernel
  can no longer hand to the pool's own reads of the file mapping. Opt-in until that has an explanation.
• STRATA_GLM_GROUP_MAX=<n> — the token-group width a chunked prefill carves for and runs at, 8 (the
  NATIVE_MMVQ_MAX_NCOLS cap) unless set, clamped to kGlmGroupMaxNtok = 512. Past the cap a native projection
  goes through prefill::mmq::dense — one GEMM over the whole group — instead of the per-column walk.
• STRATA_GLM_NO_CLAMP — make both SwiGLU limits read 0, which is exactly "no clamp". The A/B arm that proves
  the clamp is what moves the output.
• STRATA_GLM_NO_GROUP — pin prefill's pre to one token instead of a group of eight. The A/B arm that proves
  the group is bit-exact.
• STRATA_GLM_NO_PIPE — run the chunked prompt's stages in order on the prompt's own thread, the schedule the
  pipeline replaced, so one binary can measure the two against each other. The two arms must produce the same
  tokens, and on a 5,294-token prompt they do, byte for byte; the schedule is the only difference between them.
• STRATA_GLM_GROUP=N — pick the group width yourself, 1..8, which is the arm that measures what a group
  buys: the card's per-token cost against N says whether a weight read is amortized or repeated. Measured on
  a 128-token chunk at 1 / 4 / 8: the whole model's card time 6.45 / 2.31 / 1.77 s, and stage 0's weight bytes
  294,474 / 73,618 / 36,809 MiB. The bytes are exactly ÷4 then ÷2 — what a read amortized over the group looks
  like — while the time stops falling in step once the launch bill is no longer the larger half. Nothing outside
  1..8 is accepted; STRATA_GLM_NO_GROUP=1 is the same thing pinned to one, kept because measurements already
  name it.
• STRATA_GLM_PREFILL_TIME=1 — print where a chunk's time goes: one line splitting the card, the CPU pool,
  post and the dense lead, then six sections of pre, from CUDA events. On stderr, and at no cost to a run that
  does not set it. It also prints a projection census — every top-level gemv_quantized call by path (native
  batched, canonical batched, canonical single) and the top fourteen tensors by weight bytes, with the column
  count and the columns a call carried. That is what said the bands were the bill, and it reports the same number
  the fix removed: 8,208 calls over 65,664 columns for one stage-chunk, 0 MiB of it repeated per column.
• STRATA_GLM_POOL_TIME=1 — one line a layer a chunk from the expert pool: the routed pairs, the distinct
  experts the grouping reached, the jobs the MAXT-token cut made of them, the ideal bytes against the bytes read
  (so the cut's cost is the ratio), and the pool's own milliseconds split gate/up, quantize and down.
• STRATA_MTP_DUMP_H=<path> — write the hidden state the draft block is fed, as raw f32, to diff against the
  oracle's own dump.
• STRATA_NO_SLICE_MT — put the CPU pool's multi-token row kernels back to the per-token dot. Bit-identical to
  --prefill 1 at every chunk size (verified at 1, 128 and 512), and 2.3x slower on the pool. It is also the
  switch that makes greedy output independent of --prefill; see Known gaps.
• STRATA_SLICE_MT_DOWN — opt in the down-rows half of the same kernels. About 4% here; off by default because
  the engine's own dispatch deliberately does not name IQ4_XS for it.

Flags that mean something different on this arch

• --prefill CHUNK — a chunk reads each projection's weights once per group of eight tokens, not once per
  token. auto and no flag both mean 512 (glm5-next has no expert-cache buffers to lend, which is what the
  first family's auto scans), any CHUNK is clamped to the pool's own limit of 4096 silently, and
  --prefill 1 is the one-token-at-a-time control arm.
• --trim-stage-weights / STRATA_STAGE_TRIM — forced on. Without the carve every stage loads all 7,662 MiB
  of dense projections, which does not fit an 8 GB card at all.
• --layer-split auto — its own placement search, not the first family's expert-profile search, and it prices
  the draft block as the last stage's extra layer.
• --batch / concurrent requests — one request at a time: the KDA layers are a recurrence and the MLA cache is
  one array indexed by absolute position, so there is no slot form to share.
• --spec — refused by name; there is no verify window on this arch yet (see Known gaps). A native pack
  normally requires --spec >= 2, and glm5-next does not.

Refused by name, because the flag has no meaning here

--spec, --gpu-stages (no captured layer graphs), --no-pool (the pool computes every expert the card's tier
does not hold), --expert-cache-remote and an auto expert profile (the first family's helper caches; this arch's
tier is --glm-gpu-experts), and --mtp (the first family's drafter — use --mtp-block).

Serve

reasoning_effort: "none" (which the front ends map to enable_thinking: False) now closes glm5-next's thinking
block. Its chat template opens the block whatever it is told and tests no enable_thinking at all, so thinking off
used to think at full length; the empty block is the shape that template writes for a past turn that did not think.

Measured

2 × Xeon E5-2680 v4, 2 × RTX 3060 12 GB + 2 × RTX 5060 8 GB (auto split 0-11 / 12-23 / 24-35 / 36-44),
UD-IQ4_XS, one request at a time, greedy, 16K context.

Decode, one card: 4.5 tok/s (222.6 ms/token) — the dense weights stream ~7.5 GiB per token, of which KDA's
q/k/v and output projections are 4.84 GiB (64%); the matvec kernels run at 319-325 GB/s, about 88% of the card's
peak, so this is bandwidth-saturated rather than launch-latency bound.

The VRAM expert tier (--glm-gpu-experts 0, all four cards): +43% decode, 6.7 → 9.6 tok/s, prompt
reading unchanged at 19-20 s. The tier took 19.1 GiB across the four cards — 61 / 53 / 22 / 30 slots a MoE layer,
1,719 of the model's 12,096 (layer, expert) pairs — and served 49.7% of the routed-expert work on the cards.
It is not bit-identical to the CPU pool: a hit differs from what the CPU would have produced by a mean of 1.6e-2 of
the row's rms (worst single value 2.8e-1), the same 1.5e-2 relative the engine's native_expert_parity records for
this path's activation rounding. STRATA_GLM_GPU_CHECK exists because a slot holding the wrong expert also
produces a plausible token: on 5,070 layer-checks the correct slot's worst layer was 0.28 rms while a wrong slot
ran 2.2-5.6, a 7.9x gap around the bar.

Chunked prefill, grouping eight tokens: a 344-token prompt reads in 32 s → 18 s (10.8 → 19.1 tok/s), one
run an arm, --prefill 256. The card's share of the prompt's two chunks went 17.57 → 4.97 s, a factor of 3.5,
while the CPU pool's is the same to 1% (11.95 vs 11.86 s) — the pool is now the larger half of a prefill, 11.9 s of
the 17.2 s those chunks take in all. A group of one is the arm that shows the band change is only the group: same
code path at width 1, 32 s where the pre-change binary read 31 s.

The absorbed-MLA bands were the chunk's launch bill, and are not any more. nsys on one 128-token chunk of
the same split: the two band loops of the MLA layers (step 3 and step 6 of mla_layer) issued 180,224
ncols == 1 GEMVs, each with its own native_quantize_q8_1 — 360,448 of the chunk's 434,144 kernel launches,
83.0%, at 2.18 µs and 1.21 µs a call for 136 KiB of weights read (512 rows of a 256-wide Q8_0 fold; the count is
exact, 128 tokens × 11 MLA layers × 64 heads × 2 tensors). A band is per HEAD but is the same weight rows for
every token, so the only thing between it and one batched call was the layout: qfull/kqv are token-major and
project_rows wants a batch's columns contiguous. One head-major transpose per loop (glm_heads_major,
[t][h][w] ↔ [h][t][w], its own inverse) makes them 22,528 ncols = 8 calls — 45,056 launches — a layer's
band tensor is 1,024 calls where it was 8,192 at one column a call. The census weighs both sides: one stage's
128-token chunk went from 51,216 calls to 8,208 over the same 65,664 columns, and a band tensor of one MLA
layer (1,088 MiB in 8,192 calls at one column a call) does not reach its top fourteen any more. On the 1,334-token
prompt at --prefill 512: 71.9 s → 65.9 s, single runs
each, against a 528 ms-a-chunk prediction from the nsys figures — 11 chunks, so 5.8 s, and the run gave 6.0. The
chunk curve below, measured the same way on the same prompt, reads 71.8 s → 65 s at that chunk size. The
output does not move: in multi_exact (the default) every column of a batched call is bitwise equal to the
single-column call it replaced — mmvq_multi_parity, 148,480 outputs, 0 differ, and its control finds 63,830 at
multi_exact off — and 48 greedy ids at --prefill 512 are byte-identical to the pre-change binary.

Chunk size. A chunk must read every expert its tokens touch before it can serve any of them, and 88 tokens
already touch 263 of the model's 288 experts, so a bigger chunk does not read fewer bytes per token — it pays that
fixed cost fewer times. 1,334-token prompt: 74 s at --prefill 128, 65 at 512, 62 at 2048 and 62 at 4096. 5,294-token
prompt: 325 s at 1024, 318 at 2048, 317 at 4096 — flat within 3%, so past 1024 on a long prompt a bigger
chunk buys nothing, while the session rows grow 241.3 / 482.5 / 965.1 MiB. That is why the default is 512 and not
something larger.

Layer split. A split costs prefill nothing at that chunk size: 9.7 tok/s on one card, 10.4 on four, 2,615-token
prompt at --prefill 4096.

The chunked prompt pipelines its stages. The chunk call used to be for i in stages with a D2H and a
cudaStreamSynchronize at every boundary, all on the prompt's own thread — the four cards took turns and three
were idle at any instant. A stage is a thread now, with two pinned buffers a boundary, so a stage waits on its own
boundary instead of on the whole chain. One binary, STRATA_GLM_NO_PIPE picking the schedule, 5,294-token prompt
in 512-token chunks: 179,129 ms (29.6 tok/s) pipelined against 337,021 ms (15.7 tok/s) serial — 1.881×, and the
two arms' greedy ids are byte-identical. At 15,964 tokens the same split reads in 659,775 ms (24.2 tok/s) against a
serial prediction of ~1,662 s, so the factor grows with the prompt.

The CPU pool is the floor, and that is not a scheduling problem. The serial arm's own stage totals are pool
32.5 / 44.5 / 45.5 / 33.9 s (156.4 s) and card 51.8 / 56.0 / 38.6 / 27.1 s (173.6 s) against a 336.3 s wall, which
is their sum; pipelined, the wall is 179.1 s, 22.7 s above the pool. GlmExpertPool is one object on one set
of cores, so overlapping the stages cannot divide those 156.4 s — 1.88× is most of what is reachable at this
context, and the ceiling rises with the prompt because the cards' share of the work is what grows.

The growth is the attention, and it is entirely in the card's pre. Per stage, over the 5,294-token run's ten
512-token context steps, the card's pre phase grows by 0.668 / 0.656 / 0.398 / 0.264 s a step (1.986 s in all),
while post and the dense lead stay flat. The 15,964-token run measures 0.645 / 0.638 / 0.394 / 0.261 — 1.938 s a
step, within 3.5% of the same fit over three times the range, under a different schedule. A chunk's only
context-dependent work in pre is the MLA attention over the cache, so a prompt's card time is quadratic in its
length: the 99,181-token prompt this was measured for projects to ~37,800 s of card time, 96% of it that growth.
That is what --dsa exists to cut, and it does: the same 15,964-token prompt, same binary, same split, --dsa
reads in 539.5 s (29.6 tok/s) against the dense arm's 659.0 s (24.2 tok/s) — 1.22×, and the two arms' greedy
ids are byte-identical over the 8 tokens they were asked for — a match worth exactly what 8 tokens are worth, and
the verification section says why.

The per-chunk picture is the honest one, because --dsa pays a cost the dense path does not: the indexer runs at
every chunk while the sparsity only starts paying once more than 2,048 cells are visible. Summed over the four
cards, the attn section runs 1,046.9 s dense against 713.0 s for --dsa over the prompt's 31 chunks — and the
slopes are the result, 1.764 s a chunk dense against 0.112 s, 15.8× flatter. --dsa is slower until chunk
10: at chunk 4 it is 22.6 s against 13.0, because the indexer is by then selecting over pools while the attention
is still capped at 2,051 cells and dense is only reading 2,560. By the last chunk the dense attn is 96.4% of
its card time against 91.7%, and it is that slope applied 194 more times that makes a 99,181-token prompt
project to hours.

The selection was the load-bearing part of that, and it was O(n_vis²) in a single thread. glm_dsa_select is
launched once per query token with n_pools = the pools so far, so at 16k context one thread walked every visible
pool once per ranked pool: 4,096² = 16.8M comparisons, 512 times a chunk, on all 11 MLA layers. One stage-chunk's
attn was 210.0 s of a 210.8 s card time against 21.3 s for the same chunk dense — a 10× regression whose
whole cause was one loop, and the comment above it said a radix top-k "would only pay at prefill lengths", which
is exactly what this model's prefill is. It is now an MSD radix pass for the top_pools-th largest score and a
stable radix sort of the picks, and at the prompt's last chunk the four cards' attn is 26.9 s against dense's
58.2 — the collapse in the slope, not a constant taken off it.

The MLA kernel's own shape, and the tier that now serves prefill

Both landed after the numbers above, measured the same way: 4-way split, UD-IQ4_XS, the 5,294-token prompt,
greedy, chunk 4096 unless stated.

The attention kernel was the one section that did not amortize with the group width. mla_attn_kernel ran one
128-thread block per (head, token) and reduced the 512-wide dot with a full block reduction — seven
__syncthreads() for every row of the cache. Timed on its own it cost 0.68 µs a cache row, dead linear in n_kv:
1 KiB in 0.68 µs is ~1.5 GB/s against the card's 391 GB/s measured D2D roofline. At a 4096-token chunk that is
4,096 launches of a 64-block grid a layer, which is why it grew while nothing else in pre did. The replacement
puts one warp on a (head, query) pair, 16 latent slots a lane, and reduces with __shfl_xor_sync — no barrier —
with a lazy rescale that runs O(log n_kv) times instead of once a row. Standalone, same-conditions A/B with the
T=1 control row identical in both builds: 2400.6 ms → 489.2 ms per layer per 4096-token chunk, 4.91×. In
situ, layers 0..11 at 8 tokens a call, the attention section goes 21.757 s → 12.639 s over 6,144 pre calls,
every other section unchanged to within noise. End to end, kernel-only A/B: 199.2 s / 26.6 tok/s → 170.9 s /
31.0 tok/s, 1.166×, both arms emitting the same token. It also deletes both transposes from the dense path's
steps 5 and 6, because the kernel's contract is already the layout band_out holds.

The VRAM expert tier serves prefill too, and that result is conditional. --glm-gpu-experts N now streams every
expert a chunk routes to through the stage's own slots — the only change that can take the pool off the critical
path, since 107.3 s of the four stages' 162.6 s at chunk 4096 is the pool. It wins by amortization and loses
without it:

chunk	CPU pool	VRAM tier
4096	27.1 tok/s	31.0 tok/s
512	28.8 tok/s	15.1 tok/s

A chunk streams ~1.6-2.6 GiB a layer whatever its size, so the tier's cost is per chunk: the card's own work is
0.25-0.52 s a layer and the host assembly of those blobs is 0.35-1.17 s at 2.3-5.3 GB/s. At 512 tokens the same
copy feeds eight times fewer. The tier cannot win on memory traffic — both paths read each blob exactly once —
so it can only win on compute, and today the copy is larger than the compute it feeds.

And it had no correctness gate, though it claimed to have one. STRATA_GLM_GPU_CHECK covered the decode path's
hits and nothing covered run_chunk, though glm_gpu_experts.hpp said it was "the gate, and it is the same one".
A chunk is 12,096 slot decisions a layer instead of 8, all of them silent when wrong, because a slot table that is
right about indices and wrong about bytes still produces a plausible token. It now walks the plan with a stride and
recomputes each sampled entry on the CPU — a full walk is 103 GMAC a layer, about six minutes at the pool's
0.28 TMAC/s, which is the very cost the card path exists to avoid — with STRATA_GLM_GPU_CHECK_N setting the
sample. On all 42 MoE layers, 462 layer-checks: max/rms median 1.09e-1, worst 2.94e-1, against the 1.0 bar.
And the =2 wrong-slot self-test says the bar separates on this path's own data rather than only on the decode
path's: 5.19 to 6.35 wrong against a worst correct entry of 2.94e-1, a 17.6× to 294.6× gap, median 51×.

What feeding a chunk costs, and the three ways the slot table was wrong

All of it on the same rig, the 5,304-token prompt, chunk 4096, 4-way split, UD-IQ4_XS, greedy, one run an arm.

The feed: the source's own bytes, not a blob this process assembled. The chunk path began with a staging ring
— the source's blob assembled in host RAM, copied into the ring, then DMA'd to the slot — which makes the host the
feed and the DMA take turns with the tensor work. The source can hand an expert's own bytes over instead: a
GGUF-in-place pack gives three ranges (gate, up, down), a source whose experts live in registered memory gives
one. Nothing is assembled, so no ring is allocated at all, and the host copies nothing.

arm	feed	source	per-layer check	wall	prefill
af-pin0	staged ring	file mapping	on	103.3 s	51.4 tok/s
af-stg2	staged ring	resident arena	off	101.8 s	52.1 tok/s
af-dir2	direct slices	file mapping	off	80.7 s	65.7 tok/s
af-dir3	direct slices	file mapping	on	280.3 s	18.9 tok/s
af-arbd1	direct slices	resident arena	on, per-wave instrument on	85.3 s	62.2 tok/s

1.26× over the staged ring, and that is the pair the change is for: same chunk, same split, same prompt, and the
per-layer check off on both. In the staged path the host assembles the blob in RAM whichever source it came
from, so af-pin0's 103.3 s against af-stg2's 101.8 s is the source not mattering — the size of the copy in front
of the DMA does.

And the per-layer check is an instrument with a price at chunk scale, which is not a detail. af-dir3 is
af-dir2's configuration with STRATA_GLM_GPU_CHECK=1 — the per-layer check, nothing else. It reads 280.3 s
against 80.7 s, a 3.5× slower run, with the card's own time unchanged (a 4,096-token chunk over layers 0..11 is
7.03 s of card in the slow arm against 7.83 s of card in the fast one). The check calls the CPU pool once per sampled entry — 16 entries a layer — and
each of those calls reads a whole expert, 11.67 MB, from the source: 187 MB a layer, 7.8 GB over a chunk, and the
host reads them one at a time. The disk was measured at 126 MB/s while that arm ran, and its own pace was
2.4 s a layer (104 layer-checks in 250 s), which is that rate against that read. When the source is the file
mapping and the DMA is streaming the same file through the page cache, those reads are not free; when the source
is registered memory the same check is cheap — af-arbd1's 85.3 s is a checked run at a per-wave instrument that
costs 27% on its own. Timing arms are =0; the check is the correctness arm and it is slow by design.

And the per-wave instrument is not free. The same direct arm with STRATA_GLM_GPU_CHUNK_TIME=1 reads
102.3 s / 51.8 tok/s, because the instrument's per-wave cudaStreamSynchronize is exactly the overlap the change
exists to create. Every T=1 figure in this PR is a floor rather than a schedule; the honest ones are T=0.

The wave: two streams, halves, and an event. A wave's copies run on their own stream, released to the
caller's by one event, and a wave's slots are split in two halves with consecutive waves taking turns between
them — so the only edge left is against the wave that used the same half two waves back. Measured standalone on a
5060 (bw/wave.cu, 283 blobs of 11.66 MB into 13 slots): the expert tier's per-layer wall 1.19 s → 0.55 s
with the instrument on and → 0.36 s without it. Same bytes, same kernels; only when the copies happen.

The plan needed to be a ring, and one copy of it was a correctness bug. The per-wave plan lives in pinned
host buffers, so the copy that uploads it is a real DMA and the host does not wait for it: a single copy is
refilled for wave n+1 while wave n's upload is still in flight, and the card then computes wave n against wave
n+1's rows and slots. Measured: STRATA_GLM_GPU_CHECK=1 on a 512-token chunk read max/rms 10.66 on layer 3
with the per-wave instrument off — the bar is 1.0 and a correct slot is ~0.2 — and 0.10-0.19 on every layer with
it on, whose per-wave sync was the only thing draining the upload. It is kPlanRing = 8 copies deep now, gated by
an event a copy; the host runs two or three waves ahead, not eight.

Two more ways the table went wrong, both found by a check rather than by a bad token. admit's LFU scan read
cnt[owner_[v]] for a slot whose owner was −1 — i.e. cnt[-1] — and could then record a freed slot against a
slot_ entry one before the layer's own array; it now takes a vacant slot first and skips negative owners.
And run_chunk filled slots without ever advancing next_[layer], so the next decode blind-took a slot the chunk
had filled, without clearing the displaced expert's slot_ entry. Both are the same symptom from different
directions: a slot table right about indices and wrong about bytes, which produces a plausible token.

The gate had a bar to re-derive, not to inherit. kChunkCheckBar is 0.5 against the decode path's 1.0,
because the chunk gate is the worst entry's mean over its own row's rms and on this path the worst column of a
right entry reads 3.08 against a wrong-slot figure of 6.93 — a 2.3× separation with the old bar sitting inside it.
A poison pass catches rows the card never wrote at all (the buffer is filled with a value no expert can produce;
a row still holding it after the layer is a wrong slot seen from the outside, whole-buffer, 67 MB at T=512), the
check records which wave carried each entry so a failure reads "the last entry of wave 19, which ends at plan
index 812" instead of "entry 812 of 4,096", and the card's own q8_1 activations are kept for the one case a row
cannot be reproduced from the pool.

One card cannot run this shape at all, and that is worth knowing before the tier is judged. On a single RTX
3060 the model does not fit at a 16K context — the native head's own upload failed with 428 MiB of 11,911 MiB
free — and at 6,144 tokens it loads with under 512 MiB to spare, which is no VRAM tier at all: the engine
refuses to make one and every routed expert falls back to the pool. In that shape a 512-token chunk is
15.0 s of CPU pool against 5.6 s of card (20.85 s in all): the card is busy 27% of the chunk. Prefill reads
19.3 tok/s against 65.7 on four cards, decode 4.6 tok/s. The tier is not what makes prefill slow there — the
pool is, and one card has no room to take any of it.

Verification

The ladder. /home/gopi/glm-ladder holds truncated rungs of the same GGUF (4 layers is the smallest that
runs) and a tagged-dump comparator against the ik_llama.cpp oracle. 30 of 32 tensor checks matched on the shipped
model before this branch's later work, and the arms below are the ones that changed output.

The SwiGLU clamp, which the port asserted neither reference applies — both do: swiglu_limit() has
LLM_ARCH_GLM5NEXT in its allow-list in ik_llama.cpp, and mainline applies it too. Four arms on the l4 rung
against the oracle with the same metadata, limits set to 1 so that the clamp binds:

• limits 10 / 10 (the shipped artifact): nothing binds, so the arms are byte-identical — ids and --dump-logits.
• limits 1 / 10 (the routed experts clamp, the dense lead does not): oracle rael有关snowflakeyesterdayoringredient,
  ours 8/8.
• limits 10 / 1 (the dense lead and the shared expert clamp, the routed experts do not): oracle
  raelites whooping少了少见的到的, ours 8/8.
• limits 1 / 1: oracle rael also knowns'-'s, ours 8/8.
• limits 1 / 1 with STRATA_GLM_NO_CLAMP=1: back to the unclamped string, ours 8/8.

The first one is why the clamp's total absence survived the original ladder match: at the model's own limit of 10
nothing binds. The last four are the evidence — each limit's own site matches the oracle when the clamp is made to
bite.

The DSA indexer, on l4 with indexer.top_k patched 2048 → 16 in a copy of the rung so that the indexer binds
from position ~16 (chunk 0, rows 256-510, -c 512, --save-all-logits), mean KL / argmax / top-10:

• oracle dense vs ours dense: 0.00032, 249 of 255, 9.82
• oracle --dsa vs ours --dsa: 0.00081, 248 of 255, 9.78
• oracle dense vs ours --dsa, and the reverse: 0.02583 and 0.02442 — the two crossed pairings, ~30x worse

At the unpatched top_k 2048 with -c 512 the indexer selects every visible pool, so its cell set is the dense
one. In ik that makes the flag a no-op outright — its --dsa against its own no---dsa compares equal over
the whole file at -b 512 and again at -b 64, PPL 91142.1138 — because the two paths there are the same
attention kernel over a different cell list, which with all pools selected in cell order is the same list.

In this engine it is not a no-op, and the reason is in the kernels rather than in the selection. Ours are two
implementations: glm_mla_attn is an online-softmax kernel (a tree-reduced dot per cell, a running max and a
rescaled accumulator in registers), while glm_dsa_attn is a two-pass one (one thread summing a cell's whole
latent serially, then a max and a normalise pass, matching ggml_soft_max). Same cells, different sums, so the
outputs agree to rounding and not bit for bit. Measured on the shipped 4-way split at 344 tokens, greedy:
--dsa reads A mixture-of-experts (MoE where dense reads # Mixture-of-Experts Routing in — a near-tie
between two plausible openings, flipped by accumulation order. That is the honest version of "the indexer is a
no-op below the budget": the selection is, and the engine is not, and one engine run at a diverging length says
nothing about which of the two is right. The oracle pairing above is what says that.

The selection rewrite was wrong on the branch the harness had no case for, and the harness is why. What
glm_parity could not see is which path it was testing: the DsaFix fixture shows 4 visible pools against 2
selections and the wide case 300 against 32, so the "fewer visible pools than selections" path had no case at
all — and that is the path every query below 2,048 cells takes, which is every decode. The rewrite's first
version wrote the cells in pool order there, reading "all of them are selected" as "so the order does not
matter". It does: the header's descending-score order is a separate promise, and the first attempt broke it.

Two things came out of chasing it, and the second is the more useful one.

The bug was found by reading, not by the run that prompted the reading. Fixing the order changed the 344-token
--dsa output not at all — the divergence from dense there is the two attention kernels, above — so that run
neither convicted nor cleared the pool-order version. What convicts it is the contract: the rank formula it
replaced ranked every pool below the budget too (rank < top_pools is true for all of them), so ascending pool
order was wrong regardless of whether 8 greedy tokens could see it.

And the missing coverage is now a sweep rather than a case: every (pool count, selection width) pair
straddling the budget on both sides, against score vectors full of equal values, at a pos that puts n_vis
under, at and over the width, tail on and off — 2,784 fixtures, 1,376 of them below the budget and 1,424 with a
shared boundary value, all exact against the transcription of the rank formula. Plus the one deliberately
written-out case, whose expected row's first twelve cells are spelled out rather than only compared, so a
pool-order emission fails on the first one. The full suite is 208 cases.

The MTP draft block: its drafted token matches ik_llama.cpp on all four oracle prompts. The trap it had to
avoid: is_qsa_layer(g, 45) is 45 % 4 == 3 → false, so the existing predicate makes blk.45 a KDA layer, and it
has MLA tensors and no ssm_* at all.

The group prefill is bit-exact at a fixed chunk size: 48 greedy ids byte-identical with
STRATA_GLM_NO_GROUP=1 on and off, and identical across a repeat of the same arm, so the engine is deterministic
and the comparison has power. A later revision of this section also claimed the ids were identical at chunk 1, 128
and 4096; that is wrong, and the counterexample is below.

The chunk size moves a token, and the CPU expert kernels are what moves it. One 344-token prompt, greedy, on
the shipped 4-way split — where 512 and 4096 are the same arm, because 344 tokens fit in one chunk either way:

--prefill	first tokens
1	In a mixture-of-experts (MoE) transformer, each MoE layer re
128	A mixture-of-experts (MoE) layer replaces the single FFN blo
256	A mixture-of-experts (MoE) layer replaces the single FFN blo — diverges from 128 later in the same 48 tokens
512	# Mixture-of-Experts Routing in Transformer Layers

The 128 and 512 arms were each run twice and reproduce themselves byte for byte, so this is a difference in the
engine and not sampling. It is not the bands and not the cards: STRATA_NO_SLICE_MT=1, which puts the routed
experts back on ggml-cpu's per-token dot, makes chunks 1, 128 and 512 byte-identical — so the cause is
measured rather than inferred. The multi-token
CPU expert kernels (native_gu_rows_slice, and gu_rows_nt/dot_rows_nt under it) take their shape from the
number of tokens sharing a job, and a job's composition is a function of how many tokens in the chunk routed to
that expert. They round ~3e-8 away from the per-token dot; glm5-next's q8_0 quantization of h between the two
projections occasionally turns that into ~1% on a row, up to 1.76 on a logit. The switch buys 2.3x on the pool and
on the short prompt it was first measured on the text was unchanged, so it stays on by default; making the output
chunk-invariant costs that 2.3x, and the default 512 is the chunk the ladder above was validated in.

Tests. glm_parity --selftest is wired into ctest: 208 cases, 0 failures. Sixteen of the new ones are the
head/token transpose — test_heads_major at four shapes ((3,2,8), (2,5,4), (4,4,3) and the model's own
(64,8,256)), each checked value by value rather than by multiset, with the round trip and two rivals that are
themselves permutations; both counts are above 1 in every case, because at ntok == 1 the move IS the identity
and the rival would be the right answer. Every case is a double-precision
transcription of the reference plus a rival reading that must disagree, because each of these choices (the 2
on post, the end of the Sinkhorn loop, the L2 floor, the sign nesting in the decay gate, one softmax over the
whole DSA pool block instead of per key-dim, the scale from kv_lora instead of qk_nope, the tail appended or
not, ties to the higher index) yields well-formed output when wrong. native_expert_parity gains an arm that
scales the activation until the SwiGLU fixture reaches all three sides of the clamp, asserts that it did, and
checks the kernel against a double transcription at the model's limit and at 1.0.

ctest -R "glm|native_expert_parity|expert_layout|file_expert_source" — 14/14 pass.

The whole suite on the merged tree: 102 cases, 99 pass, 3 skipped, 3 fail. The three failures are not from
this branch, and this time that is measured against upstream rather than argued: a separate worktree was built at
origin/main (fb58e0d, the exact commit this branch was merged with) from the same CMake configuration, and

• qsa_select_bench (warm: unspecified launch failure) and pdl_parity (pdl_supported() = 0 on this
  sm_120 RTX 5060, then CUDA: edges: invalid argument) fail there with byte-identical messages. Both were
  touched by upstream's own commits since this branch forked (dfa26c1 tightened the qsa accuracy floor to 1e-5;
  c788e2a reworked the PDL prefetch below sm_70), and neither is a file this branch touches.
• ple_parity wants a Qwen Q2_0 fixture that is not on this box: required PLE table is missing or incompatible: ../../Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf.

Also worth knowing before reading those two: the build directory here is configured
CMAKE_CUDA_ARCHITECTURES=75 while the cards are sm_86 and sm_120, so the GPU-level failures may be as much a
property of this build as of the tests. No GLM target is among the failures.

Not in this PR

• Vision. Text generation only.
• --batch and concurrent requests. The KDA layers are a recurrence and the MLA cache is one linear array
  indexed by absolute position, so there is no slot form to share across requests.

Known gaps

• There is no verify window yet, so the MTP block is inert in a real run. It drafts tokens and nothing checks
  them, which is why --spec is refused rather than defaulted. --mtp-block is a measurement switch;
  --mtp-probe is what turns it into a number. The first family's window does not transfer — it is captured layer
  pairs, expert_pool_dispatch_multi and a state commit with its own layout.
• The MTP step refuses a multi-token scratch rather than half-reading one (glm_mtp_step returns an error if
  b.ntok != 1): the block is a single-token step and a chunked prefill would otherwise hand it a group.
• Greedy output depends on --prefill, on a long enough prompt. --prefill 1, 128, 256 and 512 each give a
  different completion of the same 344-token prompt, deterministically, because the CPU expert kernels' shape
  follows the chunk (see "The chunk size moves a token" in Verification). STRATA_NO_SLICE_MT=1 makes the output
  chunk-invariant at 2.3x on the pool; the default 512 is the chunk the ladder was validated in, and the card
  side is not implicated. Closing it properly means making the multi-token expert kernel bit-identical to
  ggml-cpu's per-token dot, which is the 2.3x, so it is a trade rather than a bug to fix.
• The CPU pool is ~60% of a prefill chunk now, and the tier that takes it over only pays at a large chunk.
  Measured with the staging ring, --glm-gpu-experts N moves that share to the cards at --prefill 4096
  (27.1 → 31.0 tok/s) and loses 2× at --prefill 512 (28.8 → 15.1), because the tier's cost is per chunk rather
  than per token — the direct feed above is what turns the 4096 figure into 65.7 tok/s, and the 512 arm has not
  been re-measured since. The pool is not RAM-bandwidth-bound and not NUMA-bound (confining it
  to one node halved the cores and doubled its time); it is bound by per-token row decode. This machine is a
  Broadwell Xeon with AVX-2 only, so this split is pessimistic for the pool — on a Zen 4 or Ice Lake box the card
  would be the bigger share.
• The DSA indexer's select_tail is a disagreement between the two oracles, documented in glm_dsa.hpp rather
  than coded as a fact: ik_llama.cpp appends the incomplete tail whenever kpool > 1 and reads no key at all,
  upstream reads attention.indexer.kpool_select_tail and defaults it to false, and this model's metadata carries
  neither key. We follow ik, the oracle the ladder is measured against.
• The RAM arena is off by default because as a default it lost 2.4×. Pinning a MoE layer's bytes into a
  registered anonymous mapping makes each DMA 2.2× faster on the probe and blocks the caller for none of it, but
  it takes the bytes out of the page cache that the pool's own reads of the file mapping depend on; measured
  245.7 s / 21.6 tok/s against 103.3 s / 51.4 tok/s for the same binary with the knob at 0. STRATA_GLM_GPU_PIN_GIB
  is the opt-in and the regression is written down where the knob is. It is not needed for the win: the direct
  feed's 65.7 tok/s is a file mapping's slices.
• The direct feed needs a source that can hand over an expert's own bytes. A GGUF-in-place pack gives three
  ranges and a source whose experts live in registered memory gives one; anything else falls back to the staging
  ring, which is the 51.4-52.1 tok/s arms rather than a failure. Checked against copy_blob for every layer at
  init before it is armed.
• A 12 GB card gets no tier at all on this model. At a 16K context the native head's own upload fails (428 MiB
  of 11,911 MiB free) and at 6,144 tokens there is under 512 MiB to spare, so --glm-gpu-experts refuses and
  every routed expert falls back to the pool — 19.3 tok/s prefill on one 3060 against 65.7 on four.
• geometry_key has no arch slot (conversation_state.cpp). It is an 18-element on-disk array written
  positionally with a hard-coded 18 * 8, so widening it is a format break that wants a version gate; as it stands
  a GLM conversation checkpoint is not distinguishable from a Qwen one by geometry alone.

🤖 Generated with Claude Code

[jerem91150]

Tried this branch on an AMD card on Windows, since nobody here has one.

RX 7900 XTX (gfx1100), Ryzen 9 7900X3D, 64 GB, Windows 11 24H2, driver 32.0.31041.1004. Built with tools\hip\build_windows.bat tests and STRATA_HIP_ARCHS=gfx1100 (TheRock 10.2.0a20260930). No source changes needed, just the usual unused hipError_t warnings, nothing in the glm files.

glm_parity --selftest passes 115/115 on the 7900 XTX.

One thing that tripped me up: the Ryzen iGPU shows up as HIP device 0 and the 7900 XTX as device 1. Without HIP_VISIBLE_DEVICES=1 the test runs on the iGPU, every glm kernel fails with device kernel image is invalid and you get 50 failures out of 115. Looks like broken math, isn't. The engine is fine since setup passes "gpu": 1. Might be worth having the tests pick the first device that matches the build arch, or at least print which one they use.

Running UD-IQ2_XXS end to end now (24 GB VRAM + 64 GB RAM), next to llama.cpp Vulkan with the same prompts. Will post numbers here when it's done. Notes: https://github.com/jerem91150/Strata/blob/glm53-hip/docs/GLM53_AMD_WINDOWS.md

[jerem91150]

Numbers from the 7900 XTX as promised. UD-IQ2_XXS, 24 GB VRAM + 64 GB RAM, so --mmap-experts (92 GiB of experts) and --prefill 1 (the chunked read wants sliceable experts, and IQ2_XXS gate/up aren't).

Decode, 256 tokens:

• llama.cpp Vulkan on the same file, 5 layers of experts on the card: 4.2 to 4.7 tok/s
• this PR as is, experts on the CPU: 2.5 tok/s

To get past that, I tried the first family's hit/miss split for glm5-next on my fork. It keeps a per-layer set of experts in VRAM, computes them with native_expert_grouped, and only sends the misses to the CPU pool, with LFU deciding who stays. With 13.8 GiB of slots it warms up to about 71% hits and 4.8 to 5.9 tok/s. I checked every hit against the CPU pool, and the gap is ~2% of the largest value, which is just q8_1 vs q8_K rounding.

Two smaller things came up on the way, possibly useful here:

• Unsloth's UD-IQ2_XXS has IQ3_XXS downs, which STRATA_D_FMTS doesn't list, so the engine refuses it at startup. Adding the IQ2/IQ3 types there works, since launch_down is generic.
• GLM's template always opens <think> and maps any effort it doesn't know to max. So reasoning_effort: "none" gives full-length thinking. I close the block empty in encode_prompt when thinking is off.

Code is on https://github.com/jerem91150/Strata/tree/glm53-hip (last three commits), notes in docs/GLM53_AMD_WINDOWS.md. Happy to open a PR against this branch if any of it is wanted.

[xyzzing]

Would be great if we can have a MOE focused inference engine to run all forms of MOE

[gopinath87607]

hey guys i have thrown some works share your feed backs and correction to this pr. i think i am still unable to finis this pr so i am stopping here for while i will work on this later after some time.

[jerem91150]

Follow-up to my numbers from the 7 October, for whoever picks this up again: same machine (native Windows 11, RX 7900 XTX 24 GB, Ryzen 9 7900X3D, 64 GB RAM, Unsloth UD-IQ2_XXS), now with Project Maya next to it, since it grew out of Strata and runs the same model.

To be clear about where this comes from: it is one piece of work done across Strata and Maya together. I run both on this box, and things moved both ways. The VRAM tier I proposed here (1f0cc02) and Strata's Windows HIP build script went into my Maya build; what I learned on Maya (below) is what I bring back here.

GLM-5.3-Flash, one card	this PR (my build of 7 Oct, --prefill 1, --mmap-experts, VRAM tier)	Maya v1.0.16
prompt read	token at a time (the chunked read wanted sliceable experts)	27K tokens in 46 to 59 s
decode	up to 5.9 tok/s	5.5 tok/s median (3.5 to 8)

Your chunk path from this morning should close most of the prompt gap; the 19.3 tok/s single-card figure in DETAILS.md is the number to beat. Decode is already on par.

Two things from the Maya side that may save time here:

• Numbers in long prompts. Maya split GLM prompts with qwen35's one-digit rule until v1.0.15, and from ~12K tokens on numbers came back wrong (240 read as 2400). This PR already reads the pack's pre (glm4, digits by three), so it is fine; worth a test with a long prompt full of numbers all the same.
• Disk reads on Windows. Each read seems to carry a fixed cost there: Maya's prompt reader went from 1.66 to 4.27 GB/s with 16 reads in flight and one chunk per expert part instead of 8.

When this moves again I can run it on the 7900 XTX under Windows.

[gopinath87607]

hey guys since https://github.com/mw00/project-maya has better prefill version and bit more support i would like to focus on some other model so thanks for all the test and report please keep on supporting it i am closing this pr for now.