* model : fold Ling 3.0 VL into the BailingMoeV3 architecture
Assisted-by: Scout
* model : keep shared NORM rope list intact when gating bailingmoe3 on mrope sections
---------
Co-authored-by: aetherbird <aetherbird@users.noreply.github.com>
* test-save-load-state : print a per-model results table in --models mode
in --models mode the output was very heavy: every model printed its
token dumps, per-test headers and PASS lines. instead, silence all
logging except the table itself (common_log_set_verbosity_thold(0)
leaves only LOG / LOG_LEVEL_OUTPUT) and print one row per model with
one column per test, colored PASS/FAIL/SKIP cells, row by row.
- run_save_load_tests_for_model returns a test_suite with a dynamic
std::vector<test_status> and continues past failures: tests 3-5 are
SKIPped when the baseline (test 1) fails, model init failure skips all
- per-test token dumps, test headers and PASS lines are demoted to
LOGV(LOG_LEVEL_INFO, ...) so they still show in single-model mode
- the table header/rows derive their columns from test_names; the
model name is printed and flushed before the suite runs so the model
currently in flight is always visible
- single-model output and exit codes are unchanged
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* test-save-load-state : print example usage on -h
add a print_usage callback passed to common_params_parse, so -h/--help
also shows example commands for the tool-specific --models option and
the -lv verbosity level
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* test-save-load-state : remove comments
ref: https://github.com/ggml-org/llama.cpp/pull/29316
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests: add backend option to test-llama-archs
* Update tests/test-llama-archs.cpp
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* remove extra space
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* jinja : parse unary +/- before variables
Lexer already emits unary_operator for -n / +n, and runtime executes
unary -. Parse them at multiplicative precedence so slices like
items[:-n] and GigaChat indent[:-indent_factor] work.
* jinja : keep filters/tests outside unary operands
Unary +/- must bind only the primary/postfix operand so -n|abs is
(-n)|abs, not -(n|abs). Add unary + and filter/test regression coverage.
Signed-off-by: sinksilk <785976238@qq.com>
---------
Signed-off-by: sinksilk <785976238@qq.com>
* vulkan: optimize IQ4_XS matmul kernels
Assisted-by: OpenAI Codex
* vulkan: address IQ4_XS review nits
- drop the dead LOAD_VEC_A != 8 branch in the IQ4_XS shmem load; iq4_xs is
in lut_load_vec_a()'s "8" list, so that path is never generated
- disable MMVQ for IQ4_XS on Intel (27.3% tg regression on A770)
- remove a stray empty line in types.glsl
Assisted-By: Claude Opus 5 <noreply@anthropic.com>
* metal : gate mul_mm_id src1 rescale behind ggml_prec
Assisted-by: Claude Fable 5.1
* ggml-webgpu: reject MUL_MAT_ID when src1 precision is F32
* cuda/vulkan: reject MUL_MAT_ID in supports_op when src1 prec is F32
fix `supports_op` to return false for failing backends when the specified src1 precision is f32
Assisted-by: Claude Fable 5.1
---------
Co-authored-by: yomaytk <yoshimura.masashi.frbs@gmail.com>
* tests/test-backend-ops : allow regex entries in the -o filter
so far -o only accepted a comma separated list of exact op names or
full test case strings. entries that are not plain op names are now
treated as regexes matched against the op name (e.g. "MUL_MAT.*"),
while plain names keep their exact-matching behavior so that
"-o ADD" does not match ADD_EX etc.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : don't print the FA vec slice log when not needed
* tests/test-backend-ops : reformat the help text
use the same style as the other tools, with separate sections for
modes, options, and examples
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* hex-dma64: enable support extended buffer mappings and 64bit dma
hex-dma64: expand binary ops to support more DMA scenarios
hex-dma64: add binary-ops.h
hex-dma64: add --hex-dma64 to run.py and fix minor issues
hex-dma64: update SSM_CONV to use dma with proper support for 64bit
hex-ops: remove obsolete gate for % 128 in binary ops
hex-l2: dont check weight tensors against dirty ranges
hex-dma64: most binary ops now support dma
hex-dma: use dma_addr_t instead of plain uint64_t to avoid overhead on older targets
hex-dma: update all dma users to use dma_data (instead of pointers)
hex-dma64: simplify lazy buffer mapping and clonning
hex-fusion: factor out try_fuse_common that checks for dma64 buffers
hex-bufs: minor cleanup for mmaping logic
hex-bufs: simplify buffer clonning
hex-ssm-conv: tighten gating checks and check vtcm size in kparams
hex-binary: fix incorred mod/wrap in scalar ops
hex-binary: make sure to call precompute kparams in support checks
hex-dma64: update addr handling in mm,concat,binary
hex-dma64: fixing up leftover of dma_addr_t conversion
hex-binary: redo the kernel selection again and fix regressions in MOEs
hex-binary: specialize per-type/per-op
hex-binary: vtcm-layout and per-src dma-queue
hex-dma64: update dma_push to transparently handle 64bit/extended
* hex-cpy: fix improper rebase with the fixes for cont. tensors
* hex-dma-cpy: update CPY to use safe dma rows/size limits
* hex-mmap: bump number of mmaps to 64 to allow avoid eviction in larger models
* hex-dma: add support for the secondary ring as a fallback for too-large transactions
* hex-rope: fix freq_factors access with 64bit dma
* hex-dma: audit all ops for proper use/gards for 64bit addresses
* hex-dma64: uninline glu-compute funcs to avoid register pressure due to 64bit addr math
* hex-dma64: refactor binary ops to separate dma loops
* hex-devel: add inspect script to help with dbg and analysis
* hex-dma: refactor dma-pipelines in unary-ops
* hex-dma: rewrite softmax to use dma
* hex-dma: rewrite GDN dma loops and improve HVX register usage
* hex-gdn: fuse GDN+CPY
* hex-mm: factor out HVX solver
* hex-mm: remove hvx-flat kernels, the chunked version now handles vtcm limits much better
* hex-buffs: reject huge buffer allocations that we cannot memory map
* hex-inspect: add logic to look for float promo calls
* hex-mm: reduce HVX register spills in HVX prompt kernels
* hex-bufs: do not double count buffers from tensors in the same op
* hex-roll: fix merge conflict
* hex-dma: reroute all matmul ddr kernels to new chunked dma/vtcm kernels
* hex-dev: update developer docs to include inspection for register spils and float promos
* hex-ops: forgot to add new headers
* hex-softmax: fix gpt-oss dims
* hex-dma64: cleanup dma_addr_t casts
* hex-dma64: add support for dma/vtcm for flash-atten with sinks
* hex-mm-add: fix MUL_MAT+ADD fusion with bias.weights in extended bufs
* hex-add-id: add support for dma for src1 (exp. table)
* hex-dma: imrpove v73 fallback paths
* hex-bufs: do not drop extended mappings during va defrag
* hex-scripts: fix flake8 warnings
* hex-docs: fix editor-config warnings
* hex-inspect: fix warnings from ty
the dsv4_hc_pre kernels hardcoded hc = 4 via a constexpr used with
simd_shuffle, so the op was rejected by supports_op for any other hc
and fell back to CPU. Kimi-K3 uses dsv4_hc_pre with hc equal to the
number of banked checkpoints in the cross-layer residual stack, which
grows with the layer index.
pass n_hc as a function constant (FC_DSV4_HC) with per-n_hc pipeline
variants, and loop over it in both pre kernels with direct loads
add test-backend-ops cases for hc = 1, 2, 3, 5, 8 and 65, gated and
not gated
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal: add F16 input to the FWHT
The Metal FWHT kernel accepts F32 input only. This change makes the source
type a template parameter, so the kernel reads an F16 source directly instead
of requiring a converted copy. The F32 instantiations are unchanged.
The pipeline name now carries the source type, and supports_op accepts an F16
src1 for the Hadamard hint at the four sizes the kernels cover. Every other
F16 src1 path still goes through ggml_metal_supports_mul_mat_op.
These are the test cases mentioned in #27779.
test-backend-ops on M5 Pro: MUL_MAT_HADAMARD 16/16, MUL_MAT 1265/1265.
* metal: ask the same FWHT question in supports_op and the dispatch
supports_op admitted an F16 src1 on the type, the hint and the width alone, but the
dispatch also requires src1 and dst to be contiguous and the same shape. A Hadamard
hinted MUL_MAT that passed the first and failed the second reached the generic path,
which has no F32 src0 by F16 src1 kernel, and aborted on a nil pipeline:
kernel not found in any metal library: base = 'kernel_mul_mv_f32_f16_4'
ggml_metal_encoder_set_pipeline: nil Metal pipeline
ggml_metal_use_fwht now holds the whole condition and both callers use it, so they
cannot drift apart again. The added test case has src1 and dst of different shapes,
which aborted before this change and is declined by the Metal backend after it.
* metal: branchless butterfly select in the FWHT simdgroup kernel
Review suggestion. Replaces the ternary in the shuffle stages with
val2 - val + 2*((lane & i) == 0)*val, which is the same value without the
select.
Measured on M5 Pro, interleaved A/B, five rounds, first discarded, on a
Hadamard matmul with block 512 and 65536 rows so the kernel rather than the
launch dominates: 1324.6 us before, 1285.0 us after, a 3.0% gain, and faster
in every round. At the shapes already in the perf suite the op runs 1.6 to
3.9 us against a 1.6 us launch floor, so the difference is not visible there.
FOR_UNROLL on the same loops was also measured and made no difference, the
delta changing sign between rounds, so it is not included.
* metal: move the FWHT dispatch predicates to ggml-metal-common
Review feedback. ggml_metal_use_fwht and ggml_metal_fwht_supported_size were
static inline in ggml-metal-device.h. They now follow the
ggml_metal_op_mul_mat_use_mm pattern: declared in ggml-metal-common.h and
implemented in ggml-metal-common.cpp, which is already the home for helpers
shared between supports_op and the op dispatch. The predicate is named
ggml_metal_op_mul_mat_use_fwht to sit alongside the _use_mm pair it parallels.
This also fixes the macos-latest-arm64 build. The header needed ggml-impl.h
for ggml_get_op_params_i32, but ggml-metal-device.h is reached from
tools/tuning through ggml-metal-tuning.h, and that target does not have
ggml/src on its include path. ggml-metal-common.cpp already includes
ggml-impl.h, so the accessor is used normally there and the header goes back
to needing nothing extra.
* metal: keep the FWHT size check internal and group the dispatch helpers
Applies the patch from the review. ggml_metal_fwht_supported_size becomes
static in ggml-metal-common.cpp since nothing outside it needs the size list,
which also drops stdint.h from the header again, and
ggml_metal_op_mul_mat_use_fwht joins the existing _use_mm declarations under
their shared comment instead of carrying its own block.
* tests: drop the mismatched-shape Hadamard case
I added a case with m != k to cover an abort, but the hint is a promise that
src0 is a Hadamard matrix, so src0 is square and dst has the same shape as
src1. Every other case in the suite holds to that. The case was not a valid
op, and on CPU it compared the FWHT against a real matmul of a non-square
src0, which cannot agree.
The supports_op and dispatch conditions still come from one predicate, which
is what keeps them from disagreeing on contiguity.
* chat: add dedicated Ling 3.0 (Bailing V3) parser
Ling 3.0 Flash templates pre-open the think block in the generation
prompt, so the model never emits an opening <think>, and a tool call can
arrive before any </think>. The generated autoparser terminated reasoning
only at the close tag, which classified such tool calls entirely as
reasoning_content: clients received content="" with no tool_calls and
agent loops died as reasoning-only turns.
Adds a specialized parser that terminates reasoning at the think close
tag or at a <tool_call> start, mirroring the hand-written Qwen3-Coder and
Kimi K3 parsers and the reference vLLM/SGLang Ling3 parser (which treats
<tool_call> as an implicit reasoning terminator). Detection is gated on
the <role>...</role> section markers, unique to this family among the
tagged-argument templates.
Adds the Ling 3.0 Flash chat template and tests covering the
unclosed-think tool call (full parse and streaming), healthy closed-think
paths, trailing prose, parallel calls, marker-like strings in argument
values, string-union and non-string argument types, and
reasoning_format=none.
Assisted-by: Kimi Code
* tests : move Ling 3.0 test
---------
Co-authored-by: aetherbird <aetherbird@users.noreply.github.com>
Co-authored-by: Alde Rojas <hello@alde.dev>
* metal : add top-k MoE fusion
Adds a Metal fusion for SOFT_MAX + ARGSORT + GET_ROWS with optional
routing-weight normalization and scale, matching the top-k MoE fusion
available in the CUDA and Vulkan backends. The fused kernel writes the
selected expert ids and routing weights directly, eliding the separate
softmax, argsort, get-rows, sum-rows, clamp, div and scale kernels.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add MoE weighted reduction fusion
Fuses MUL(experts, weights) plus the expert VIEW/ADD chain into one kernel
that computes the weighted sum directly. The graph_optimize hook keeps the
expert and weight buffers alive until the fused output so the allocator cannot
reuse them while the kernel is still reading them.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* tests : expose MoE weighted reduction in fusion baseline
Use 2 experts per token in the generated MoE test models so the Metal
MoE weighted reduction fusion (MUL + ADD) is exercised by test-fusion.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : fuse RMS_NORM + SCALE
Adds NORM/RMS_NORM + SCALE fusion to the Metal backend by reusing the
norm+mul kernel with a scalar scale flag. Adds test coverage for both
NORM+SCALE and RMS_NORM+SCALE and regenerates the fusion baseline.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constant for RMS_NORM + SCALE
Replaces the runtime use_scale karg with a Metal function constant. The
norm+mul kernel is compiled with FC_norm_use_scale=false for MUL fusion and
FC_norm_use_scale=true for SCALE fusion, so the fused kernel has no runtime
branch.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constant for top-k MoE with_norm
Replaces the runtime with_norm karg with a Metal function constant. The
top-k MoE kernel is compiled separately for the normalized and non-normalized
routing variants, removing the runtime branch.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename moe_weighted_reduction suffix to moe_reduce
Shortens the MoE weighted-reduction fusion identifiers, kernel, pipeline,
matcher, args struct, and test op name from moe_weighted_reduction to
moe_reduce.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion
Adds dense mat-vec activation fusion for sigmoid/silu and bias+softplus.
The mat-vec kernels apply the activation/bias epilogue via function
constants, avoiding the separate unary/add passes.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : revert MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion
The mat-vec activation fusion regressed decode throughput on Qwen3.6-35B-A3B
by ~8% (tg32 81.5 vs 88.5 t/s). The regression is caused by loss of
concurrency: the standalone unary kernels previously overlapped with other
mat-vec work, while fusing the activation into the mat-vec kernel serializes
it on the critical path.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add SSM_CONV + UNARY (silu) fusion
The SSM_CONV kernels apply silu directly via a function constant, eliding
the separate unary pass. Regenerates the fusion baseline.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : address fusion review comments
- Fix declaration/table alignment
- Rename top-k MoE kargs fields to val_clamp / val_scale
- Move moe-reduce alloc-deps handling into a general fusion helper
- Remove the public moe-reduce matcher API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : fix unused parameter in top-k MoE fusion check
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : guard SSM_CONV fusion lookup behind use_fusion
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : track all fused outputs in graph reorder
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : keep top-k MoE logits alive until fused output
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : refactor alloc deps to pattern-driven approach
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : check fused kernel destination in concurrency tracking
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* meta : forward graph_optimize to underlying backends
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use vector for fusion table
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* meta : keep graph_optimize unimplemented
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* parallel : fix non-deterministic prompt selection
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* parallel : support dummy models and add global logits run hash
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : sync cross-device copies with destination completion event
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : avoid const_cast in fusion alloc deps
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : skip fusions with aliased sources
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : hide fusion pattern definition
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use vector fusion op sequences
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : drop redundant struct keywords
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add alloc deps comment separator
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : generalize fusion output memory ranges
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename fusion out_offsets to outs
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : avoid dst vector in memory range check
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : optimize fusion matching and multi-output handling
- use pointer arithmetic for fusion info count lookup
- avoid heap allocations in top-k MoE and MoE reduce pattern matchers
- use fusion outs for multi-output subgraph checks
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* Revert "parallel : support dummy models and add global logits run hash"
This reverts commit 57c7caf941c1b43c270fd5009c9f175063522e96.
* fusion : update MTL.csv
* metal : unroll constant loops in top-k MoE kernel
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constants for top-k MoE n_expert and top_k
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename fusion kargs to scale and clamp
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constants for moe_reduce and ssm_conv
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* fusion : update MTL.csv
Allow HMX flash-attention to run with head_dim not a multiple of 64
(e.g. SigLIP head_dim=72), by operating on DK/DV rounded up to 64 with
zero-filled tail lanes.
* vulkan: raise the hoisted row-id limit for mul_mat_id to 512 experts
The expert-count shader (count_experts.comp) sizes its shared arrays
with BLOCK_SIZE, which is 256. Because of that, row-id hoisting is
switched off for any model with more than 256 experts, and every
mul_mat_id workgroup has to rescan the whole ids tensor on its own.
Qwen3.8-Flash-Next has 512 experts and was quietly running on that
slow path.
This change sizes the arrays with a separate MAX_EXPERTS constant (512),
clears them in a loop instead of one entry per thread, and raises the
matching limit on the host side.
On Strix Halo at batch 2048 the expert matmuls drop from 12.5 to 9.5 ms
(iq3_s) and from 14.0 to 7.5 ms (iq4_nl) per op, and prompt processing
gets about 19 % faster at 8k tokens. test-backend-ops MUL_MAT_ID passes
(891/891) with new 512-expert test cases.
Assisted-by: Claude Fable 5.1
* vulkan: raise the hoisted row-id limit for mul_mat_id to 1024 experts
Follow-up to review feedback: 1024 matches LLAMA_MAX_EXPERTS instead of
stopping at 512. The three shared arrays in count_experts.comp grow to
3 * 1024 * 4 = 12 KiB, which fits the 16 KiB that Vulkan guarantees for
maxComputeSharedMemorySize.
Adds mul_mat_id test cases at 1024 experts alongside the existing 512
ones. test-backend-ops MUL_MAT_ID passes on Vulkan (RADV, Strix Halo,
Radeon 8060S): 889/889.
* gguf : align the data section relative to the GGUF start, not the file
gguf_init_from_file_ptr reads a GGUF from the current file position, but padded
the data section from file offset 0, so a GGUF embedded at an offset that is not
a multiple of the alignment loaded without error and returned wrong tensor data.
Also adds llama_adapter_lora_init_from_file_ptr, and disables mmap with a warning
when an embedded data section is not aligned, instead of asserting in ggml.
Assisted-by: Claude Opus 5
* llama : load lora from path through the FILE* variant
The test now checks that mmap is disabled only for an unaligned offset.
Assisted-by: Claude Fable 5.1
* Update ggml/src/gguf.cpp
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Update include/llama.h
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* llama : error on unaligned mmap of an embedded GGUF, drop test-load-file-ptr
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* model : add support for HrmTextForCausalLM (DFM Mimir 1B)
HRM-Text runs two transformer stacks (low, high) in an alternating cycle over the same token stream. The low-cycle state z_l starts from a learned [n_embd] tensor and is broadcast over positions.
- conversion: new writer for the fused gqkv projection (order gate,q,k,v) remapped to llama.cpp q/k/v plus a separate sigmoid gate tensor
- loader: block_count = lps * h_cycles * (l_cycles + 1) cache slots aliasing 2*lps physical blocks via struct copies
- graph: looped build with sigmoid-gated attention, SwiGLU FFN and parameterless RMS norms; learned embedding_scale applied in build_inp_embd
- saver: pointer-deduplicated layer loop (looped archs alias tensors)
- tests: hrm_text fixture (lps 1, h 2, l 3) in test-llama-archs
Limitations:
causal attention only - the upstream prefix-LM mode is not implemented (the prefix_lm GGUF key round-trips unused).
The KV cache holds one entry per pass: 128 layers for Mimir 1B, i.e. 4x a same-width 32-layer model - about 3072 MiB at ctx 4096 in F16 (halves with q8_0 KV + FA).
Every token runs all 128 block passes, so decode cost is roughly 4x a dense model of equal width (2.65 t/s BF16, 8-thread desktop CPU).
Verified against the HF reference: identical argmax at 334/334 positions across 20 prompts (BF16 GGUF vs FP32 golden).
q8_0 requant: 95.8% top-1, all remaining misses inside the HF top-5 (accumulated error over 128 sequential blocks).
AI usage disclosure: YES
Used GLM-5.3 for the majority of code AI-generated under my direction, all gates verified locally.
All in all I could say that I have written less than 20% of the code and most of the heavy lifting has been done by the model. As such, this should be considered experimental.
* Update conversion/hrm_text.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* convert : add gguf_writer methods for hrm_text metadata
replace raw add_uint32/add_bool calls with dedicated GGUFWriter methods, following the add_embedding_scale pattern
Assisted-by: GLM-5.3
* convert : map regular hrm_text tensors via tensor_mapping
delegate unfused checkpoints to the base tensor mapping; training-style attn. names are renamed to self_attn. so the patterns match
Assisted-by: GLM-5.3
* model : format hrm-text build_* calls as in other models
one argument group per line, matching sibling model files
Assisted-by: GLM-5.3
* llama : move hrm z_l_init table entries out of the nemotron group
place the name and tensor-info entries with the other global input tensors
Assisted-by: GLM-5.3
* convert : slim down hrm_text comments
Assisted-by: GLM-5.3
* convert : build hrm_text block tensor names from the {bid} template
The tensor map holds concrete per-block names, so format the template
with the computed layer index before handing it to super().
* llama : name hrm metadata keys in their own hrm. namespace
The four keys are arch-independent, unlike the arch-substituted
Keys.LLM entries, so group them under Keys.HRM (like Keys.Split) and
rename the llm_kv entries to LLM_KV_HRM_*. Only our own GGUFs carry
the old hrm_text.* keys; they are regenerated.
* Update src/llama-model-saver.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* llama : keep hrm metadata keys arch-substituted
Per review: the GGUF keys stay "{arch}.h_cycles" style, so the Python
members drop the LLM_KV_HRM_ prefix and keep arch templates; C++ keeps
the LLM_KV_HRM_* enums. GGUF output is unchanged - existing files and
HF uploads stay valid.
* Update gguf-py/gguf/constants.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* convert : rename hrm writer methods to add_hrm_*
Generic names like add_h_cycles/add_prefix_lm are too broad on the
shared GGUFWriter; prefix them with hrm_ like the metadata keys.
* model : fix meta-split lookup for archs with aliased cache slots
Cache tensors of archs that alias physical blocks across looped slots
(hrm_text, nanbeige with num_loops > 1) can reference block indices
without weight tensor names. Take the output projection from the layer
array instead of asserting; all other lookups are unchanged.
* model : replicate hrm_text tensors on meta devices instead of splitting
The aliased cache slots rotate split states differently from their
physical weights, so the meta-split execution invariants (set_rows
requires the cache state to match the token indices) cannot hold for
any device count. Replicate all hrm_text tensors on every meta device
instead; single-device and non-meta paths are unchanged.
Assisted-by: Claude Sonnet
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* test-backend-ops: reproduce MUL_MAT_ID NaN for activations beyond f16
The Metal mul_mm_id path narrows src1 to `half` for the simdgroup MMA
(`S1 = half` in every instantiation; ggml-metal.metal:10582 and :10595,
mirrored at :10643/:10654 in the tensor-ops path). f16 saturates at
65504, so a model whose activations exceed that produces inf, and
`simdgroup_multiply_accumulate` then turns the whole 8x8 accumulator
tile into NaN. The mul_mv_id path used below `ne21_mm_id_min` (32)
carries the same values in f32 and is correct, as is every CPU path.
This was untestable before: `init_mul_mat_id_tensors` initializes
uniform [-1, 1], so no existing case can drive an operand out of f16
range. `test_mul_mat_id` gains an `amax` parameter (default 1.0f,
preserving the historical init exactly) that scales only the f32
activations, leaving the quantized weights in their normal range.
Six cases: n=16 sits below the mul_mv_id -> mul_mm_id switch and is the
control that must stay green; n=32 and n=64 are above it and fail on
Metal today. Two shapes, because this is not model- or size-specific —
q4_K at 128 experts / 4 active / 4096x2048 mirrors a real model, and
q8_0 at 8 experts / 2 active / 512x256 shows the same failure at
minimal size.
Observed on Apple M2 Max, macOS, llama.cpp b10156:
MUL_MAT_ID(type_a=q8_0,...,n=32,k=256,amax=100000.000000):
[MUL_MAT_ID] NaN at index 0 (MTL0=nan CPU=583442.375000) FAIL
The real model behind this is Mistral Small 4 (arch mistral4, 128
experts / 4 active), one of whose layers reaches ~1e5 activations: on
Metal every prefill of >=32 tokens returns an entirely NaN vocabulary,
while <32 tokens is correct.
Note kernel_mul_mm (dense) has the identical conversion at :10273 and
:10286 and is expected to fail the same way; it is not covered here.
Found and written by Claude Opus 5 (via Claude Code).
* metal: fix NaN in mul_mm_id when activations exceed f16 range
kernel_mul_mm_id narrows src1 to `half` for the simdgroup MMA operands
(`S1 = half` in every instantiation). f16 saturates at 65504, so a model
whose activations exceed that produces inf on load, and
simdgroup_multiply_accumulate then propagates NaN across the whole 8x8
accumulator tile. The result is an entirely NaN output — not a precision
loss, a total loss. The mul_mv_id path taken below ne21_mm_id_min (32)
keeps the same values in f32 and is correct, as is every CPU path, so
the same model produces correct logits for short inputs and NaN for
long ones.
Fix: rescale src1 by a power of two so it fits, and undo the scale on
the f32 accumulator at the store. A two-stage reduction computes
max(|src1|) and writes the pair (1/scale, scale) into scratch chained
off the destination buffer, in the same style as the existing tpe/ids
id-mapping scratch. The matmul multiplies on load and on store.
This is exact, not approximate, for two reasons: the dot product is
linear, so one tensor-wide factor commutes through the accumulation;
and the factor is a power of two, so both multiplications are exact in
binary floating point. When max(|src1|) already fits — every model that
works today — the factor is exactly 1.0 and the output is bit-identical
to before. Accumulation was already f32 and is unchanged; only the
operand narrowing was ever the problem.
The reduction is two-stage (256 threadgroups into partials, then one
threadgroup folding them) specifically so it stays bandwidth-bound. A
single-threadgroup version was measured first and cost up to +451%
median on prefill — the scan serialized against an otherwise idle GPU.
It is also dispatched only on the mm path, so decode never pays for it.
Measured on Apple M2 Max, `test-backend-ops perf -o MUL_MAT_ID -b MTL0`,
99 cases, versus the same build without this change:
n=1/4/8 (mul_mv_id, decode) : -0.8% / -0.8% / -0.4% median (noise)
n=32 (mul_mm_id, prefill) : +1.73% median
n=64 : +1.30% median
n=128 : +1.80% median
n=256 : +3.98% median
n=512 : +3.74% median, +7.20% worst
overall : +1.14% median
Correctness, same machine:
- the six new test-backend-ops cases go from 4 FAIL / 2 OK to all OK,
with the n=16 controls (mul_mv_id path) unchanged;
- `test-backend-ops -b MTL0` full run: 0 failures, no regression;
- Mistral-Small-4-119B (arch mistral4, 128 experts / 4 active) now
generates correctly at the default n_ubatch of 512, in both
UD-IQ3_S and UD-Q4_K_XL quantizations. Before this, every prefill of
>= 32 tokens returned an all-NaN vocabulary and only n_ubatch <= 31
(forcing the mul_mv_id path) worked.
Likely fixes#25722 (mistral4 empty output on Metal above ~300 tokens,
FA on and off, generation degenerating to a single control token — the
signature of argmax over an all-NaN distribution). #20668 may be the
same defect attributed to a bad GGUF.
Note kernel_mul_mm (dense) has the identical narrowing at the
corresponding load sites and is expected to fail the same way; it is
left alone here to keep this change reviewable. Also possible, and left
for later: scaling per output column rather than per tensor, which
would preserve more precision when a single token is the hot one.
Found, diagnosed and fixed by Claude Opus 5 (via Claude Code).
* metal : make requested edits
- remove verbose comments
- explain rationale as requested
Generative AI disclosure: Claude made the edits as requested.
* metal : stack mul_mm_id map0 with amax_part
Implement @ggerganov suggestion to stack amax_part + map0. Mean 2.6% faster (worst -0.7%, best -4.1%). Win grows with batch size. Benchmarked on a hot M2 Max after reboot.
Generative AI disclosure:
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* cont : fix var scope
* cont : comment out tests temporarily
Comment out tess to not break CI temporarily
Assisted-by: Claude Fable 5.1
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* add self-hosted vulkan and webgpu to hf-jobs
* try t4-medium
* cont : adjust cpu backend threads
* try t4-small again
* restore cm jobs
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* cuda: support row-contiguous SUM_ROWS
* organize the code and add GGML_OP_MEAN to support row-contiguous tensors using the same shared kernel, and add a test to MEAN permute/slice
* Keep original comments and add if/else branch
* metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3)
MiniCPM3 sets attention.key_length to 96 and does not set
attention.value_length, which defaults to n_embd / n_head = 64. Metal had no
(96, 64) instantiation, so -fa auto aborted on the missing
kernel_flash_attn_ext_vec_f16_dk96_dv64.
Instantiate the tile kernel at (96, 64) for every K/V type that already has
(96, 96), and the vec kernel for the NE=4 configurations. Of the NE values the
vec dispatch considers, only NE=4 works here, because NL = 32/NE has to divide
both DK/4 = 24 and DV/4 = 16.
* tests : avoid redundant FA vec slice coverage
This commit updates cmake to use PROJECT_SOURCE_DIR instead of CMAKE_SOURCE_DIR for paths in function calls.
The motivation for this is that when using add_subdirectory,
CMAKE_SOURCE_DIR is fixed to the top-level projects source directory,
that is the caller of add_subdirectory and not the llama.cpp root
which means that common/common.h header will not be resolved.
Refs: https://github.com/ggml-org/llama.cpp/pull/28091#issuecomment-5636106377
This commit removes the precompiled headers that I added in Commit
3bcfeb700 ("cmake : add PCH and unity build to improve build times
(#28091)").
The motivation for this is that this looked good when developing this
but has caused multiple issues that I had taken into consideration and
we have decided to remove it and only keep the unity builds from the
above commit.
Refs: https://github.com/ggml-org/llama.cpp/pull/28882#issuecomment-5662272126
* tests : add README for updating the per-backend fusion baselines
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* ci : trigger fusion on changes to test-llama-archs.cpp and src/models
the dummy models and their architectures drive the fusion baselines, so a
change to either can alter the per-fusion counters and should re-run the
fusion job.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : merge the fusion build commands in the README
assisted-by: pi:llama.cpp/Qwen3.8-27B
* pi : require explicit permission before posting PR/issue comments
assisted-by: pi:llama.cpp/Qwen3.8-27B
* gguf-py: add Maple tensor constants
Add MODEL_ARCH.MAPLE, its "maple" name, and the tensor list for the
Maple 20B-A1B ternary MoE architecture: token embeddings, output,
attention with Q/K RMS norms, and per-expert FFN tensors.
* convert: add Maple HF->GGUF converter
Register MapleForCausalLM in the HF architecture map and add the
converter for the Maple 20B-A1B ternary MoE model: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, partial rotary factor 0.5, and
per-expert weight stacking into merged 3D tensors.
* llama: add Maple architecture (20B-A1B ternary MoE)
Add the Maple 20B-A1B ternary MoE architecture: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, and ternary TQ1_0/TQ2_0
quantization support.
- register LLM_ARCH_MAPLE between MAMBA2 and JAMBA
- implement llama_model_maple: Q/K RMS norms after projection (GEMMA4
style), rope applied only on SWA layers (nope_on_global_attention),
ISWA KV cache, and MoE FFN with swiglu gate clamp at +7 (DEEPSEEK4
style)
- mark MAPLE as unsupported by the model saver (roundtrip skipped)
* tests: mark Maple as MoE-mandatory
Maple is always-MoE: the model throws when n_expert == 0, so the
test harness must only run the MoE config for LLM_ARCH_MAPLE.
* maple: apply review feedback (n_ff_exp_arr, get_arr, rope params)
- load_arch_hparams: use n_ff_exp_arr + n_ff_exp() accessor (upstream
changed these from a scalar member during the rebase)
- sliding_window_pattern: get_arr, the pattern is mandatory for this arch
- partial_rotary_factor: read only from rope_parameters (base.py mirrors
the top-level key automatically)
- document why TOKEN_EMBD/OUTPUT are forced to F16 (they are the two
dense tensors in Maple, and the reference GGUFs ship them as F16)
- add @ModelBase.example("deepgrove/maple-preview")
* tests: add Maple to the SWA pattern array list
get_arr for maple.attention.sliding_window_pattern requires an array, but
the harness only emitted a per-layer array for the arches in its list, so
test-llama-archs -a maple failed to load the model.
Assisted-by: DeepSeek Harness
* maple: move swiglu_clamp_exp to the converter
The loader prefilled 7.0 and read the key optionally. The converter now
writes it and the loader reads it as required, because llama-graph.cpp
skips the clamp when the limit is 0 and an optional read would silently
run unclamped. The test harness provides the key for the same reason.
Also drops tensor_force_quant: base.py already forces FFN_GATE_INP to F32
and TOKEN_EMBD/OUTPUT to F16 for ternary file types.
Assisted-by: DeepSeek Harness
* convert: fix the LazyBase func signature in the Maple converter
ty flagged the stack() closure: it takes no argument, while LazyBase is
annotated with func: Callable[[Any], Any]. Pass the tensor list through
args instead of closing over it, the same way kimi_k3 does, so the
callable shape matches.
Assisted-by: DeepSeek Harness
* sycl: GPU-resident TOP_K for large k, parallelised over the device
The SYCL backend refused GGML_OP_TOP_K above k = 32 and let it fall back to
the CPU, a backend round-trip per call. The limit was not conservatism: the
scan-merge kernels keep (split_block + 1) * k candidate (value, index) pairs
in SLM, so at k = 128 a work-group already needs 132 KB and cannot launch.
qwen4exp's sparse-attention indexer asks for k = 2048 in 12 layers on every
token, so this fired at every context length.
Add a radix select for large k. The k-th largest is found by four
most-significant-first passes over an order-preserving unsigned key: histogram
the digit over the candidate set, walk the buckets from the top, and recurse
into the one where the running count reaches what is still needed. SLM holds
the histogram rather than candidates, so the footprint is independent of k.
A final pass emits every column beating the pivot plus exactly as many
pivot-equal columns as are still missing, so duplicate keys still yield
exactly k distinct indices. Output order is not required and is not paid for:
ggml-cpu/ops.cpp swaps its first two outputs to say so.
The key folds -0.0 onto +0.0 so its equivalence classes match the reference
comparator, under which the two tie. NaN has no defined order in the reference
(its comparator is not a strict weak order there); here +NaN keys above +inf
and -NaN below -inf, which at least makes the result deterministic.
One work-group per row leaves the device idle whenever a graph has fewer rows
than it has cores, which at batch size 1 means one work-group full stop:
qwen4exp tops-k a tensor of shape [n_kv, n_tokens/n_stream, n_stream], so
token generation gives nrows == 1, and the backend sampler reshapes logits to
a single row as well. Measured, ne=[200000,1] and ne=[200000,16] cost 358.0 us
and 363.4 us -- sixteen rows for 1.5% more wall-clock.
So also spread a row over several groups when there are too few rows to cover
the device. Per-pass state moves to global memory and each digit pass becomes
its own launch, since a work-group barrier can no longer span the row. Groups
accumulate in SLM and contribute 256 global atomics each, keeping global
traffic per-group rather than per-element, and the last group of a row -- the
one whose fetch_add returns G-1 -- performs that pass's scan, holding the
launch count at one per digit plus one emit. The group count comes from the
device and is floor-divided by nrows, so a row count that already covers the
device is left whole and pays nothing. Below 64K columns the single-group
kernel finishes inside the cost of the extra launches and stays in charge.
Reading the row's prefix/mask/need through a device-scope atomic_ref costs
more than the sweep it guards: those loads are uncached, so passes 2-4 ran at
49 us against 12 us for pass 1. One lane reads them into SLM and the group
takes them from there -- 208 us -> 44.6 us at ne=[131072,1], k=2048.
The block size now takes the device's max_work_group_size instead of a cap of
512. The cap was never a floor, so a device reporting 512 is unaffected; one
allowing 1024 was being given half its width.
Finally, put the scan-merge gate where the two paths actually cross. That
kernel's cost climbs with k while the radix select's does not; measured over
widths from 2 to 200K columns and row counts from 1 to 8192, radix is ahead
everywhere from k = 8 up and behind at k <= 2, where scan-merge's smaller
fixed cost wins. The short-row corner (ncols=2, nrows=65536, as in bailingmoe2
group selection) is exactly where radix loses at low k, and the gate keeps it
on scan-merge.
Op-level against the CPU-fallback path this replaces, and against the
single-group radix select for the split: 4.98x at ne=[131072,1] k=2048,
6.65x at ne=[151936,1] k=40, 13.35x at k=20, 118x at ne=[65000,16] k=32.
No measured shape regressed. End to end on 3x Arc Pro B60 with
Qwen3.8-Flash-Next UD-IQ4_XS, llama-bench tg64, the parallelisation is worth
5.91 -> 6.05 t/s at d=131072 and a wash at shallower depths. Perplexity over
wikitext-2 is unchanged within noise at both 512 and 81920 context.
test-backend-ops: 525/525 TOP_K (previously every k > 32 case was refused),
880/880 MUL_MAT_ID. Perf coverage added for k > 32 at large widths and for the
short-row corner, neither of which was exercised before.
* move topk-select to topk-radix.{cpp|hpp}
---------
Co-authored-by: cwriter <cwriter@localhost>
Corrects a typo in `tests/test-quant-type-selection` for the
Nvidia Nemotron 3 Nano 30B A3B model, which was referred to as
*nvidia-nemotron-nano-3-30b-a3b*.
The error made the test skip that test case, rather than failing
the test.
[no release]
kurquhar