Several buffer sizes and pointer offsets in examples/perplexity compute counts as
n_ctx / n_token / n_chunk / index times nv (or n_vocab) using int operands. At 16k
context with a >=131k-vocab model, n_ctx*nv exceeds INT_MAX (16384*131076 = 2^31+),
overflowing to a negative int that becomes a huge size_t in resize/ctor -> std::length_error
crash. Affects:
- --kl-divergence-base / --kl-divergence: log_probs.resize, the token/logits writes,
the compare-path base pointer (segfault at 32k) and inner offsets, and the read reserve
- hellaswag_score / winogrande_score / multiple_choice_score: batch_logits(n_vocab*n_ctx)
Cast the counts/indices to size_t at each site. Byte-identical below the overflow threshold;
the offset accesses that were already size_t (i_logits, eval_pairs.first) are unchanged.
Repro: llama-perplexity -c 16384 --kl-divergence-base out.dat on any >=131k-vocab model
(the hellaswag/winogrande/multiple_choice allocs overflow the same way at large -c).
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
For IMROPE models (Qwen3.5/QWEN35MOE) the RoPE op reads 4 position sections per
token, but the legacy null-pos decode fallback (llama_batch_get_one path, used by
llama-perplexity/llama-cli/llama-eval-callback) sized the position vector to n_tokens.
llama_set_inputs then copies n_tokens*4 out of that n_tokens-sized vector, an
out-of-bounds read that feeds garbage into 3 of the 4 rope sections, giving wrong and
run-to-run non-deterministic RoPE. Build the sections like the explicit-pos path (text
t,t,t,0). Standard-rope (NEOX) unchanged; VL image input uses the explicit-pos builder.
Qwen3.5-4B self-vs-self same-top: 91.3% -> 100.0% (CUDA), single-thread CPU likewise.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The Metal backend lacked GGML_ROPE_TYPE_MROPE/IMROPE support, so models
whose GGUF carries rope_sections (e.g. Qwen 3.5 hybrids) could not run
fully offloaded on Apple Silicon.
Add rope_multi_f32/f16 kernels with section-based position handling
(t/h/w/e blocks, 4 position ids per token), imrope's interleaved section
selection, and the corresponding dispatch in ggml-metal.m. Vision-mode
mrope is not implemented and is asserted out explicitly.
Validated on M2: kernel output matches the CPU backend, and Qwen 3.5-9B
(Q4_K_M, -ngl 99) WikiText-2 perplexity over 145 chunks lands within
0.006 of the same model's CPU baseline.
Faster fp32-class replacement for the tile_f32 flash-attention inner path on P100
(sm_60). Bit-identical fp32 arithmetic to the un-retiled kernel (same-top 99.28% =
the float-reorder floor; QK differs only in accumulation order, ~1 ulp; P.V
bit-identical), restructured via half2 K/V shared-memory staging that both cuts
shared memory and leans the inner loop.
Measured +4-9% vs the un-retiled fp32 tile kernel, back-to-back. The speedup is a
COMPOUND of two effects the staging produces together (shares not isolated):
(1) occupancy: smem 36992 -> 28800 B/block admits 2 blocks/SM where the un-retiled
kernel fits only 1 (2*28800=57600<=65536; 80 regs would allow 3, smem is the
ceiling); a genuine 1->2 gain, corroborated by a cross-family perf panel;
(2) a leaner inner loop: ~2x fewer QK-loop smem loads, the dropped score round-trip,
one fewer barrier.
__launch_bounds__(...,2) only PINS the target the smem reduction already reaches (a
(...,1) bound compiles bit-identically), so the directive is not itself a lever, and a
3rd block does not help (cost-curve + an implemented regs-80->124 attempt confirm smem
caps occupancy at 2).
Validated on P100 as the drop-in for the carve-out path (pr-p100-fp16). Reviewed by a
6-model Claude council + a cross-family cloud perf panel; the panel corrected an earlier
'not occupancy' framing to this compound one.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
On P100 (GP100, sm_60) the fp16 vec kernel used for decode (batch<=8)
accumulates the online-softmax denominator and the P*V product in fp16,
flipping ~3-4% of decode top-1 tokens vs an all-fp32 reference
(llama.cpp#25593). Decode is memory-bandwidth-bound on P100, so routing
sm_60 decode to the in-tree vec_f32 kernel is free (tg128 ~identical).
Gated on cc == CC_PASCAL && Q->ne[1] <= 8 (decode only) inside the
!fp16_mma_available block, so the prefill tile_f16 path, the D=256 prefill
vec path, and fast_fp16_available() are untouched, and the
is_pascal_mla_absorbed_decode early-return (MLA) is unaffected.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ggml_graph_node_has_matching_properties exempted every GGML_OP_CPY node from the source-address check, for both operands. The exemption is needed for KV-cache write copies, whose destination advances each step through the indirect-destination path. It also skips a CPY whose read source (src[0]) moves between graph replays while the surrounding subgraph stays shape-stable, so the captured kernel replays against a stale source pointer and reads the previous step's bytes.
Keep the exemption for the destination operand (src[1]) only, and compare a CPY's read source (src[0]) like any other node. Stable-source copies are unaffected and force no additional re-captures; a CPY whose read source genuinely moves now re-captures instead of reading stale memory.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
llama-dflash.cpp calls ggml_backend_is_metal() and
ggml_backend_metal_set_n_cb() inside an #ifdef GGML_USE_METAL block but
never includes ggml-metal.h, so any Metal-enabled build fails to
compile. Add the same guarded include llama.cpp already uses.
* perplexity: add signal-driven hot-swap mode for persistent KLD/PPL benchmarking
llama-perplexity can now stay resident and be driven through control/status files (reload/compute/exit): it reloads only the tensors that changed on disk and recomputes PPL/KLD without ever reloading the full model. File-based signalling works on Windows, macOS and Linux. The reload returning-to-original path now refreshes tensor data from disk instead of
reattaching stale weights.
* Not Cygwin specific
Loading a deepseek2 model with -sm graph across 2 or more GPUs segfaults
during model load, in the per-head split of the MLA attn_k_b/attn_v_b tensors.
In the ik-converted GGUFs these ship 2D with the heads folded into ne[1]
(DSV2-Lite attn_k_b is [128, 8192]). distribute_mla_tensors_for_split_mode_graph
splits them with prepare_split_tensors(2, ...), which takes the head count
from ne[2]. With the heads still in ne[1] each split comes out n_head times
too large and the upload reads past the end of the host buffer. It is
host-side (compute-sanitizer is clean) and happens at any device count >= 2.
Fix: reshape the 2D layout to 3D [head_dim, inner, n_head] before the split,
the same head-outermost interpretation ggml_reshape_3d already uses on the
non-split path. Metadata-only, guarded on ne[2]==1 so it is a no-op for GGUFs
that already ship these tensors 3D.
Tested on DeepSeek-V2-Lite-Chat Q4_K_M (only what fits across the GPUs here):
-sm graph PPL matches -sm layer and CPU within stderr (7.26 / 7.24 / 7.24,
wikitext-2 n_ctx 2048), coherent generation at 2, 3 and 4 GPU splits, and a
REAP-pruned DSV2-Lite in parity too. GLM-DSA, Mistral4 and larger deepseek
GGUFs share the split path so I expect the same fix to cover them, untested here.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* WIP: indexer_topk on CUDA
* Forgot these
* WIP
* WIP
* This seems to work
* Minor
* Fix bug. Fix suggested by @sayap using GLM-5.2
* GLM-DSA: much better PP long context performance (CUDA)
* DSA: Better way to build the attention mask
* llama-swap latest version
- Sample config still works
* llama-swap latest version ik_llama-cuda.Containerfile
- Sample config still works
* Update README.md
- The prebuilt docker images are fully working, move their instructions to the front
This, to avoid an intermediary loading of the split tensors in RAM.
Commit c32c3819f7 (PR 2102) introduced mmap for split-mode graph tensor
loading to reduce memory usage via MADV_DONTNEED. However, on Windows
dontneed_fragment() is a no-op (no madvise equivalent), so the entire
mmap'd file stays resident in RAM until load_all_data() returns.
This causes excessive RAM usage with larger models than the available RAM on Windows.
This PR fixes it by guarding the mmap path with !defined(_WIN32) and using the original read-buffer approach on Windows, where per-tensor memory is bounded to the largest single tensor size.
* WIP: indexer_topk on CUDA
* Forgot these
* WIP
* WIP
* This seems to work
* Minor
* Fix bug. Fix suggested by @sayap using GLM-5.2
* GLM-DSA: much better PP long context performance (CUDA)
* ggml: add fused sinkhorn op (eps + output-layout params); use it for openPangu mHC
* openpangu: call ggml_sinkhorn directly from mhc_post
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
Route DSA indexer selection through ggml_indexer_topk when -fidx is set: one op computes the weighted-relu head sum plus causal mask and returns top-k rows without materializing the [n_kv, n_ihead, T] score tensor. Off by default; the unfused chunked/full paths are unchanged. Composes with gathered DSA, deferred attention-chunk masks, and the set_rows mask path.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* IQK AVX2: Replace MM256_SET_M128I(x, x) identity broadcasts with _mm256_broadcastsi128_si256
## Summary
Replace all 132 instances of the identity-broadcast pattern
`MM256_SET_M128I(x, x)` with `_mm256_broadcastsi128_si256(x)` across
8 files in ggml/src/iqk/.
## Correctness
The transformation is bit-exact. The `MM256_SET_M128I(a, b)` macro
expands to:
_mm256_insertf128_si256(_mm256_castsi128_si256(b), a, 1)
which places `a` in the high 128-bit lane and `b` in the low lane.
When `a == b == x`, the result is x replicated to both lanes:
[x_lo: x_hi] = {x, x}
`_mm256_broadcastsi128_si256(x)` produces the identical register state:
it copies the 128-bit input to both lanes in a single micro-op.
Perplexity was verified identical on llama-3.2-1b-Q8_0 and
google-gemma-3-4b-Q4_0-IQ4_XS before and after the change.
The transformation is a pure intrinsic substitution with
zero semantic difference.
## Performance
### Micro-architecture analysis
The old pattern compiles to:
vinsertf128 ymm, ymm, xmm, 1 -- 3 uops, port 5, 3-cycle latency
The new pattern compiles to:
vbroadcasti128 ymm, xmm -- 1 uop, port 5, 1-cycle latency
Both instructions execute on port 5 (Intel), but vbroadcasti128:
- Uses 1/3 the dispatch slots (fewer pipeline stalls)
- Has 3x better latency (1 vs 3 cycles)
- Is not lane-crossing (no bypass delay between 128-bit halves)
### Measured results
#### Test 1: llama-3.2-1b-Q8_0, 2048 ctx, -b 128 -ub 128
Compiler | Metric | Before | After | Change
--------------|--------|---------|---------|-------
MSVC 19.44 | PP t/s | 596.89 | 604.57 | +1.29% (noise imo)
MSVC 19.44 | TG t/s | 55.13 | 55.72 | +1.07% (noise)
Clang 19.1.5 | PP t/s | 645.72 | 700.93 | +8.55% (systematic gain)
Clang 19.1.5 | TG t/s | 54.26 | 54.06 | -0.37% (noise)
#### Test 2: google-gemma-3-4b-Q4_0-IQ4_XS, 4096 ctx, -b 512
Compiler | Metric | Before | After | Change
--------------|----------|----------|----------|-------
Clang 19.1.5 | PP t/s | 262.40 | 314.79 | +19.96% (systematic gain)
Clang 19.1.5 | TG t/s | 30.78 | 32.32 | +5.00% (noise, TG oscilates between 31.5 and 33 t/s before / after)
Clang 19.1.5 | Total ms | 48881 | 44696 | -8.56%
Both tests ran on Intel Core Ultra 265K, 18 threads, flash_attn=1
The IQ4_XS result shows a dramatic PP improvement (+20%) because this
quantization format uses significantly more identity broadcasts in its
dequantization path (lookup-table expansion, scale duplication). The
compressed 4-bit representation requires more setup per block, making
the broadcast-to-256 step a measurable bottleneck that the 1-uop
vbroadcasti128 eliminates.
### Why these 132 instances matter
Every dequantization path (IQ2_XXS through IQ6_K, Q4_0 through Q8_1,
MXFP4) starts by broadcasting a 128-bit lookup table or scale vector
to 256 bits. These are in the inner loop of every quantization format's
dot-product kernel. Reducing each broadcast from 3 uops to 1 uop
cumulatively reduces port-5 pressure across the entire dequant
pipeline.
## Scope
This change touches only the identity-broadcast case (both MM256_SET_M128I
arguments are identical).
* IQK AVX2: Introduce MM256_SET1_M128I(x) wrapper for identity-broadcast pattern
Per Ikawrakow's review: replace direct _mm256_broadcastsi128_si256(x) with
a new macro MM256_SET1_M128I(x) wrapping the intrinsic, so that if the
broadcast turns out harmful on some CPU, only the macro definition needs
changing, not 132 call sites.
#define MM256_SET1_M128I(x) _mm256_broadcastsi128_si256(x)
The 132 identity-broadcast MM256_SET_M128I(x, x) call sites across 8 files
now use MM256_SET1_M128I(x) instead of the raw intrinsic.
* openpangu: Stage-1 converter probe for openPangu-2.0-Flash
Add OpenPanguV2ForCausalLM conversion support (converter-only; runtime graph
is Stage-2). Registers a new LLM_ARCH_OPENPANGU on the Python/gguf-py side:
- gguf-py/constants.py: MODEL_ARCH.OPENPANGU + name, indexer KV keys, 22 new
tensor enums (DSA indexer x4, MoME convs x3, param-sink x2, mHC/Hyper-
Connections x12, block-post-norm), and the full MODEL_TENSORS list reusing
the deepseek MLA + MoE + NextN bricks.
- tensor_mapping.py: arch-specific block mappings that disambiguate the
sandwich norms (post_attention/pre_mlp/post_mlp) and pin every Pangu-only
tensor; non-block global mHC merge module.
- convert_hf_to_gguf.py: OpenPanguV2Model (subclasses DeepseekV2Model) with
set_gguf_parameters (MLA/MoE/indexer/mHC/param-sink/DSA+SWA metadata),
modify_tensors (expert merge, kv_b split, no MTP skip), and the
OpenPanguV2Tokenizer pre-tokenizer hash.
Validated offline against the real 50-shard safetensors index: all 37,587
tensors map to a GGUF target (0 unmapped), and set_gguf_parameters reads only
hparams present in config.json. No weights downloaded; no GPU. Pinned on the
ik/dsa_loop_hadamard_blend DSA substrate.
* openpangu: Stage-2 arch scaffold (LLM_ARCH_OPENPANGU) — loadable, compiles
New arch on main (DSA-decoupled). Declares openPangu-2.0-Flash to the runtime so
the model loads into memory; the compute graph is the next step.
- llama-arch.{h,cpp}: LLM_ARCH_OPENPANGU + name; 3 KV keys (mhc_num_stream,
mhc_recur_norm, param_sink_number); 18 tensor enums (mHC x12, MoME conv x3,
param-sink x2, block-post-norm).
- llama-model.cpp: OPENPANGU tensor-name block, strings matched to the converter.
- llama-model.h: layer + model struct fields (mHC / conv / sink / block-post / merge).
- llama-hparams.{h,cpp}: reader (MLA + MoE + sigmoid gate + indexer + mHC +
param-sink + NextN); n_layer_kv_from_start = n_layer - nextn (MTP skipped).
- llama-load-tensors.cpp: create_openpangu_tensors (GLM-DSA MLA/MoE base + Pangu
tensors; indexer loaded-but-unused for dense fallback); dispatch + is_mla_attn.
Builds clean (CPU-only libllama). Dense-fallback design: no DSA indexer / SWA
windowing / MTP for first generation (exact <=512 tokens). Graph is Stage-2b.
* openpangu: fix compresskv_conv dim (kv_lora_rank, not +rope); pin attention order in spec
* openpangu: end-to-end runtime — build_openpangu graph runs, generates (garbled)
First full forward pass of openPangu-2.0-Flash on ik_llama. Pipeline works end to
end: new LLM_ARCH_OPENPANGU loads the Q4 GGUF, the graph executes, and llama-cli
generates 40 tokens (EXIT=0). Output is currently garbled (tensor-layout bug to
debug), but the structure is proven.
graphs/build_openpangu.cpp: dense decompressed-MHA attention + 4-stream mHC
(Hyper-Connections) with 20-iter Sinkhorn + MoE(sigmoid+shared) + sandwich norms
+ entry stream-repeat/tail-merge + inp_out_ids selection.
Bring-up fixes to load+run:
- llama-vocab.cpp: register 'openpangu' pre-tokenizer (QWEN2 family)
- llama.cpp: OPENPANGU -> LLAMA_ROPE_TYPE_NORM (was defaulting to NONE=-1)
- llama-load-tensors.cpp: full wkv_b load; k_b/v_b as flattened 2D; block_post_norm
dim = S*H (10240); conv weights 2D {3,C}; mHC alpha/beta/gamma + param_sink +
merge params use bare (no-.weight) tensor names
- llama-model.cpp: OPENPANGU is NOT is_mla_attn (decompressed MHA, standard KV cache)
- graph loop bounded to base layers (skip NextN/MTP)
v0 deferrals (need conv-state cache / manual attention path, all documented):
MoME convs (passthrough), o_conv, param_sink. Next: fix the layout bug to coherence.
* openpangu: COHERENT generation — NEOX rope, Sinkhorn orientation, MoME convs, param_sink
Four correctness fixes on top of the end-to-end scaffold, verified checkpoint-by-
checkpoint against a Python golden reference running on the GGUF's own dequantized
weights (block-0 activations now match to rounding at full fidelity):
- rope: NORM -> NEOX. Pangu config rope_interleave=false; the Infer source maps it
as is_neox_style = not rope_interleave (rotary_mode='half').
- mHC Sinkhorn: the flat h_res block is torch-[r,c] row-major, so a bare ggml
reshape lands column-fastest; the doubly-stochastic iteration ran transposed
(Sinkhorn is not transpose-symmetric). One transpose at input fixes the whole
chain including the mhc_post application.
- MoME convs (qa/compresskv/o): were passthrough stubs. Implemented as
out = x + causal_conv1d(x) (every Infer call site uses residual_connection=1;
tap stats confirm the perturbation form). Taps cast f16->f32 for ggml_mul.
Batch-local v0: exact for fresh-sequence prefill; decode steps miss the
t-1/t-2 taps until a conv-state cache exists.
- param_sink: 128 learned latent-KV entries prepended per layer via a manual
attention path (kv_store + explicit soft_max over [sinks ++ cache]); huge
effect at short context. o_conv now applied pre-o_proj on the same path.
flash_attn forced off for OPENPANGU (FA kernel cannot see the sinks).
- converter: add_bos_token=true (HF prepends <|pangu_text_start|> via the
post-processor; the key was absent so ik dropped BOS).
Greedy Q4_K_M smoke, chat template + <think>: coherent CoT reasoning and a
correct answer. Layer-0 instrumentation (opg0_* names) kept for now.
* openpangu: MoME conv-state cache — decode steps get real t-1/t-2 taps
Allocate a per-layer cache_s_l tensor for OPENPANGU base layers holding the last
two pre-conv latents of the three MoME sites, packed
[qa 2*1024 | compresskv 2*512 | o 2*6144] f32 (~60KB/layer). The conv helper
reads the [C,2] history window (zeros at sequence start, kv_head==0), builds
xx = [hist ++ x], and writes the last two columns back each ubatch — the concat
naturally handles both prefill chaining and the T==1 shift. Read precedes write
in graph order; the fixed-offset copy is graph-reuse safe.
Verified: prefill anchors unchanged (bit-path identical, zero-history branch);
-ub 1 token-by-token run matches the golden reference at t4 (qlora_conv 0.084,
R_block 0.008 rel; attn_out 0.15 on one channel = f16 KV-cache rounding, washes
out by post-norm); final logits differ from full-batch only by a common-mode
shift that softmax cancels. Chat-template greedy smoke: think-block repetition
is gone — clean structured CoT and correct answer.
v0 limits documented in the helper: one state slot (single sequence); cache
rewinds leave the state stale.
* openpangu: NextN/MTP speculative decoding — 1.7-1.8x TG on CPU
Wire the three NextN layers (46-48) into ik's MTP speculative framework
(--spec-type mtp). v0 drafts with head 1 (layer 46), self-chained by the
framework.
- llama.cpp: add OPENPANGU to the cparams.mtp arch allowlist (it was silently
zeroed, which left the target context without a logits buffer once the server
enabled embeddings -> GGML_ASSERT(lctx.logits) in speculative_is_compat).
- load-tensors: MTP layers carry no mHC tensors (tail_use_mhc=false in the
reference) — create them only for base layers. nextn.* tensors were already
wired by the Stage-1 probe.
- build_openpangu: extract the attention sublayer into
build_openpangu_attention (shared base/MTP); add build_openpangu_mtp:
eh_proj(cat(enorm(embed), hnorm(prev_hidden))) -> one plain-residual Pangu
block (sandwich norms, convs+param_sink, MoE+shexp, no mHC/block_post_norm)
-> shared_head norm+head. MTP branch returns the draft graph when
mtp_op_type != NONE; main graph keeps all-token outputs under cparams.mtp.
MTP convs run batch-local (no conv-state slot) — affects acceptance only.
A/B (Q4_K_M, CPU, greedy, 192-token chat CoT completion, warm back-to-back,
medians of 3, bracketed B/A/B):
no-spec: 2.44 t/s (2.34-2.86)
--spec-type mtp:n_max=3: 4.23 / 4.49 t/s (brackets) => ~1.7-1.8x
Draft acceptance 34% on CoT prose (46% on repetitive text); spec and no-spec
greedy outputs are byte-identical. Headroom: conv-state for MTP drafts, n_max
tuning, true 3-head chaining (spec_step_idx).
* server: include draft_n/draft_n_accepted in /completion timings
get_formated_timings() (the /completion path) omitted the speculative
counters that get_timings() (the OAI path) already reports; add them,
guarded by n_draft_total > 0 like the OAI path.
* openpangu: position-indexed MoME conv-state ring — rollback-safe spec decoding + MTP draft chaining
The v0 single-slot conv state held the last-2 pre-conv latents of the most
recent batch, so any speculative draft rejection left latents of REJECTED
positions in the state and every later decode ran with wrong t-1/t-2 taps
(3 conv sites x 46 layers). At 192-token greedy runs every spec config
diverged from no-spec, each differently (rejection-pattern dependent).
Replace it with a per-layer ring cache_s_l [n_lora_q+n_lora_kv+n_head*v_dim, 16]:
column pos%16 holds position pos's pre-conv latents ([qa|ckv|o] packed).
Invariant: reads target only positions before the first batch token, which
are committed, and committed latents depend only on the committed prefix -
rollback-safe by construction, no checkpointing. Writes cover the last
min(T,16) batch positions in <=2 contiguous cpy segments; the copy sources
are views of the [hist ++ x] concat so the history read is an ancestor of
every write (read-before-write by graph dependency).
The ring is also allocated for the NextN/MTP layers, so the draft head
chains real conv taps across WARMUP -> sequential DRAFT_GEN steps (was
batch-local zero-history per draft token).
graph_reuse is forced off for the arch: ring view offsets are position-
baked and the reuse patcher only updates the standard K/V-store copies.
Measured cost on the CPU server path: none visible. ggml_set_rows driven
by an input index tensor is the future reuse-safe shape.
Verified (Q4_K_M, CPU, greedy 192-tok chat-CoT, warm single process):
- no-spec output byte-identical to pre-ring build
- spec output byte-identical to no-spec below the n_predict cap, for all
of n_max in {1,2,3,4,6} x p_min in {0,0.3,0.6} (old build: all diverged)
- acceptance n3-p0: 33.9% -> 60.9%; n3-p0.3: 58.1% -> 68.9%
- TG medians: no-spec 3.19-3.32 t/s; mtp:n_max=3,p_min=0.3 6.97 t/s (~2.1x)
* openpangu: DSA lightning indexer + SWA schedule — long-context correctness past the dense fallback
The dense fallback was exact only <=512 tokens (SWA window). This wires the real
DSA/SWA hybrid schedule, self-contained from GGUF keys the converter already
writes (openpangu.swa_layers + sliding_window_list; absent keys keep the old
dense fallback):
- SWA layers (30 base @512): the generic inp_KQ_mask_swa path, per-layer mask
choice in the builder. The NextN/MTP layers are SWA @2048 in the checkpoint
schedule; MTP graphs run in their own context, so the mask fill picks
hparams.n_swa_mtp when built with an MTP op type.
- DSA layers (16, every 3rd): lightning indexer implemented in-graph from the
Infer reference semantics (jointfix _pangu_torch_calib): q_idx = wq_b on the
post-conv post-norm q-lora latent (24x128), k_idx = rms-normed wk(x) shared
across heads, both NEOX-roped on the FIRST n_rot channels; score =
sum_g w_g * relu(q_g . k) in f32, causal-masked, exact top-k via
argsort + ggml_set_rows scatter into a -1e30 base -> additive selection mask
on the existing manual soft_max seam. Selection engages only when the causal
window exceeds index_top_k (2048); below that the layer is exactly dense.
- Indexer keys cached per position (cache_idx_l, f32 [128, kv_size], DSA layers
only) with the same committed-position invariant as the conv-state ring, so
speculative rollbacks stay safe.
- param sinks remain outside both the window and the selection budget, matching
the reference.
Verified (Q4_K_M, CPU):
- <=512 tokens: byte-identical to the dense build (96/160-token greedy)
- indexer scores vs a GGUF-dequant golden reference at 2101 tokens: 1e-3 rel
(f16 weight rounding); top-3 selection indices exact on all compared queries
- >512 coherence clean; 3.4K-token needle retrieval through active selection
(needle outside every SWA window, ~1300 positions pruned) answers exactly
* openpangu: MLA-latent KV cache — attention absorbed into the 512-latent, 14x smaller cache, ~2.2x TG
Store per position only [ckv_norm 512 | roped k_pe 64] (f32, k_l) plus the
transposed 512-latent (f32, v_l, v_trans layout); per-head K/V are never
materialized. q_nope is absorbed through attn_k_b (loaded 2D from the
converter split for base layers; derived at load via llm_prepare_mla for the
NextN layers - now guarded for layers without attention weights, e.g. the
idle NextN heads 2/3). The value side is the latent itself, up-projected
through attn_v_b after the weighted sum, matching the Infer _forward_dsa
reference. param sinks are native latent-space entries, which removes the
per-step full-cache concat+cast that dominated long-context decode.
llama_state row sizes now come from llama_kv_k_row_embd/llama_kv_v_row_embd
(arch-aware), fixing an out-of-bounds crash in the server prompt-cache save
path (hparams-derived 9216-wide rows vs actual 576-wide latent rows).
Verified (Q4_K_M, CPU): layer-0 attention output matches an f32 golden
reference computed from the same GGUF weight encodings (~1e-2 on O(1)
values); MTP spec output byte-identical to no-spec; 3.4K needle retrieval
through active DSA selection exact under greedy. Output differs from the
materialized build at the token level because attn_k_b/attn_v_b are
independently quantized tensors - both are legitimate Q4-fidelity encodings.
Perf (CPU, warm): no-spec TG 3.2-3.3 -> 6.9-7.1 t/s; mtp:n_max=3,p_min=0.3
-> 11.1 t/s (byte-exact, 67% acceptance); prefill 30.5 t/s at 3.4K; KV self
size at 4K ctx: 5.5 GiB -> 391 MiB. Not yet supported on the latent cache:
K-shift/defrag (context shifting) - unreached in current usage.
* openpangu: fence unsupported serving modes, truth-pass comments, drop dead weight/keys
Post-audit hardening. The cache's position-indexed side state (MoME conv ring,
DSA indexer keys) made several generic serving paths silently unsound; they are
now fenced loudly instead of documented as unsupported:
- s_l_position_ring flag on llama_kv_cache: the qnext-state predicate no longer
claims the conv ring, so per-seq state save, seq_cp and the s_copy graph skip it
- state save/restore refused for the arch at every llama_state_* entry (the ring
and idx_l are not in the state format; restoring without them diverges silently)
- K-shift/self-extend assert, defrag skips with a warning, server ctx_shift off
via new llama_model_supports_ctx_shift()
- single sequence enforced at context creation (n_seq_max > 1 refused)
- server prompt-cache reuse limited to pure extension via new
llama_model_supports_partial_kv_reuse(): mid-cache divergence reprocesses from
scratch (the 16-column ring cannot rewind); multi-turn continuation stays fast
- MTP draft length clamped to 13 via new llama_model_max_draft_tokens() so a
rejected draft can never overwrite the ring columns the next decode reads
- K/V cache types forced to f32 for the arch so the KV size log reports the truth
- cache_size(): real latent-cache branch (was falling through to the ~14x larger
materialized estimate used for offload planning)
- unused fused wkv_b no longer loaded (TENSOR_SKIP; the graph runs entirely on the
pre-split k_b/v_b), llm_prepare_mla openPangu special-case removed (it was a no-op)
- stale v0 comments rewritten to describe the shipped graph; converter stops
writing dead keys (dsa_layers, block_post_layernorm_idx) and the tokenizer
pre-hash is registered in convert_hf_to_gguf_update.py
Gates on this build: greedy spec output byte-identical to no-spec (EOS-terminated,
sha-equal); 3.4K needle retrieved exactly; -np 2 / state save / n_max=20 / stale
prefix reuse all refused or clamped with clear messages.
* openpangu: assert kv_head == first batch position at graph build
The ring, indexer and latent stores are addressed by absolute position through
kv_head; the fences make append-only decode the only reachable mode, but the
invariant was unchecked. Assert it at both graph entries (base and MTP) so any
future cache plumbing that breaks it fails at build instead of corrupting
output. Worst-case measurement builds pass pos = null and are exempt.
* openpangu: cont h_pre before the mHC broadcast mul (CUDA binbcast misreads strided views)
h_pre is a row-slice view of the fused mixes tensor. The CPU mul handles the
strides; the CUDA broadcast path reads the view as if contiguous, so token 0
mixes correctly and every later token gets h_post/h_res rows instead. First
divergent node in the whole graph (oracle rel 0.36 at opg0_attn_mhcpre_x,
fixed to 7.5e-5). Sibling views h_post/h_res were already cont-wrapped, which
is why only h_pre was exposed.
* openpangu: keep DSA zero-trick sources finite (CUDA clamp propagates the 0*(-inf) NaN)
The selection-mask base and zeros were built by scaling the MASKED scores by
zero, but post-mask sc contains -inf and 0 * -inf = NaN. The CPU clamp launders
NaN back to -1e30 (fminf/fmaxf ignore NaN); the CUDA clamp propagates it, so
every DSA layer emitted NaN masks at n_kv > top_k and logits collapsed
(observed: eval-callback CLAMP sum -1.3e36 on CPU vs nan on CUDA, 11748 NaNs
downstream). Scale the pre-mask finite scores instead, which is correct on any
backend regardless of clamp NaN semantics. Also defensively cont the strided
KQ_mask slice feeding the score add (same strided-view kernel class as the mHC
h_pre fix; unproven here but cheap). Gates after fix: 2600-token probe coherent,
3.4K needle exact ('7391') with and without MTP speculation, PP ~120 t/s.
* openpangu: f16 latent KV cache option (explicit -ctk/-ctv f16 halves cache memory, f32 stays default)
Track explicit cache-type requests through CLI/env; openPangu resolves no-request
to f32 (unchanged), accepts explicit f32/f16, warns and falls back to f32 for
BF16/quantized. Sink and cached-token KQ paths stay separate until after KQ so
the latent cache is read directly without the f32-only concat; value is the sum
of the sink and cache matmuls. Ring and DSA indexer caches stay f32; cache_size()
follows the resolved types.
* openpangu: enable graph reuse
* openpangu: wire multi-head MTP drafting
* openpangu: add MTP heads override
* openpangu: keep MTP update logits last
* openpangu: scope MTP warmup heads
* speculative: apply per-request MTP heads before warmup
* openpangu: fix multi-head MTP warmup computing on unwritten inputs
Each chained head called the build_inp_* helpers itself, so the warmup and
update graphs held one inp_tokens/inp_pos/inp_out_ids/KQ_mask tensor per
head while llama_set_inputs only fills the tensors the lctx pointers
reference, i.e. the last head's copies. Every head but the last read
unwritten compute-buffer memory: with heads=3 active even head 1's ring,
latent cache, and cached one-token draft were computed from garbage, which
is why depth-1 acceptance measured 4% against 98% for the heads=1 control.
Create the batch inputs once in build_openpangu and pass them to every
build_openpangu_mtp call, and fix the two chaining errors that were hiding
behind the garbage inputs:
- Shift the chained hidden: head k+1's row at position p consumes head k's
output row at p-1, the same convention head 1 uses for the target's
conditioned hidden rows. The predecessor of a batch's first row lives in
the previous warmup/update, carried across decodes through a new
inp_mtp_carry input backed by lctx.mtp_carry (written back per ubatch,
zeroed when a prompt warmup restarts from position 0).
- Fill head 3's cache row at draft step 2: each draft step runs one head,
so head 3's own decode at step 3 attended over a never-written row at
the step-2 position. Pre-write it from the committed carry.
Also include the active head count in the graph-reuse key next to the
existing step index (reuse stays forced off for this arch).
* speculative: default MTP drafting to a single head
A stage without an explicit heads= override previously resolved to 0,
meaning all model heads, so multi-head drafting was silently on by
default for models that carry more than one NextN layer. Keep it opt-in
(heads=N or heads=0 for all) until multi-head measures a win over the
single-head config; single-head models are unaffected either way.
* speculative: fence MTP head upshift over a warmed prefix
Deeper NextN heads only hold valid cache rows for spans that were warmed
with them. A request drafting with more MTP heads than the cached prefix
was warmed with (e.g. a heads=1 conversation continued with heads=3, a
pure extension the divergence fence deliberately allows) would read
never-written deeper-head rows: verification keeps the output correct,
but acceptance quietly collapses and any measurement taken there is
misleading.
Track the minimum head count the committed context has been warmed with
since position 0 and have the server reprocess from scratch when a
request asks for more. Also announce the model's NextN head count and
the single-head default once at MTP context setup.
* openpangu: skip dead MTP chain compute and stall-free carry readback
The update chain's last head and the draft-time row fill only matter for
their latent-cache and conv-ring writes; their FFN, norms, and shared
head fed nothing. Add a cache-writes-only mode to the MTP block builder
that returns after the attention block, and use it at both sites.
The multi-head carry readback previously synchronized the scheduler
after every warmup/update decode, a hard stall on CUDA. Issue the
device-to-host copy async on the backend stream instead (stream order
protects the source buffer from later graphs) and synchronize lazily
when the host buffer is next consumed or resized.
* openpangu: stop emitting fused kv_b tensor
* openpangu: default latent cache to f16
* openpangu: refuse unsupported latent cache types
* Window OpenPangu SWA cache reads
* Gather OpenPangu DSA decode reads
Gather DSA decode attention over the selected latent rows for OpenPangu base-model decode and verify graphs. The gathered branch now uses ggml_top_k order directly, runs maskless softmax over sinks plus selected rows for T <= 14, and derives values from the gathered k_l rows instead of the transposed latent cache.
* Chunk OpenPangu indexer prefill scoring
* Chunk OpenPangu prefill attention
* Gather OpenPangu sparse prefill attention
* Drop OpenPangu value cache
* Add OpenPangu indexer cache type flag
* Add OpenPangu q8_0 latent cache type
Store the OpenPangu MLA latent K cache as q8_0 via -ctk q8_0 (about 0.53x of
f16); the default stays f16 so behavior is unchanged without the flag. Latent V
stays f16/f32.
The q8 latent cache is a storage format only: it is dequanted to F32 before all
compute. K reads go through openpangu_build_k_latent_for_read, V derivation
through openpangu_build_v_latent_from_k (full 576-wide row to F32, then slice),
and the DSA gather paths already dequant via get_rows. Feeding a q8 latent view
directly into the KQ mul_mat corrupts large-context prefill, so that path is
removed for quantized caches. The cache write stages ckv and kpe through F32 and
writes one full 576-wide q8 row per token.
Verified on a small discriminator model: the default f16 path is byte-identical
to the prior code; the first-DSA-layer attention envelope is within 0.6% of the
f16 cache (linf_rel 0.0057); top-k selection is bit-identical between cache
types; the q8 latent cache is 0.531x the f16 size at 8K and 32K context; and
generation stays coherent on both the dense and DSA-gather paths at all tested
context lengths.
* Remove OpenPangu debug trace env knobs and redundant DSA_TOPK override
Drop the five LLAMA_OPENPANGU_*_TRACE debug-logging knobs (DSA_GATHER_TRACE,
IDX_CHUNK_TRACE, ATT_CHUNK_TRACE, PREFILL_GATHER_TRACE, SWA_WINDOW_TRACE) and the
LLAMA_OPENPANGU_DSA_TOPK override, which duplicated the -dsatk / --dsa-top-k CLI
flag; top-k now comes solely from cparams.dsa_top_k. The five perf-tuning knobs
(DSA_GATHER, IDX_CHUNK, ATT_CHUNK, ATT_KQ_MAX_MIB, PREFILL_GATHER) are retained
pending the perf battery. No change to default behavior.
* Subchunk OpenPangu DSA prefill gather to fit CUDA grid limit
The prefill gathered-attention ggml_get_rows produced dst rows = topk *
token_chunk (2048 * 256 = 524288) mapped to the CUDA grid.y dimension, which
caps at 65535, crashing with GET_ROWS invalid argument at long context (N_KV
around 10.5K with the natural topk of 2048). Split the prefill gather into token
subchunks so topk * subchunk_tokens stays within the grid limit, and guard the
decode gather with the same fit check (falling back to the dense masked path if
a pathological topk would not fit). The subchunking is over the token dimension
only, so per-token attention is unchanged and the result is numerically
identical. Verified: the GPU sweep runs past the old crash boundary to 22K+ with
zero CUDA errors; CPU and -ctk q8_0 paths unaffected.
* openpangu: fix scheduler node budget for chunked DSA prefill; drop unused attn_kv_b; remove env tunables
- Size the scheduler graph node budget for the chunked DSA prefill so 32K/ub2048 no
longer trips the hash-set reservation assert; derive the extra budget from the
builder's chunk/top-k/window structure with a fixed safety margin.
- Remove LLAMA_OPENPANGU_* environment tunables from both the node-budget estimator
and build_openpangu.cpp; use fixed constants in both so they stay in sync.
- Converter: emit only the split attn_k_b/attn_v_b projections and drop the unused
fused attn_kv_b tensor.
* openpangu: restore DeepSeek converter kv_b; drop trace env + dead code; fix dense-fallback node budget
- convert_hf_to_gguf.py: restore fused attn_kv_b in DeepseekV2Model (shared
parent); openPangu subclass keeps split-only k_b/v_b. Stops newly-converted
DeepSeek GGUFs from failing to load.
- src/llama.cpp: remove LLAMA_GRAPH_REUSE_TRACE getenv, hit/miss counters, and
the unconditional destructor log (no getenv or behavior change for any arch);
node-budget estimator now covers the dense-fallback (n_swa==0) attention-chunk
loop while skipping absent idx/top-k terms, preserving a strict overcount;
remove unreachable openPangu split-cache block.
- src/llama-context.h: drop now-dead graph_reuse_hits/misses members.
- include/llama.h: move type_k/type_v/idx_type_k *_explicit bools to struct end
to avoid a mid-struct ABI shift for out-of-tree consumers.
- src/graphs/build_openpangu.cpp: replace vestigial env-struct singletons with
the OPENPANGU_* constants; drop a redundant Sinkhorn permute round-trip
(one transpose; greedy output verified byte-identical).
Decode output unchanged (byte-identical greedy generation verified); shared-file
changes are openPangu-gated or restore the pre-PR baseline.
* openpangu: chat-parser support (reasoning split + thinking toggle)
Two openPangu-only fixes, both gated on the arch-unique token
<|pangu_text_start|> so no other model's parsing changes.
- chat-diff-analyzer: add a workarounds entry that force-sets TAG_BASED
reasoning with an empty start and a </think> end. openPangu prefills
<think> in the generation prompt, so the output is delimited only by
</think>; the differential detector otherwise learns start="<think>"
from the assistant-history form and fails to split, leaking reasoning
into content. Same shape as the existing Laguna prefill patch.
- chat.cpp: bridge enable_thinking to the template's `thinking` variable.
openPangu's template gates reasoning on `thinking` rather than the
ecosystem-standard `enable_thinking`, so the standard toggle was inert.
An explicit `thinking` chat_template_kwarg still overrides via the
extra_context merge.
Blast radius: test-chat-auto-parser 437/437 unchanged; the sole
test-chat-template diff is a pre-existing GLM trailing-newline.
* openpangu: use ggml_cast for latent dequant reads
Replace ggml_cpy(view, ggml_new_tensor_2d(F32, ...)) with ggml_cast in the MLA
latent V-from-K and K-read helpers. ggml_cast emits the identical GGML_OP_CPY
node into a fresh f32 tensor, so behavior is unchanged; it is the idiomatic
form. Per review.
* openpangu: narrow SWA reuse-key fields to 32-bit
The openpangu_swa_window_view reuse key stored n_kv/n_tokens/window/pad as
int64_t, but these are bounded well under 2^31 (window/pad are uint32_t at
source; n_kv/n_tokens <= context length). Narrow to int32_t/uint32_t and drop
the widening casts. w_view/win_off stay int64_t: they feed ggml view
dims/offsets. Per review.
* openpangu: precompute param_sink derived tensors at load
The per-layer attention-sink block (sink_blk [576,NS]) and its transposed
latent (s_lat_t [NS,512]) are pure functions of the layer weights, yet were
rebuilt every eval across all 49 layers (RMS-norm + cast + concat + transpose).
Compute them once at load, mirroring the wk_b derived-weight precompute, and
read the stored tensors in build_openpangu_attention. Numerically identical;
removes per-token work at decode.
* openpangu: replace conv position-ring with ggml_ssm_conv + spec-rollback checkpoint
Migrate the MoME depthwise causal conv (three sites per attention sublayer:
qa-lora, compressed-kv, attn-out) from the bespoke 16-column position-indexed
ring onto the core ggml_ssm_conv op with a recurrent conv-state slot.
Cache: s_l becomes [2*conv_col_ne, qnext_state_slots], holding the (d_conv-1)=2
history taps per channel for the three sites (float offsets 0 / 2*n_lora_q /
2*(n_lora_q+n_lora_kv)). Drops the conv_hist_idx / conv_write_idx graph inputs
and their fill in llama_set_inputs; adds one single-sequence sq input for
ggml_ssm_conv shared across the three sites and the MTP head.
Speculative rollback: the position ring self-healed rejected draft columns by
absolute position; a recurrent slot does not, since seq_rm is a no-op for
recurrent state. openPangu is admitted at the three spec-checkpoint save/init
gates so the whole-slot shadow checkpoint (gpu-fallback) snapshots the conv
slot before drafting and restores it before the accepted-token replay. The
restore path is already keyed on ckpt.valid, so no gate change is needed there.
Per-step checkpoint mode is declined for openPangu, which has no SSM recurrent
term, so auto mode resolves to the whole-slot shadow.
Gated: non-spec needle unchanged; MTP-spec needle correct with healthy draft
acceptance (rollback verified via the acceptance canary).
* openpangu: single ggml_concat copy for the latent cache store
The non-quantized latent store split the [ckv | roped k_pe] row into two views
and two cache copies, with a base_offset field on the CacheCopy struct to place
the second one. Match the quantized path: concat the two parts and do one copy
into the cache row. This drops the second cache-copy slot (OPENPANGU_COPY_K_KPE)
and removes base_offset from CacheCopy entirely.
Cache contents are unchanged: the concat writes the same [ckv 512 | k_pe 64]
bytes to the same row. Gated on the needle for both the f16 latent path (the one
that changed) and the q8 latent path, plus coherence.
* openpangu: reuse the shared kr_l indexer cache instead of a separate idx_l
The DSA lightning indexer stored its per-position keys in an openPangu-only idx_l
cache, parallel to the kr_l indexer cache GLM-DSA already uses. Both have the same
storage contract: [indexer_head_size, kv_size], idx_type_k dtype, one row per KV
cell, written at kv_head and read [dim, n_kv] from zero. openPangu now allocates
its indexer keys into kr_l and shares the dsa_cache_copies graph-reuse fixup.
The fixup patch is factored into a helper that both the generic path and the
openPangu update_cache_copies branch call, so the openPangu indexer copy is
repointed to the current kv_head on graph reuse like every other cache write.
This drops the idx_l vector, its allocation and memory accounting, and the
openPangu third cache-copy slot (now one latent copy per layer).
Per-arch allocation predicates stay separate (GLM uses indexer_is_full, openPangu
uses the window==0 DSA schedule); only the kr_l storage and the copy fixup are
shared. openPangu keeps its no-shift/no-defrag/no-state-I/O behavior, and the GLM
Hadamard/k-shift logic stays GLM-gated.
Gated: needle correct on f16 and q8 latent caches and under MTP speculation
(acceptance unchanged at 0.67), plus coherence.
* openpangu: discard pos-0 graphs from reuse; retire stale conv-state comments
The ggml_ssm_conv refactor bakes the pos-0 conv-state reset into graph
topology (a scale-by-zero node on the state view). A graph built at pos 0
could be reused at pos > 0 when the batch shape and padded n_kv match (a
1-token prompt followed by TG is the concrete case), zeroing the conv
history on every reused decode. Admit openPangu at the existing
reset_previous gate so pos-0 graphs are discarded from reuse, the same
guard the qnext recurrent state relies on.
Also retire the internal phase-plan comments the conv refactor left
behind: they claimed the spec-checkpoint wiring had not landed in the
commit that landed it, and misdescribed the s_l slot as awaiting rollback
support.
Gated: needle 8457 on f16 and q8 latent, MTP-spec needle (drafts fully
accepted), coherence.
* openpangu: drop the _explicit cache-type plumbing; validate unconditionally
Review follow-up (item 1 of the second review). The explicit/default
distinction carried less than claimed: the latent K/V fallback was f16,
which is already the -ctk/-ctv and API default, so distinguishing unset
from set-to-the-default bought nothing, and the two bools were behaviorally
redundant. The only load-bearing use was the indexer cache, where openPangu
defaulted to f32 while -ictk defaults to f16. Gating the f16 indexer
directly (needle on f16 and q8 latent paths, MTP speculation, coherence)
shows no quality difference, so openPangu now takes the standard f16
indexer default and the f32 special case is gone. Default indexer cache
memory halves (64 -> 32 MiB at c 8192).
Removes type_k_explicit/type_v_explicit/idx_type_k_explicit from llama.h,
the cparams/mparams plumbing, and common; the resolve helpers become plain
unconditional validators, so -ctk q8_0 is honored and an unsupported type
errors out at load instead of silently coercing.
Gated: needle 8457 on the new f16-indexer default, on q8 latent with MTP
speculation, and with -ictk f32 explicitly honored (64 MiB f32 buffer in
the load log); -ctk q4_0 and -ictk q4_1 refused with a clear error.
* openpangu: keep MTP draft decodes position-contiguous under speculation
The MTP framework's one-token draft shortcut caches a prediction one row
past the accepted prefix during the accepted-token update, then skips
re-decoding the last sampled token at the next draft round. A
mask-addressed cache tolerates the resulting position gap; openPangu's
position-addressed append-only cache (cell == position) does not: after a
rollback the next draft decode lands one cell behind its position, and
after a full acceptance the cache head sits one row ahead of the next
draft base, either way tripping the kv_head == pos[0] invariant and
aborting the server. The checkpoint admission in the conv refactor made
this the standard openPangu speculative flow; the needle-first gates
never generated enough draft rounds against a short prompt to reach it.
Decline the shortcut re-seed for openPangu in mtp_accept_batch (restoring
the drafting behavior all measured acceptance numbers were taken on) and
trim rows at or beyond the draft base in mtp_speculative_gen_draft, so
every draft decode stays position-contiguous with the cache head.
Gated: the crashing flow (short prompt, 512-token spec generation, then a
second request) completes with acceptance 0.60 prose / 0.87 code,
matching the pre-checkpoint baseline profile; needle 8457 plus coherence
on f16+spec and q8+spec.
* openpangu: remove stale ring limits and fix MTP graph reuse
* cli: preserve speculative carry on fallback
Decode an already-emitted pending token when a draft cannot be used instead of sampling unchanged logits and duplicating output. Document single-head MTP as the default and multi-head modes as experimental.
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* Add --prefetch-experts to stream mmap'd MoE experts into page cache
* Drop fds, fault experts in with MADV_POPULATE_READ instead of pread
* Remove stale note about pread workers
* Move MoE prefetch behind ggml_backend_prefetch_* wrappers
* Cleanup stale comments
* Add --prefetch-experts-threads, drop GGML_MOE_PREFETCH_THREADS env var
When building with clang-cl (MSVC + Clang), the CMake MSVC branch defined
__AVXVNNI__ as a preprocessor macro alongside /arch:AVX2, but clang-cl
requires the actual -mavxvnni target feature flag to enable AVX-VNNI
codegen. Without it, clang-cl refused to inline _mm256_dpbusd_avx_epi32
and _mm256_dpwssd_avx_epi32 into functions compiled under /arch:AVX2,
causing 'requires target feature avxvnni' errors in:
- ggml-quants.c (mul_sum_us8_pairs_float)
- iqk_gemm_iquants.cpp (mul_mat_iq3_xxs_r4_q8_k)
- iqk_gemm_kquants.cpp (mul_mat_q3_k_r4_q8_k)
- iqk_gemm_legacy_quants.cpp (dot, accum_q4_0_quants, operator())
Fix: Detect clang-cl via CMAKE_CXX_COMPILER_ID STREQUAL 'Clang' and
append -mavxvnni to ARCH_FLAGS instead of manual __AVXVNNI__ define.
Also add missing GGML_AVXVNNI handling for the non-MSVC (GCC/Clang on
Linux) branch, passing -mavxvnni as expected.
The DSA lightning indexer hardcoded LLAMA_ROPE_TYPE_NEOX for indexer Q/K
positional encoding, but the GLM-5.2 model config.json explicitly sets
"indexer_rope_interleave": true, meaning the indexer uses interleaved
RoPE (LLAMA_ROPE_TYPE_NORM).
GLM-5.1 config.json also sets "rope_interleave": true.
The mismatch caused incorrect indexer scores, wrong top-k key selection,
and single-character token errors and other corruption that could grow
worse with context length.
Here the setup is changed to use the model's given rope_type instead of
hardcoding NEOX. This fixes GLM-5.2, GLM-5.1, and should be less brittle
for future GLM variants.
* CUDA: use the mask tensor stride (nb31) in the tile FA kernels, not ne11
The tile flash-attention kernels (no-tensor-core GPUs, e.g. Pascal/sm_60) indexed
the KQ mask using ne11 (= K->ne[1]) as the row stride. The SWA windowing in
ggml_cuda_flash_attn_ext (the n_swa branch) re-points K/V/mask to the last nton
tokens, setting ne11 = nton while the mask keeps its original row stride nb31.
Indexing by ne11 then reads across mask rows and yields NaN once the window slice
engages (context past nton). Use the mask's own stride nb31/sizeof(half); in the
non-sliced case this equals ne11, so it is a no-op there. Matches how the vec
kernel already computes its mask offset, and how mainline llama.cpp fixed the same
latent bug in its unified tile kernel.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* CUDA: apply attention sinks in the tile FA kernels (gpt-oss)
The tile flash-attention kernels took a sinks argument but never applied it, so
gpt-oss (which has a per-head sink logit) got a wrong softmax denominator with
-fa 1 while the -fa 0 path (ggml_soft_max_add_sinks) matched CPU. Apply the sink
after the KV loop, mirroring the vec kernel: the sink joins the running max and
adds exp(sink - max) to the denominator once, rescaling kqsum and VKQ. Only ip==0
adds it so it is counted once across the parallel_blocks KV split (the epilogue
writes the sink-inclusive max/denominator into dst_meta before the combine). The
per-head index is blockIdx.y (the same index the slope uses). Guarded by non-null
sinks, so non-sink models are unaffected. This is the tile-kernel equivalent of
mainline llama.cpp PR #15178; the wmma kernel is left untouched (no Turing/Volta
hardware to validate here, but the same defect likely applies).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
On GPUs without FP16 tensor cores (Pascal / sm_60, e.g. Tesla P100) MLA
flash-attention decode falls back to the CPU. The !fp16_mma_available path
routes decode to the f16 vector kernel, whose is_supported check requires
K == V head sizes; MLA's absorbed head sizes are 576/512 (asymmetric), so it is
rejected and attention runs on the CPU. With --cpu-moe that recomputes the full
MLA attention on the CPU every decoded token, which dominates decode at long
context.
Route Pascal MLA decode (Q->ne[1] <= 8 && K == 576 && V == 512) to the f32
vector kernel and enable that kernel for the 576/512 case, including Q8_0 KV.
Scope: decode only (batch <= 8). Prefill (batch > 8) and -fa 0 are untouched;
tensor-core GPUs never reach this branch. Aligned head sizes are byte-identical
(the asymmetric/Q8_0 work folds to a no-op at compile time), so no other model
or configuration is affected.
- fattn.cu: route 576/512 decode to vec_f32 in the !fp16_mma dispatch and its
is_supported mirror.
- fattn-vec-f32.cu/.cuh: accept + instantiate 576/512 (F16 and Q8_0); fix latent
issues exposed by the first asymmetric/large-head use (KQ-row granularity uses
FATTN_KQ_STRIDE not Dv; guard the dst store to tid < Dv; guard the softmax exp
on the KV tail; size Q_i32 by ceil; only convert K/V to F16 when the type is
F16). All are no-ops for the previously-exercised symmetric cases.
- fattn-vec-f32.cuh / fattn-vec-common.cuh: guard the Q8_0 ragged tail (Dk=576 is
144 int32 lanes = 4.5 warps) with the ragged-dim idiom; compile-time-constant
for aligned head dims, so it folds away.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* GLM-DSA: improve TG performance even more
* Fix crash when using MTP
* GLM-DSA: much better PP performance (CPU-only)
* GLM-DSA: allow for quantized indexer cache
