* chat : honor json_schema in Ling 3.0 parser
Ling 3.0 only built a grammar for tool calls and did not handle inputs.json_schema, so response_format requests were left unconstrained.
Add an eager response-format grammar path with precedence over tools, following the existing parser patterns. Require </think> before JSON when thinking is enabled and do not allow trailing prose after the JSON response.
Fixes#29652.
Assisted-by: Claude Opus 5.5
* chat : require Ling 3.0 think block for response formats
build_rs gathered the extra states (n_rs - n_seqs rows) with their own
get_rows. The worst-case reserve has n_rs == n_seqs, so that node was
sized at zero rows, and any ubatch whose cells are not contiguous forced
a graph reallocation at an unchanged node count, which aborts under
GGML_SCHED_NO_REALLOC.
A single get_rows now gathers the n_rs states: the ubatch states and the
extra states are views of it, and its size only depends on n_rs, which
the reserve already sets to the maximum. A custom getter (mamba ssm_scan)
gathers from the second state, so a single sequence ubatch copies no
state. The views are built once per graph in the input to keep the host
overhead of the graph unchanged.
* ggml-openvino : Qwen3.5 MoE perf (#312)
Squash of ravi9/llama.cpp#312:
- ggml-openvino: add detailed inference profiling (Yu, Zijun)
- ggml-openvino: use remote output tensors by default (Yu, Zijun)
- ggml-openvino: optimize single-sequence recurrent state (Yu, Zijun)
- opt1: remove recurrent reset for single sequence, opt2: direct gdn outputs (break parallel sequence) (Yu, Zijun)
- fix parallel sequences (Yu, Zijun)
- ggml-openvino: simplify graph cache key (ynimmaga)
- enable stateful for qwen35 single sequence (Yu, Zijun)
- Fix after rebasing (Yu, Zijun)
- Add k-requant option q4_asym64 (Yu, Zijun)
- Fix qwen35 llama-bench -p 0 (Yu, Zijun)
- Simplify RESHAPE translation (Yu, Zijun)
- openvino: fuse MoE routing (Yu, Zijun)
- openvino: fuse GDN qk normalization (Yu, Zijun)
- openvino: enable GPU MoE fusion by default (Yu, Zijun)
- ggml-openvino: add cache_only mode to import cached compiled model on disk directly (Yu, Zijun)
- openvino : report the device allocation limit to ggml (Łukasz Ślusarczyk)
- Fix windows build (Yu, Zijun)
Co-authored-by: ynimmaga <ynimmaga@users.noreply.github.com>
Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
* ggml-openvino: Update doc of compiled model cache
* openvino: implement PRD-compliant device enumeration and memory reporting
* openvino: fix multi-device listing issues from review
- Only the device selected by GGML_OPENVINO_DEVICE reports as GPU; the
other OpenVINO devices report as IGPU so llama.cpp does not offload to
them. Initializing a non-selected device logs a warning.
- Name devices OPENVINO<i> again and show the OpenVINO id in the
description. Raw "CPU" names shadowed the ggml CPU backend.
- Support GPU.N: create the OpenCL queue on OpenVINO's own context for
the selected device, and replace "GPU"/"NPU" string comparisons with
ggml_openvino_is_gpu()/ggml_openvino_is_npu().
- An unavailable GGML_OPENVINO_DEVICE is now an error that lists the
available devices, instead of silently falling back to CPU.
- Memory: cap iGPU/NPU free memory at system available memory, fall back
to system memory instead of 0/0 when the plugin lacks memory
properties, and ignore host USM allocations in GPU usage.
- Initialize the device config once under a lock, even if OpenCL setup
fails.
- Fix supports_op return type for non-selected devices (build error).
* openvino : take USM entry points from the selected device platform
clGetExtensionFunctionAddressForPlatform was called on the first platform
returned by clGetPlatformIDs. The address it returns is only valid for the
platform it was queried on, and the first platform is not always the one that
holds the device OpenVINO selected.
On a host whose first platform comes from another vendor the lookup returns
null, and then every read, write and memset on a GPU buffer fails with
"clEnqueueMemcpyINTEL not available".
Look both entry points up in init(), on the platform of the device OpenVINO
picked, and keep them in the device config next to the command queue.
Assisted-by: Claude Opus 5
* openvino: fuse MoE experts for models with a fused gate_up weight
FuseMoeCompressed only matches models whose gate and up projections are
separate GatherMatmul ops. gemma-4 packs both into one expert weight and
splits the result after the GEMM, so its MoE block stayed unfused and ran
the expert GEMMs as per-token GEMVs.
Add FuseMoeCompressedFusedGateUp, which matches that shape
(one GatherMatmul -> Slice/Slice -> Gelu(ERF) -> Multiply) and folds it into
the same MOECompressed op, using GEMM3_SWIGLU with GEGLU_ERF. The fused
weight, scale and zero point are split into gate/up halves by copying raw
bytes, since a graph Slice would be rewritten to StridedSlice and constant
folded, whose reference evaluator crashes on sub-byte types.
gemma-4 also applies a per-expert output scale to the down projection before
the router weights. MOECompressed takes only one per-expert weight, so that
scale is folded into the routing weights, which is exact.
The op reads the zero point straight off a weight port and needs an integer
Constant there, so the matcher requires one and leaves natively quantized
experts (exact f16 zp) to the unfused path.
gemma-4-26B-A4B on Arc B390, GGML_OPENVINO_REQUANT_KQUANT=q4_asym64_all,
llama-bench -p 512 -n 128 -r 2, against a GGML_OPENVINO_MOE_OP=0 baseline:
pp512 66.16 -> 1608.73 t/s, tg128 25.94 -> 26.46 t/s. Perplexity over 12
chunks is unchanged (1451.3 +/- 177.9 unfused vs 1427.6 +/- 175.1 fused).
No effect without that requant option, on models with separate gate/up
weights, or on CPU. test-backend-ops -b OPENVINO0 is unchanged by this
commit: two MUL_MAT_ID m_v cases fail, the same two on the unmodified base.
* openvino: fix rank-3 axis handling so MoE works under stateful execution
Stateful execution drops the leading size-1 batch dim, so OV tensors are rank
3 while GgmlOvDecoder::get_shape/get_stride still report GGML_MAX_DIMS=4
reversed entries. Several MoE ops derive OV axis indices straight from that
metadata, so they picked the wrong axis. A MoE model with
GGML_OPENVINO_STATEFUL_EXECUTION=1 aborts while building the graph:
Check 'is_axis_valid(axis, r)' failed at src/core/src/validation_util.cpp:336
While validating node 'opset11::TopK ... _ffn_moe_probs ...'
Axis 3 out of the tensor rank range [-3, 2].
Fix idiom throughout: take the axis from the real OV rank, or shift a
metadata-derived axis down by metadata_rank - actual_rank.
argsort.cpp the router top-k axis is 2 on rank 3, not 3. This is the
abort quoted above.
add.cpp the MoE expert-sum bypass collapses the 8-ADD chain into one
ReduceSum on hardcoded axis 2, which on rank 3 reduces n_embd
instead of the expert axis. Now rank-2, with the following
Unsqueeze at rank-3.
get_rows.cpp squeezing a hardcoded {0,1} also strips the batch dim
whenever it is 1, which is every decode step. Squeeze down to
the trailing two dims instead.
mul_mat_id.cpp pick the reshape dims by actual rank, and skip the trailing
Unsqueeze that re-adds the batch dim.
view.cpp the expert-plane slice had the Slice axis, dst_ov_axis, the
ShapeOf+Gather index and the Reshape target all rank-4.
utils.cpp process_view_input_new's "translate_view already resolved
this VIEW, skip re-slicing" shortcut required equal ranks. 4
vs 3 never matched, so every resolved expert plane got
re-sliced. Now compares the common trailing dims. Same axis
shift for the Slice in the view-chain walker.
Stateless is unchanged by construction: every edit is gated on the actual
rank, so axis_shift == 0 reproduces the previous code exactly. Checked on
OV-CPU by diffing greedy output against the unmodified base for dense
gemma-4-E2B, granite-1b-a400m and gemma-4-26B-A4B; all identical.
granite-1b-a400m on OV-CPU aborts with the error above before this change;
after it, it generates and is byte-identical to stateless. Dense gemma-4-E2B
is identical stateless vs stateful both before and after. test-backend-ops
-b OPENVINO0 is unchanged: two pre-existing MUL_MAT_ID m_v cases fail, the
same two on the unmodified base.
gemma-4-26B-A4B is a poor correctness vehicle here. On OV it already drifts
into degenerate repetition a few tokens in, in stateless as much as stateful,
and the two modes diverge somewhere inside that degenerate region instead of
matching token for token. Each mode is self-reproducible across runs.
Known limitation: FuseMoeCompressedFusedGateUp does not match the rank-3
graph, so a MoE model run with GGML_OPENVINO_STATEFUL_EXECUTION=1 loses the
prefill fusion while gaining decode. gemma-4-26B-A4B on Arc B390,
GGML_OPENVINO_REQUANT_KQUANT=q4_asym64_all, llama-bench -p 512 -n 128 -r 2:
unfused (GGML_OPENVINO_MOE_OP=0) pp512 66.16 tg128 25.94
fused, stateless (default) pp512 1608.73 tg128 26.46
fused, stateful pp512 66.18 tg128 29.91
Stateful is opt-in and off by default, and MoE did not run there at all
before this, so nothing that previously worked regresses. Making the pass
match rank 3 is the follow-up.
* OpenVINO Backend: Upgrade graph cache to use node_idx, src_idx, node type
* ggml-openvino : enable more comprehensive conv fusion
* enable conv ops
* Reject kernel size 0 and support IM2COL_3D
* openvino : abort when the GPU remote context cannot be created
init() logged the error and returned, which left the device name a GPU but
remote_context empty. The remote buffer and tensor paths assert only on the
device being a GPU and then dereference that empty optional.
Those paths have no host fallback, and a device that OpenVINO listed should
have a working OpenCL context, so stop instead of continuing. An OpenCL stack
that is broken as a whole is still caught earlier by the device availability
check, which falls back to CPU.
Assisted-by: Claude Opus 5
* openvino : fix build warnings
The single-argument form of the OpenVINO RTTI macros is the intended one, but
their selector macro leaves __VA_ARGS__ empty, which -Wpedantic reports on
every pass and op header. Turn that warning off for this backend only, the
way ggml-cuda and ggml-sycl already do for their own third-party warnings.
Also drop a break and a dead assignment around a GGML_ABORT, which is noreturn.
Assisted-by: Claude Opus 5
* OpenVINO Backend: Support common MTMD ops
* ggml-openvino: give a reshaping view its own ov::Tensor
* ggml-openvino : compute HARDSIGMOID and EXPM1 in f32
HARDSIGMOID used a 1/6 constant in the input type, which is not exact
in bf16, and EXPM1 lost precision for small inputs in f16. Both now
compute in f32 and convert back, except on NPU where the f32 path
gives wrong results.
Fixes the HARDSIGMOID/EXPM1 test-backend-ops failures on GPU.
* ggml-openvino : update device selection and --list-devices
Show the selecting GGML_OPENVINO_DEVICE value and active device in
--list-devices, startup logs, and backend tests.
Clarify OpenVINO selection uses GGML_OPENVINO_DEVICE, not -dev.
* openvino : remove unreachable OpenCL queue checks
A remote buffer exists only on a GPU device, and init() aborts there if the
queue cannot be created, so the queue is never null at these call sites.
Assisted-by: Claude Opus 5
* openvino : update OpenVINO to 2026.4.1 and GPU drivers to 26.35.39758.10
* docs : update OpenVINO validated models and GPU driver version
* ggml-openvino : skip empty views when giving a reshaping view its own tensor
A zero-size view can sit at the end of a GPU USM buffer (Qwen3.5 recurrent cache). Wrapping it as a remote tensor throws "shared USM buffer has smaller size (0)".
Assisted-by: Claude
* ggml-openvino : rebind the cached decoder when llama passes a different graph
llama keeps separate graphs for batches with and without outputs. llama-server splits the prompt into chunks for context checkpoints, so a cached decoder could be reused with a graph built in other memory and bind the previous chunk's input tensors. SWA and recurrent models then lost most of the prompt in llama-cli and llama-server.
Assisted-by: Claude
* docs : update OpenVINO validated models
Smoke test on Lunar Lake (32 GB) with the two fixes above. Re-add the Qwen3.5 and gemma models.
Assisted-by: Claude
---------
Co-authored-by: Yu, Zijun <zijun.yu@intel.com>
Co-authored-by: ynimmaga <ynimmaga@users.noreply.github.com>
Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
Co-authored-by: haarika-madaka <haarika.madaka@intel.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
* qwen4exp : halve the indexer score memory
The indexer scored all heads in one product and rectified a copy of it,
so two [n_pool, n_idx_h, n_tokens] f32 tensors were live at once, the
largest buffers of the graph at long context. Each head now gets its
own product, rectified and summed in place into one [n_pool, n_tokens]
score.
* qwen4exp: let the allocator reuse the indexer score buffers
Address review from CISC: use plain ggml_add and ggml_relu in the
indexer head loop. The graph allocator already runs them in place when
their source has no other consumer, so the _inplace variants are not
needed. The compute buffer and the speed are unchanged.
* cuda: support 4 heads in the lightning indexer
Dispatch 4 heads to the vector kernel, too few for a wmma tile, and
accept them in supports_op. test-backend-ops covers 4 heads.
* metal: take the lightning indexer head count as a function constant
The kernel reads the head count from a function constant and zero fills
the last head tile, so any head count runs and 64 heads is unchanged.
* qwen4exp: compute the indexer score with the lightning indexer
Address review from am17an: the unweighted sum of the rectified head
scores scaled by 1/sqrt(head_dim) is the lightning indexer with every
head weight set to that scale, so the indexer calls
ggml_lightning_indexer on the pooled keys with an f16 pool mask. The
keys are read once for all heads and no per head score is
materialized.
* vulkan: tile the lightning indexer over keys and tokens
A workgroup scores 64 keys against 8 tokens: the keys are staged once
in shared memory, the queries one head at a time, and each invocation
owns one key for two tokens, so no dot product needs a cross invocation
reduction. The subgroup variant and the flat dispatch are gone, the grid
is keys x tokens x streams.
* vectorize vulkan loads and use fp16 dot product
---------
Co-authored-by: Ruben Ortlam <rortlam@redhat.com>
* Make the drafter probabilistic and the target verify by rejection sampling
* Drop stale spec_draft_q before drafting
* Fallback to argmax sampling for grammar-constrained requests and adding flag for enabling probabilistic draft sampling. Default flag value is greedy.
* Support grammar-constrained requests in rejection sampling
* Fix - renormalize distribution after masking
* copy rng on sampler copy and re-accept drafted tokens on replay
* Fix draft sampler sharing the target's rng stream
* Simplify the rejection sampler's inputs and move replay to the server
* Truncate the draft candidates along with the draft
---------
Co-authored-by: praneshgo <227579474+praneshgo@users.noreply.github.com>
Co-authored-by: Pranesh Gonegandla <pgonegandla@nvidia.com>
* ggml-quants : avoid invalid rounding in qkx3 scale search
The imatrix scale search can produce an infinite, NaN, or otherwise out-of-range value when the fitted minimum collapses to the maximum or makes the range extremely small. That value is then passed to nearest_int and can trip its assertion in Debug builds.
Clamp the quantization level to [0, nmax] before rounding so valid in-range values behave the same as before while invalid scale-search results no longer reach nearest_int.
Add regression coverage for degenerate imatrix groups across q2_K, q4_K, q5_K, q4_1, and q5_1.
Fixes#29804.
Assisted-by: Claude Opus 5.5
* tests: print degenerate imatrix quant types
* ggml-cpu : fix soft_max_back wrong output when dst aliases src1
GGML_OP_SOFT_MAX_BACK is listed in ggml_op_can_inplace, so the graph
allocator may assign dst to alias either src0 (dy) or src1 (y).
The result was built in several steps:
ggml_vec_cpy_f32 (nc, dx, dy);
ggml_vec_acc1_f32 (nc, dx, -dot_y_dy);
ggml_vec_mul_f32 (nc, dx, dx, y);
ggml_vec_scale_f32(nc, dx, scale);
When dst aliases src1, the first step overwrites y and the third step
then reads the overwritten values, so the output is silently wrong.
Aliasing dst with src0 is unaffected. The CUDA kernel completes its
reduction before writing and is already safe.
Replace the sequence with a single fused loop that reads both sources
before writing, which is correct under either aliasing.
Add a regression test that marks dy as a graph output so the allocator
is forced to alias dst with y, asserts that the alias actually
happened, and compares against values computed on the host.
* cont : remove comment
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* metal : add tensor API flash attention kernel for F16 KV
* cont : add tensor FA kernels for DK=DV=512 and DK=576, DV=512
* cont : support attention sinks, ALiBi and logit softcap in the tensor FA kernel
* cont : add tensor FA kernel for DK=192, DV=128
* init conversion
* convert: ok
* model loaded
* add server code
* improve conversion script
* support shared prompt prefix
* add docs, imorove UX a bit
* add vision support
* add openjev tiny model for testing
* add dev docs
* support lev & kev
* clean up
* fix lev noul
* fix py lint
* nits docs
* clarify about not supporting date_facts
* Adding wide-load mmvq for Q8_0 and esimd dmmv for q8_0
Assisted-by: Codex
* remove guard for q8_0
* remove docs
* Simplify by committing to clean code without fallback
* Add feature flag as requested
Assisted-by: Claude Opus 5
---------
Co-authored-by: cwriter <cwriter@localhost>
* ggml : add `alloc_buffer_n` to buffer type interface
Add alloc_buffer_n method to ggml_backend_buffer_type_i
interface, with a public API ggml_backend_buft_alloc_buffer_n.
- Default implementation in ggml-backend.cpp handles multi-buffer
splitting and tensor allocation via ggml_tallocr
- Meta buffer type provides custom implementation that creates
per-device sub-contexts and delegates to simple buffer types
- ggml_backend_alloc_ctx_tensors_from_buft now collects tensors
into a list and delegates to the new API
- Remove temporary ggml_backend_meta_alloc_ctx_tensors_from_buft
- Add NULL alloc_buffer_n to all existing buffer type
interfaces (cpu, metal, openvino, hexagon, webgpu, zdnn, virtgpu, repack)
Assisted-by: llama.cpp:local pi
* cont : fix `cur_buf_size` init after flushing a buffer
* ggml : add TODO tag for shared buffer split logic
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* tests : add alloc_buffer_n coverage
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* cont : fix compile warnings
* tests : add descriptions for alloc_buffer_n tests
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* ggml : address review comments on alloc_buffer_n
- restore GGML_LOG_ERROR on buffer alloc / tensor init failure in the
default impl (name the failing tensor)
- check the malloc result and drop the _impl indirection in
ggml_backend_alloc_ctx_tensors_from_buft
- remove comments that restate the code
- fix the TAG_ALLOC_SHARED_BUFFER_SPLIT typo
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* ggml : add get_alloc_size_n to buffer type interface
- Add ggml_backend_buft_get_alloc_size_n public API
- Add optional get_alloc_size_n callback to ggml_backend_buffer_type_i
- Share tensor->buffer planning between alloc_buffer_n default and get_alloc_size_n default
- Replace unchecked realloc with std::vector in alloc_buffer_n default
- Make ggml_backend_alloc_ctx_tensors_from_buft_size use the new API
- Add test-alloc coverage for get_alloc_size_n
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* cont : report malloc failure
In tool.uv.sources, torch was unconditionally pinned to the custom
pytorch CPU index, which lacks macOS Darwin wheels and causes uv sync
to fail on macOS. Add the sys_platform == 'linux' marker to match the
existing Poetry dependencies configuration.
Assisted-by: Antigravity
Resolves: https://github.com/ggml-org/llama.cpp/issues/29176
* hexagon: add q2_k and q3_k quant type support
* hex-qk: consistent allocation of src1_row_size
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* tests : simplify function signature
* llama : clamp kpool re-pool bound to existing pools
The n_tokens/kpool + n_seqs_unq bound on n_new_g overshoots when a batch
fills the whole cache: n_ctx tokens complete exactly n_ctx/kpool pools, so
the +1 pads new_pool_idxs/new_pool_rep one entry past n_pool_real. Graph
reserve only covers n_pool_real entries, so the first full-context decode
builds bigger tensors than reserved and ggml-alloc demands a graph
reallocation (abort under GGML_SCHED_DEBUG_REALLOC=1).
Clamp the bound to n_pool_real: a ubatch can never mark more pools than
the cache holds, and reserve's n_pool_max already covers that.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-MOPD
* cont : cap to n_pool_max
* hexagon: shared strided DMA copy for CPY and CONCAT, any-dim CONCAT via DMA
* hex-cpy: various fixes on top of the concat optimizations
Removed CONCAT_DMA_MIN_ROW logic, it was broken with 64-bit DMA.
While it's kinda silly to use DMA for tiny stuff if that tensor gets mapped to an extended buffer the only way to read it is DMA.
Added missing dma_queue_flush() calls.
Added additional guards for conditions we don't support.
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* convert : write Gemma embedding scale for DFlash drafts
A DFlash draft shares the target's token embeddings. Gemma scales them by sqrt(hidden_size) in the forward pass, and the draft config does not state that scale, so the converted draft read unscaled embeddings.
Take the scale from the target config when the draft config has none.
Assisted-by: Claude
* convert : check with get_model_architecture for gemma models
* cuda : route sm70 to the Turing MMVQ nwarps table
Volta (sm_70) has no MMVQ parameter table of its own and falls through
to GENERIC, which launches K-quant batch-1 decode (ncols_dst == 1) at
nwarps=4. sm_70 shares TURING's tuning: the K-quant vec_dot prefers
nwarps=2 there. Route sm_70 to the existing MMVQ_PARAMETERS_TURING
table in both the device and the host table selector.
Measured on one Tesla V100 32GB PCIe (PG500-216, driver 580.178.04,
CUDA 12.0.140) with Qwen3.8-27B Q4_K_M, tg128, interleaved A/B in 6
ABBA blocks with paired per-block deltas: +1.091 t/s = +3.17 %
(t = +49.0, all six per-block deltas positive); perplexity
bit-identical (6.3697 +/- 0.04066 both builds, wiki.test.raw). The
patched build's K-quant mul_mat_vec_q kernels launch at nwarps=2
(cubin EIATTR_MAX_THREADS) while Q4_0/Q8_0 stay at nwarps=4, and the
same measurement on the September master base gave +3.84 % (t = 85).
The tuning originates from the V100-focused fork anyei/llamacpp-v100
(MIT), commit b912d1b1e, which carries a dedicated
MMVQ_PARAMETERS_VOLTA table; a cubin-level comparison confirmed that
routing sm_70 to the existing TURING table is equivalent for the
K-quant batch-1 path this change affects, so this is the minimal
2-line form. https://github.com/anyei/llamacpp-v100/commit/b912d1b1e
Original-patch-by: anyei <angelyoelroblesmercedes@gmail.com>
* Update ggml/src/ggml-cuda/mmvq.cu
---------
Co-authored-by: tkittich <tkittich@gmail.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* metal : release temporary private transfer buffers
Assisted-by: OpenAI Codex
* metal : fix order and formatting
---------
Co-authored-by: Niklas Wenzel <dev@nikwen.de>
* CUDA: Handle compute type for NVFP4 on cublass path
Signed-off-by: ynankani <ynankani@nvidia.com>
* Use BF16 compute type for quantized models if HW allows
Signed-off-by: ynankani <ynankani@nvidia.com>
* Set acc prec to bf16 for nvfp4 as it needs atleast bf16 range
Signed-off-by: ynankani <ynankani@nvidia.com>
* Update ggml/src/ggml-cuda/ggml-cuda.cu
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* preserve op_params for per-expert matmul
Signed-off-by: ynankani <ynankani@nvidia.com>
---------
Signed-off-by: ynankani <ynankani@nvidia.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
Pinning to >= 3.4.3 is required to enable DeviceTopK, which was affected by
a race condition https://github.com/NVIDIA/cccl/pull/10627.
We will relax this for future CTK versions which will bundle CCCL >
3.4.X (CTK 13.5 will bundle CCCL 3.5.0 for example)
* BLAS : Document AOCL-BLAS build and label the device AOCL-BLAS
* AOCL-Blas : Add an AOCL-BLAS Quick Start and drop the fixed version path
* AOCL-BLAS doc : Note on ZenDNN
LLM-jp-4.1 uses the GPT-OSS format, but its tokenizer decodes a space
after every special token and parallel tool calls are separated by
<|end|>. The GPT-OSS handler rejects this output, so add a dedicated
handler, selected by the chat_format=llm-jp-harmony-v1 declaration in
the chat template.
Assisted-by: Claude Fable 5.1
* vocab : honor BOS/EOS settings for PLaMo-2 and PLaMo-3
The original tokenizer configs for PLaMo-2 and PLaMo-3 have
`add_bos_token: true` and `add_eos_token: false` , but
_set_vocab_plamo() did not write the BOS/EOS metadata. The
PLAMO2 tokenizer path also ignored add_bos/add_eos during
tokenization.
Write the settings from tokenizer_config.json and honor them in
the PLAMO2 tokenization path. GGUFs without these keys keep the
previous behavior.
* Update conversion/base.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
The committed docs/ops/CPU.csv is out of sync with the current
test-backend-ops suite: 11 ops with CPU support (COL2IM_1D,
MUL_MAT_HADAMARD, SWIGLU_CLAMP, MUL_MAT_W4A4/W4A8, MUL_MAT_ID_W4A4/W4A8,
DSV4_HC_COMB/PRE/POST, LIGHTNING_INDEXER) are missing entirely, and
many other ops have fewer test cases than the suite generates now.
docs/ops.md (which CI requires to match the CSVs) therefore
understates CPU support.
Regenerated with:
test-backend-ops support -b CPU --output csv > docs/ops/CPU.csv
scripts/create_ops_docs.py
Note: ADD1 now reads unsupported on CPU because ggml_add1 is
GGML_DEPRECATED and the suite no longer generates test cases for it;
the CPU implementation itself is still present.
Assisted-by: Xing
#27941 disabled -sm tensor for qwen4exp because test-llama-archs asserted on the
Meta device once the fixture carried a PLE layer:
GGML_ASSERT(ggml_backend_buffer_is_meta(tensor->buffer)) at ggml-backend-meta.cpp:476.
With host-resident embeddings the PLE gather is a CPU node and hc_init (the REPEAT
that fans the embedding out to the hc streams) was first reached through layer 0's
PLE path, after that gather. ggml_backend_sched_split_graph pass 2 expands a device
assignment upwards only until it meets a CPU node, so the REPEAT stayed on the CPU
and the later reshape of hc_init inside the meta split viewed a host-resident node.
Expanding hc_init right after it is built puts the REPEAT directly before the first
device node, where pass 2 assigns it; the embedding reshape stays in the CPU split
and is copied in as a split input, as in deepseek4.
dequantize_block_iq4_nl writes QK_K values per block, but a row can be shorter than that (an IQ4_NL row is only guaranteed to be a multiple of QK4_NL). Threads whose 32-value sub-block starts at or past k currently read and write past the end of the row. Skip those sub-blocks; for rows that are a multiple of QK_K the check never fires.
* hex-allreduce: add support for safe scatter mode
* hex-allreduce: pare down excessive comments
* hex-allreduce: re-write to remove register spills
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* llama: preserve original batch order for layer inputs
Assisted-by: Codex
* tests: cover layer-input order across KV layouts
Assisted-by: Codex
* tests: exercise layer-input ordering on CUDA devices
Assisted-by: Codex
* llama: make layer input reordering compatible with tensor split
Copy each microbatch tensor from offset zero and restore original row order after synchronization. Extend the layer-input regression to cover tensor split and repeated reads and decodes.
Assisted-by: Codex
* llama: restore token order for unmasked NextN embeddings
Use the original-token mapping for unmasked NextN rows, including when
layer-input capture is disabled. Keep masked NextN rows on the logits
output mapping and preserve offset-zero tensor copies.
Extend the existing regression to cover NextN alone, combined layer
capture, and masked outputs with repeated decodes and getters.
Validation: all 256 CPU/CUDA/tensor configurations pass. Qwen3.8-27B
Q4_K_M MTP completes MT-Bench at concurrency 16 before and after.
Assisted-by: Codex
* ggml: fix WebGPU reservation and OpenVINO hidden-state capture
Reserve WebGPU vector attention scratch across batch sizes and refresh reservations when NextN capture settings change. Preserve requested OpenVINO outputs, dynamic shapes, sequence counts, and current graph bindings.
Extend existing WebGPU regression coverage and enable strict allocation checks.
Assisted-by: Codex
* llama: defer regression test and backend fixes to follow-ups
Keep this PR focused on restoring token order for layer inputs and unmasked NextN embeddings. Remove the added regression test, OpenVINO and WebGPU changes, and the separate NextN reservation change.
Assisted-by: Codex
* llama: keep n_embd declaration in its original position
Assisted-by: Codex
* llama : pass token count to layer input extraction
Assisted-by: Codex
* llama : name original batch indices batch_idxs
Assisted-by: Codex
* llama : name extracted embedding indices embd_batch_idxs
Assisted-by: Codex
* llama : tag target embedding reordering
Assisted-by: Codex
* llama : tag extraction and name the index capture flag
Assisted-by: Codex
After the device decode, flip causal_attn off, decode n_ubatch/2 then
n_ubatch tokens. Both have the same node count, so a shape that depends
on the flag makes the second reallocate at an unchanged graph size,
which aborts under GGML_SCHED_NO_REALLOC. Skipped for the encode archs.
* convert: fix LoRA conversion crash for Qwen3.5 V-head reorder
_reorder_v_heads does reshape+permute+reshape to reorder V heads from
grouped to tiled order. LoraTorchTensor.reshape() cannot split its
row dimension (A matrix), so converting Qwen3.5 LoRA adapters that
target out_proj crashes with NotImplementedError.
Fix: detect LoRA tensors and apply the equivalent index permutation
directly — column reorder (dim=last) permutes A's columns, row
reorder (dim=0) permutes B's rows. This is mathematically identical:
(B @ A)[:, perm] == B @ A[:, perm]
(B @ A)[perm, :] == B[perm, :] @ A
Verified: both paths produce exactly zero diff against the full-tensor
reorder on random (rank=32, 4096×4096) matrices.
Fixes#21125
Signed-off-by: Radu Swigler <radu@swigler.com>
* convert: add ty: ignore for hasattr-guarded LoRA call
Assisted-By: Claude Opus 4.6 <noreply@anthropic.com>
* fix comment
* nowrap
---------
Signed-off-by: Radu Swigler <radu@swigler.com>
Co-authored-by: Radu Swigler <radu@swigler.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Assisted-by: Claude Opus 4.6 <noreply@anthropic.com>
The sparse indexer mask is built with a set_rows scatter. Padded pools,
absent sequences and missing tail cells all pointed to the same n_kv
sentinel row, and invisible pools picked by top_k to fill the selection
overlap the tail cells of the token, so several CPU threads wrote the
same element (ThreadSanitizer data race in the sanitize CI).
Allocate the slot mask for both selection paths and route every dead
slot to its own dump row n_kv + slot. Live slots address disjoint cells,
so the scatter indices of a token are unique.
* ggml: fix integer overflow guard for zero-element tensors
* ggml: validate number of elements in tensor to prevent integer overflow
* ggml: fix error print
* cli: exit on stdin EOF and drop the console wide Ctrl+C broadcast
On Windows the simple input reader sends CTRL_C_EVENT to every process
attached to the console when stdin reaches EOF, killing unrelated
processes such as a supervising agent. The CLI only stopped on EOF
because of that self inflicted SIGINT; on POSIX, and with the advanced
reader, it spins forever printing prompts.
Drop the broadcast so both platforms just return an empty read, and
treat an empty read as EOF in the chat loop and the model selection,
since a submitted line always ends with a newline.
* cli: keep the newline of a trailing "/" and stop mtmd-cli on EOF
A lone "/" came back as an empty read and was taken for EOF, and
mtmd-cli only stopped on EOF through the removed broadcast.
The fixture recycles its two blocks over 8 cache slots, so the fp16
error builds up past the 1e-4 NMSE bound on the Vulkan T4 and WebGPU
jobs of Models Backend. Two l-cycles keep every branch of the cycle
loop and halve the error.
* llama: llama_prefetch_rows
* llama: support row prefetch on Windows
Apply the Windows port contributed by @praneshgo unchanged.
Source: https://github.com/ggml-org/llama.cpp/pull/29599#issuecomment-5887721014
* avoid exposing llama-mmap in model code, route via llama-impl
* add windows check, only prefetch in lazy mode
* cont : clean-up
* cont : fix build
* cont : clarify padding token for gemma4
---------
Co-authored-by: Pranesh Gonegandla <pranesh.iitp@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* ggml : add BF16 unary, GLU, binary and scale ops (CPU, CUDA)
* ggml-cpu : use per-op _bf16 functions for BF16 unary and GLU ops
Assisted-by: Claude Opus 5.5
* CUDA: use ggml_cuda_cast in binbcast and unary kernels to fix the HIP bf16 build
* ggml-openvino : reject BF16 SCALE and mixed-type BF16 ADD/MUL/SUB
* cpu: accept BF16 in src1 of mul_mat
ggml_conv_1d_dw builds its im2col in F32 when the kernel is BF16, then
calls ggml_mul_mat(im2col, kernel), which puts F32 in src0 and BF16 in
src1. The CPU backend refused that combination, so it was reported as
unsupported on every backend and never compared against anything.
Widen BF16 into the F32 work buffer, next to the existing packing of F32
into vec_dot_type. This is the arithmetic the Metal mat vec kernel
already uses, both operands promoted to float and accumulated in float,
so the two agree exactly rather than approximately.
Cover it with a conv_1d_dw test over F32, F16 and BF16 kernels, plus
three mul_mat cases with BF16 in src1.
* vulkan: reject BF16 in src1 of mul_mat unless src0 is BF16
supports_op only checked the src1 type for non contiguous tensors, so
a contiguous BF16 src1 was accepted and the pipeline lookup asserted.
The only BF16 src1 path is the BF16 x BF16 multiply, every other src0
type now reports the op as unsupported and the scheduler keeps it on
the CPU.
The BF16 kernel case of the conv_1d_dw test needs the f32 x bf16
mat vec variants of the Metal backend, which land separately.
This commit adds an optional --add-bos token command line option to the
run-org-model.py script.
The motivation for this is that there are models, for example Gemma4,
that explicitely set the add_bos value to true in llama-vocab.cpp even
if the original model does not set this value to True.
It would be nice to be able to force the models to agree on the bos
token so that logit verification can proceed.
Refs: https://github.com/ggml-org/llama.cpp/pull/21500
* ui : shared model display primitives
Extract ModelCapabilityIcons (canonical Tools/Reasoning/Vision/Video/Audio
order) out of ModelId and reuse it there, add the shared DialogConfirmDownload
for destructive download actions, the discover org avatar with dark-mode
inversion and the thin download progress bar, and rework ModelId badges to take
thinking/tool-use support directly.
Assisted-by: pi:GLM-5.3-Flash
* ui : remember hub avatars that failed to load
Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash
* ui : render shared model row hints as native titles
Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash
* ui : fix badge guard for draft sidecars, keep parameter precision
hasBadges now counts draft sidecar badges, so a sidecar-only model still
renders. Billions keep one decimal for hub counts and stay bare for whole
values. Avatar failures track the org instead of the instance, and the
download progress bar no longer pulses while determinate.
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : model download pipeline
Track HuggingFace downloads end to end: the server download/cancel endpoints,
a status manager fed by the /models/sse download progress events, and a
models-discover store holding the catalog and detail state for the discover
view. Downloaded and in-flight entries are excluded from the loadable model
list.
Assisted-by: pi:GLM-5.3-Flash
* ui : route sidecar tag lookup through the sidecars util, validate the paused list
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : model memory-fit estimation
Replace the raw runtime-memory estimate with the app's compatibility check:
the smallest real Mac memory tier that fits a model file, budgeted as
RAM x 0.75 minus fixed overhead with headroom on the file size. The constants
move to lib; the unused runtime-memory estimate is dropped. browser-info's
OS detection is exported for reuse.
Assisted-by: pi:GLM-5.3-Flash
* ui : cover the memory-fit and tool-use heuristics in tests
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : Hugging Face Hub data layer
Add HuggingFaceService and its constants/enums/types: GGUF repo search, file
tree and model detail fetching, quant/sidecar filename analysis, shard-set
collapsing and the llama.app catalog feed, plus an orgOf() helper on the model
name utils.
Assisted-by: pi:GLM-5.3-Flash
* ui : strip provider tilde prefix from hub avatar urls
Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash
* ui : trim redundant comments in the HF data layer service
Per review: drop JSDoc that restates the method name and inline comments
that restate the code; keep only comments carrying non-obvious context.
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : harden the HF data layer error typing, cover the helpers in tests
Carries the HTTP status on retryable fetch errors instead of matching the
message text. Marks expand-dependent catalog fields optional and documents
the data/models index pairing. Adds table tests for the pure helpers.
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : model id grammar for sidecars, quants and capability parsing
Extend the shared model id parser with sidecar tokens (draft variants and
auxiliary imatrix/mmproj files), weight-file and custom-quant regexes, and add
the tools capability to ModelCapabilities; the selector option row picks it up
from the model's declared capabilities.
Assisted-by: pi:GLM-5.3-Flash
* ui : escape sidecar tokens in the regex alternation
Assisted-by: pi:zai-org/GLM-5.3-Flash
* ui : type-safe API types, fetch helpers and download-ready models store plumbing
Assisted-by: pi:GLM-5.3-Flash
* ui : document the model list index pairing, fix an em-dash
Assisted-by: pi:zai-org/GLM-5.3-Flash
* Update tools/ui/src/lib/components/app/chat/index.ts
Co-authored-by: Pascal <admin@serveurperso.com>
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* openvino: serve GET_ROWS on a weight view from the base Constant
Resolve view_src when collecting weight Constants so a view over a
quantized weight no longer becomes a dynamic typed Parameter, and fold
the row offset of the view into the gather indices instead of slicing
the dequantization subgraph.
* openvino: lift the quantized GET_ROWS view rejection
The supports_op rejection of a quantized src0 view with a nonzero
offset keeps the vs0 GET_ROWS cases of #28253 away from OpenVINO.
The weight view now resolves to the base Constant with the row offset
folded into the gather indices, so the rejection goes away.
#27773 adds the glm5-next arch without its rows in the Metal fusion
baseline, so test-fusion --check fails on it. The rows come from
test-fusion --record on an M5 Max, and --check passes 270/270.
The MUSA vendor header never defined __CUDA_ARCH__, so every architecture
test in the shared ggml-cuda sources evaluated to 0. Kernel bodies gated on
the architecture therefore compiled to nothing, for example the q8_0 -> f16
dequantization kernel in convert.cu, whose NO_DEVICE_CODE fallback expands to
an empty body in host code.
Report the newest architecture like the HIP backend does and exclude the
NVIDIA-only features explicitly, as they are not usable on MUSA. Define it
for device passes only: CUB uses defined(__CUDA_ARCH__) to detect device
compilation, which is also how nvcc behaves.
Drop the now-redundant defined(__CUDA_ARCH__) checks in the architecture
comparisons: __CUDA_ARCH__ is undefined in host passes for CUDA and MUSA, and
HIP defines it for every pass, so both forms select the same branch.
* Rebase GLM-Next support onto master, and migrate to llama-memory-hybrid-idx
* Add initial MTP support
* Merge branch optimizations. Reduce allocated compute buffer size, speed up long context decode, fla, and slight MTP improvements.
* Review driven changes, remove env vars, protect tensors
* Strip MTP for initial PR
* Clean up after mtp strip
* Clean up after mtp strip
* Update speculative.cpp
* Update llama-context.h
* Clean up after mtp strip
* Fix tokenizer ignore merges
* Improve quantization protection selection
* Refactor mhc helpers, graph base
* Lint Fixes
* Apply suggestions from code review
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Skip glm5-next in model saver, fix CRLF
* Skip glm5-next in sweep
* Remove T4 fallback
* Review cleanup
* Review suggestions
* Defer separate MTP gguf handling to MTP PR, drop filter
* Repad n_head_kv
* kpool init apply
* Order by descending score
* Drop guard
* read kpool from hparams, clarify kpool cache flags, remove kpool_build_state(nullptr)
* Add glm5-next support to model saver and add arch test fixture
* Review cleanup
* Kpool pooled caching clarify
* Add multi stream support
* Finish Rebase
* Sparse FA fir DSA prefill
* Const
* Update llama-model.cpp to fix rebase error
* gguf-py : merge tensor map entries for HC tensors
* model : use build_gdn_l2_norm in GLM5_NEXT implementation
* chore : remove trailing whitespace
* model : use new OP precision setting API in GLM5_NEXT implementation
* mtmd : use ggml_swiglu_clamp in GLM5V and apply the image token limit
The two clamps around swiglu_split are what ggml_swiglu_clamp already does,
so the clamp bounds collapse back to one value. GLM5V also never called
set_limit_image_tokens(), so --image-max-tokens had no effect.
Assisted-by: Claude Opus 5
(cherry picked from commit 46d18e12d422be4cc04a70e4a9a9e0168bb3d5b7)
* llama : keep the GLM5-Next k-pool layout across ubatches
The layout was rebuilt from a full cell scan on every ubatch. Pools are fixed
by the positions relative to the sequence's first one, so the layout now lives
on the memory and a ubatch only appends to it.
A sequence edit no longer stales every pooled key either, only the ones at or
after the edited position, which makes a tail seq_rm free. The pooling subgraph
is built unconditionally so the graph shape no longer changes every kpool
tokens, and the pool axis is folded into rows before soft_max, which otherwise
exceeds the CUDA gridDim.y limit past n_kv 262144.
Assisted-by: Claude Opus 5
(cherry picked from commit 5d1c40b93e17fddbf73b785efe43e0d02ccb3977)
* model : write the GLM5-Next recurrent rollback checkpoints
The conv state and the delta net state were only written to the live row, so a
rollback restored whatever the checkpoint rows happened to hold. Take the same
route as kimi-k3: build_recurrent_attn for the state, and write all K_rs conv
groups. That also drops a state view that assumed contiguous rows.
Enroll the arch in test-recurrent-state-rollback, which catches this under its
garbage-filled cache pass.
Assisted-by: Claude Opus 5
(cherry picked from commit 5ace37e86d5d448e83ef5dde5632c748185b18cd)
* llama: fix PR #27773 test-save-load-state restore failure
Clear the attention and indexer cache data after a failed hybrid state restore so restored NaNs cannot affect a later sequence.
Assisted-by: Codex
* llama: fix PR #27773 gpu-rocm graph reallocation
Reserve the full GLM5-Next pool capacity and dirty pool count. The gpu-rocm Test step aborts when n_new grows while the graph node count stays fixed; CUDA, Vulkan, Metal, and WebGPU checks report the same error.
Assisted-by: Codex
* llama : fix GLM5-Next k-pool layout staleness after edits and shared teardown
Two defects in the cross-ubatch k-pool layout added by the k-pool commit:
1. Wrong results. An edited sequence only rebuilt its pool layout when its cell
count changed, so if the first ubatch after an edit added back exactly as many
cells as were removed, the stale position-to-cell list survived. With a unified
cache and more than one sequence, where another sequence takes the freed cells,
the reused layout points at the wrong cells (CPU: large logit drift, CUDA: NaN).
Rebuild whenever the sequence is stale, not only on a size mismatch.
2. Slowdown. "shared" mode was assumed to end only with an edit that forces a
rebuild, but sharing also ends when the other sequence is removed. The survivor
kept shared = true, pinning cache_safe off and re-pooling every pool on every
ubatch (server trigger: n>1 completions with -kvu, via the seq_cp in
copy_state_to). In seq_rm, if the layout has shared cells, stale every sequence
so one rebuild re-derives sharing and cache_safe returns to 1.
Assisted-by: Claude Opus 5
* llama : fix build_attn_mha stream stride for non-contiguous q
build_attn_mha split the batch into streams with a stream stride of
q->nb[3]/n_stream. That only equals one stream's span, (ne[2]/n_stream)*nb[2],
when q is contiguous. GLM5-Next is nope-only, so it does not concat a rope part
and passes the permuted q_absorbed straight in, where nb[3] != ne[2]*nb[2]; the
stride was then n_head times too large and every stream s >= 1 read another
head's queries. Split-KV (-np N without --kv-unified) multi-stream prefill was
wrong for every stream past the first. Unified KV and decode were unaffected
(n_stream == 1, and decode takes the gather path). Other MLA models concat rope
so q is contiguous and the computed value is unchanged for them.
Compute the stride from the token dimension, which is identical for a
contiguous q.
Assisted-by: Claude Opus 5
* llama : re-derive GLM5-Next k-pool sharing on state_read/state_drop
The shared-cell teardown added to seq_rm (stale every sequence when the layout
has shared cells, so a survivor does not keep shared = true and pin cache_safe
off) was missing from the other paths that can free shared cells: state_read
and state_drop staled only the one sequence. Apply the same re-derivation there
and correct the comment that claimed sharing ends only via an edit or seq_rm.
Assisted-by: Claude Opus 5
* quant : drop duplicate GLM5-Next hc_ filter
The hc_ name filter was listed twice in the GLM5_NEXT protection block.
Assisted-by: Claude Opus 5
* glm5-next: scope K-pool cache access to indexed operations
* glm5-next: keep K-pool access in hybrid index memory
* glm5-next: keep mHC graph builders model-local
* glm5-next: mark only touched pools per ubatch
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
Co-authored-by: Piotr Wilkin <ilintar@gmail.com>
* hexagon: add F16 support for activation ops (SILU/GELU/GELU_QUICK/GEGLU/SWIGLU)
Widens ggml_hexagon_supported_activations() to accept F16 (src0/dst/src1
must agree on type), and adds F16 per-thread worker functions in
act-ops.c mirroring the existing F32 workers, backed by new HVX f16
kernels (hvx_sigmoid_f16_aa, hvx_tanh_f16_aa, hvx_mul_mul_f16_aa,
hvx_min_scalar_f16 family).
SILU, GELU, GELU_QUICK, GEGLU, and SWIGLU are verified correct on-device
(QRD8850) via test-backend-ops CPU-diffed correctness tests. SWIGLU_OAI's
F16 path is code-complete and builds clean on host + all 4 DSP arch
variants (v73/v75/v79/v81), but has no F16 test-case coverage in
test-backend-ops and is therefore unverified on-device in this change.
* hex-ops: align macros
* hex-ops: minor formatting
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
GGML_PAD(nbytes, alignment) wraps to 0 when nbytes is within
(alignment - 1) of SIZE_MAX, which silently bypassed the size
overflow guard in gguf_init_from_reader. Reject the tensor before
padding when nbytes + (alignment - 1) would overflow.
Adds a test-gguf handcrafted case (F32, ne = [4, 2^30-1, 2^30+1, 1])
whose ggml_nbytes = 2^64 - 16 lands in the wrap window. Fails on
master, passes with the guard.
* model : support classifier_pooling for ModernBERT rerankers
Assisted-by: Claude Opus 5.5
* model : read classifier pooling type in load_hparams
Write classifier.pooling_type from _try_set_pooling_type whenever the
config has classifier_pooling, and read it in
llama_model_base::load_hparams. ModernBERT falls back to mean when it
is unspecified.
Assisted-by: Claude Opus 5.5
* conversion : only accept cls and mean for classifier_pooling
Assisted-by: Claude Opus 5.5
* model : rename classifier_pooling_type to pooling_type_cls
Assisted-by: Claude Opus 5.5
* hex-concat: reduce pkts in gather/transpose hot loop
gather directly into dst buffer, use special instruction for gather sync
* hex-concat: use fastdiv
replace calls to sw divide with fastpath
* hex-concat: optimize DMA-HVX pipeline and add transpose helpers
Without CUB (HIP, MUSA) argsort ran the bitonic kernel with one thread
per padded column, so any row above 1024 entries launched an invalid
block configuration. Each thread now owns several columns, every stage
of the network runs all owned columns before the barrier, and the block
is capped at 1024 threads. Shared memory becomes the only bound, which
supports_op checks against the device instead of a fixed 1024.
Rows up to 1024 run the same work as before. Bit-exact with the CUB
path on rows of 2048.
It turns out Intel doesn't particularly like loading F32s one at a
time and we already have the _2aliagned load logic in mul_mat_vec,
so here we use it.
While we do already check all the requirements to load elements 4
at a time across [B]F16 and F32, it turns out [B]F16 loading 4 at a
time is sometimes slower on very specific shapes on Intel BMG.
Loading 4 at a time is a bit faster on F32, but its not material
and I assume might be slower on other platforms.
Note that we also need to validate `a_offset` is 2-aligned in
`mul_mat_vec.comp`, which was missing in the original 2-way-load
patch.
Some selected speedups from `test-backend-ops perf` on a B60.
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=1,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1704 runs - 767.17 us/run - 117.44 MFLOP/run - 153.08 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=1,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 2556 runs - 529.81 us/run - 117.44 MFLOP/run - 221.66 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=2,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1704 runs - 727.13 us/run - 234.88 MFLOP/run - 323.03 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=2,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 2130 runs - 528.84 us/run - 234.88 MFLOP/run - 444.15 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=3,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1704 runs - 702.19 us/run - 352.32 MFLOP/run - 501.74 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=3,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1988 runs - 532.14 us/run - 352.32 MFLOP/run - 662.08 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=4,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1278 runs - 919.50 us/run - 469.76 MFLOP/run - 510.89 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=4,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1917 runs - 543.69 us/run - 469.76 MFLOP/run - 864.03 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=5,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1197 runs - 892.12 us/run - 587.20 MFLOP/run - 658.21 GFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=5,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1881 runs - 575.17 us/run - 587.20 MFLOP/run - 1.02 TFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=8,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1498 runs - 716.40 us/run - 939.52 MFLOP/run - 1.31 TFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=8,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 1819 runs - 576.36 us/run - 939.52 MFLOP/run - 1.63 TFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=512,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 134 runs - 7467.09 us/run - 60.13 GFLOP/run - 8.05 TFLOPS
MUL_MAT(type_a=f32,type_b=f32,m=4096,n=512,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): 134 runs - 7478.12 us/run - 60.13 GFLOP/run - 8.04 TFLOPS
mut_mul_id selected its matmul tile with total token count.
For MoE dispatch grid the true N per workgroup is per-expert rows.
At pp128 on Sarvam 30B that is 6, not 128, so the picker took the l-tile for ~6 live rows.
Most workers in each group had nothing to do.
This wasted time. The slow part was 55% of the whole job.
* fix c++ odr by properly using GGML_COMMON_DECL_CPP
* using actual field rather than macro
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Co-authored-by: XZiar <xziar@xziar.xziar>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* vocab : keep </s> NORMAL in PLaMo-2 and PLaMo-3
The PLaMo-2 and PLaMo-3 vocabularies mark </s> as NORMAL. Current
EOG token heuristic matched it by text and added its attribute
to CONTROL.
Skip this heuristic for the PLAMO2 vocab type so </s> stays NORMAL
and is not treated as EOG.
* use <|plamo:eos|> for detection
With --path or --no-ui, /sw.js returned 404, and a 404 does not remove a service worker, so browsers kept showing the cached built-in UI. Serve a worker that unregisters itself, clears its caches and reloads open tabs. A sw.js in the --path folder is still served first.
Assisted-by: Claude Opus 5.5
- Check for buffered write errors when closing downloaded files.
- Use UTF-8 paths when writing ETag files on Windows.
- Write in binary mode on Windows.
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
graph_inputs was populated while splitting the graph, so it only
contained the inputs that are used as srcs of some node. With pipeline
parallelism (n_copies > 1), each graph input contributes n_copies leafs
to graph_copy, so switching between batches that consume different
inputs (e.g. token batches that do not use the embeddings input vs
image batches that do) changed the graph composition. This shifted the
input copies in graph_copy, making the backend ids comparison report
spurious changes and forcing the scheduler to re-reserve. The
re-reserve could then record smaller input sizes (e.g. out_ids with
n_outputs = 0) and abort later on a graph with an unchanged size via
GGML_SCHED_DEBUG_REALLOC.
Collect the inputs after the split instead, from all input leafs of the
graph, so that the graph composition depends only on which inputs
exist, not on which inputs are used.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* musa: build the docker image and CI container from the MUSA SDK images
Use registry.mthreads.com/mcconline/musa_sdk:5.2.0-{devel,runtime}-ubuntu22.04-s5000
instead of registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64
for the MUSA docker image and the MUSA CI container, and let the runtime stage use the
runtime image instead of reusing the devel one, which drops the MUSA toolchain from the
published images.
* musa: install the MUSA headers and loader path the SDK images omit
musa_sdk:5.2.0-*-s5000 does not ship the cub and thrust headers that the MUSA
backend builds against, and its runtime image does not register
/usr/local/musa/lib with the dynamic loader.
Install both header packages in the build stage and in the MUSA CI container,
and write the loader path in the runtime stage.
* musa: install libmthreads-compute for the MUSA runtime library
The MUSA SDK images do not install libmthreads-compute, which provides
libmusa.so.1 in /usr/lib/x86_64-linux-gnu, so linking anything against the
MUSA backend fails.
* musa: install libmthreads-compute in the runtime stages
The MUSA runtime image does not install libmthreads-compute, so the published
images would have no libmusa.so.1 at run time.
---------
Co-authored-by: yeahdongcn <yeahdongcn@users.noreply.github.com>
* ggml : speed up model loading
A crafted model could hang the server for a very long time, try with:
llama-cli -hf angt/test-gguf-1Mkv -hff model.gguf
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
* Avoid empty keys
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
* Fix
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* ci : update the oneAPI toolkit to 2026.1
oneDNN is removed from Intel Deep Learning Essentials in 2026.0, so
staying on the deep-learning-essentials path would silently lose oneDNN
support when the toolkit version is updated. Switch both the Ubuntu and
Windows CI jobs to the new unified Intel oneAPI Toolkit installer,
which still includes oneDNN (until 2027.0) and keeps the component IDs
unchanged for the Windows install script.
Measured with the same code (b10899) built with oneAPI 2026.1 vs the
2025.3-based release build on Arc B570: prompt processing 1331 vs 434
t/s (3.1x), token generation 50.1 vs 45.3-48.0 t/s.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* docs : update the SYCL backend build requirements for oneAPI 2026.1
With the 2026.0 release the Base toolkit and the HPC toolkit are
combined into the oneAPI Toolkit, and oneDNN is removed from the Deep
Learning Essentials package. Update the install instructions, the
verified release table and the news section accordingly.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* ci : update the release workflow for oneAPI 2026.1 and Level Zero SDK 1.33.1
Align the release package build with the CI build update:
- oneAPI toolkit 2025.3.3 -> 2026.1 (the unified oneAPI Toolkit)
- Level Zero SDK 1.28.2 -> 1.33.1, and the Debian package names
(level-zero/level-zero-devel -> libze1/libze-dev)
- The Windows DLL copy list for the 2026.1 runtime: sycl9.dll and the
.6/.3 MKL library versions
Assisted-by: GLM (z-ai/glm-5.3-flash)
* ci : remove the removed .spv fallback files from the Windows DLL copy list
oneAPI 2026.1 no longer ships libsycl-fallback-bfloat16.spv and
libsycl-native-bfloat16.spv (the OpenCL fallback mechanism changed), so
the copy step failed with exit 1.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* devops : update the oneAPI toolkit image in the Intel Dockerfile
Assisted-by: GLM (z-ai/glm-5.3-flash)
---------
Co-authored-by: Asahi-Prv <Asahi-Prv@users.noreply.github.com>
* server : support multimodal input for /v1/embeddings (Qwen3-VL-Embedding)
Accept the OpenAI-style wrapped content array format for multimodal
embedding requests. Each {"content": [...]} object is one input that
produces one embedding; text parts are concatenated and image_url parts
are decoded via handle_media then spliced with process_mtmd_prompt.
The legacy formats (plain string, token arrays, mixed arrays, and the
{prompt_string, multimodal_data} object) continue to work unchanged via
tokenize_input_prompts. Bare content arrays (the unwrapped shape) are
rejected with a migration message.
Also disables KV prefix reuse for stateless embedding/rerank tasks so
that repeated inputs do not incorrectly share cached KV across requests.
Assisted-by: Opencode Qwen3.8 27B
* clean up comments and docs
* refactor
* add tests
* support video and audio inp
---------
Co-authored-by: timothywang21 <timothywang21@users.noreply.github.com>
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* models: pad on the left with ggml_pad_ext
The Parakeet, LFM2-Audio, Granite Speech and Gemma 4 audio encoders
build a left padding as a right pad followed by a roll, and DFlash2
concatenates a zero filled block in front of the previous tokens.
ggml_pad_ext does both in one node now that every backend supports a
left padding. The Gemma 4 audio embeddings are bit identical.
* models: skip the DFlash2 taps that only read padding
A tap at or past block_size shifts every row out of the block, so its
term is zero. The loop runs min(kernel_size, block_size) taps.
* adapt common
* add common_batch
* wip
* wip: spec
* cont
* common_speculative_process
* server_batch to use common_batch
* rm some stale calls
Assisted-by: Claude Fable 5.1
* migrate mtmd
* handle imrope, handle return val of add()/add_embd()
* add spec zeros vector
* add warning on zero fill path
* tests : use llama_context_ptr in test-recurrent-state-rollback
Replace raw llama_context pointers with llama_context_ptr and drop the
manual llama_free calls and cleanup lambda.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* tests : run test-recurrent-state-rollback over all dummy models
Add a --models DIR mode that mirrors test-save-load-state: iterate every
dummy model, report PASS/FAIL/SKIP in a table and fail only when a model
fails. Register a single ctest entry with ARGS --models instead of the four
per-model registrations.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* cont : fix typo
* metal : allow fusing 0-element nodes to keep graph packing shape-independent
The fusion packing in ggml_metal_fusion_max excluded 0-element tensors and
the topk_moe/moe_reduce checks rejected n_tokens == 0, so graphs decoding
batches with no outputs packed differently from the worst-case reserved
graph. The Metal optimizer then reordered the nodes differently and
ggml_gallocr_needs_realloc failed on the layout mismatch, forcing an
unexpected graph re-reserve (caught by GGML_SCHED_DEBUG_REALLOC).
Treat empty tensors like their non-empty counterparts: match them in the
pattern sequence and only reject genuinely malformed shapes. Fused kernels
dispatch zero threadgroups for empty graphs, which is a legal no-op.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* tests : run test_multi_seq_split_replay as a separate test
test_multi_seq_split_replay was invoked at the end of test_rollback,
so its result was folded into the rollback status and it only ran when
the rollback part passed.
Give it its own test_status return, run both tests independently over
both cache fills via a shared run_tests helper, and report them as
separate rollback / split replay columns in the --models table with
per-test summaries. The exit code fails when either test fails.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* tests : loosen the split replay nmse bound to 1e-4
test-generate-models seeds its weights from std::random_device, and some
generated lfm2 models drift up to ~1.7e-5 nmse on the split replay due to
rounding noise, tripping the previous 1e-5 bound. Raise the bound to 1e-4
so the random generations stop flaking.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* tests : reuse run_tests_for_model in single-model mode
The single-model path duplicated the model init and the non-recurrent
check from run_tests_for_model; route it through the shared helper
instead. Model load failures now return FAIL rather than SKIP so that
--model with a broken file still exits non-zero, and the helper loads
with model_only like the --models loop does since the tests create
their own contexts.
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
* ggml-cpu: enable tiled flash attention for non-vector-multiple head dims on x86
* add AVX2 support for masked loading and storing in simd_gemm_ukernel_tail
* ggml-cpu: fix FA softcap handling for padded KV tiles
* metal: support left and circular padding in GGML_OP_PAD
Align Metal with CPU, CUDA and Vulkan: shift the source coordinates by
the left paddings, wrap them around with the same wrap_around when
circular, and read the source through nb00, which also fixes a right
padding of a permuted source. A test case covers it.
Drop the f32_4 kernel: its selection is disabled as slower, and it
fails two pad cases once enabled.
* metal: use a function constant for the circular pad variant
Address review from ggerganov: replace the bool template with FC_PAD,
as FC_upscale_aa does, so the pad kernel is compiled once and
specialized per pipeline.
* context : do not re-reserve the scheduler when toggling causal_attn
`llama_context::set_causal_attn()` marks the scheduler to do a full re-reserve on every change of the flag. For vision inputs, this flag is flipped twice around each non-causal image chunk for Gemma models, resulting in two expensive `sched_reserve()` passes per image. This is especially slow for multi-image or video inputs.
The cost of a re-reserve scales with context and ubatch configurations, so larger settings pay more per image (see table below).
The re-reserve is unnecessary in this case because `causal_attn` only changes the values written to KQ mask, not tensor shapes or any other buffer sizes.
Note: `causal_attn` is a graph reuse key (`llm_graph_params` via `cparams`), so a new graph is built regardless of `sched_need_reserve`, so this doesn't change the graph rebuilding behaviour.
llama-server with gemma-4-26B-A4B Q4_0 + BF16 mmproj, 130-token images,
cache_prompt=false, prompt_ms median of 3 (before -> after):
| images | config | H200 before -> after | RTX 4090 before -> after |
|-|-|-|-|
| 1 | `-c 8192 -ub 512` | 134 -> 105 ms (1.27×) | 201 -> 119 ms (1.69×) |
| 24 | `-c 8192 -ub 512` | 2278 -> 1562 ms (1.46×) | 3559 -> 1748 ms (2.04×) |
| 24 | `-c 32768 -ub 2048` | 5379 -> 1584 ms (3.40×) | 13377 -> 1759 ms (7.61×) |
Generated output remains identical before and after.
* qwen4exp : make the indexer bias shape independent of causal_attn
The block/cell bias path was selected on cparams.causal_attn, so the
causal and non-causal graphs differed in tensor shapes and ops. With the
re-reserve removed (previous commit), a runtime flip resulted in
reallocating the compute buffers, which would fail under
GGML_SCHED_NO_REALLOC.
This commit selects the block path from the mask shape only, independent
of causal_attn. causal_attn is instead passed to set_input_qsa.
causal_attn is fixed per graph as it's part of the reuse key. Causal
values are unchanged. Non-causal values now follow the reference rule,
where every visible block competes on score and only unpooled cells are
always selected.
* context : state the causal_attn shape rule in the comment
* cont : add TODOs
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* vulkan: read the batch stride of an in place src0 from nb[2]
A dim01 contiguous tensor can still be a view whose batches are
strided by more than ne[1] rows, the first rows of a KV cache for
example. Both the mat-vec and the matrix paths read such a tensor in
place but passed ne00*ne01 as the batch stride, so every head past
the first read the wrong rows. The same applies to src1. The stride
now comes from nb[2] whenever the tensor is used in place; the value
is unchanged for a contiguous tensor.
test-backend-ops gets an m_v parameter on test_mul_mat, the number of
rows of a in memory, and two cases at the shapes of a decoder self
attention over a cache.
* vulkan: size the in place A and B ranges by their strided extent
The matrix path bound src0 and src1 to the shader with a range of
elements times type size, which ends before the batches of a strided
view. Pipelines with bounded access read zero past that range, so the
same view that the mat-vec path already handles gave wrong results
on Intel and on NVIDIA without coopmat2. The range now comes from
ggml_nbytes when the tensor is read in place.
* vulkan: address review from jeffbolznv
Bind the in place A and B of the matrix path with ggml_vk_subbuffer,
which spans to the end of the buffer, so a strided view is in range
without computing its extent.
mul_mat_id reads the batch stride of an in place src0 and src1 with
the same helper as mul_mat. test_mul_mat_id gets an m_v parameter,
the number of rows of as in memory, and a case whose experts are
strided by more rows than it uses.
* vulkan: read the batch stride of an in place src0 in mul_mat_vec_id
The single token path of mul_mat_id passed ne00*ne01 as the batch
stride of A, so a strided expert view read the wrong rows. The stride
now comes from ggml_vk_batch_stride like the other three paths, and
src1 follows the same rule.
test_mul_mat_id gets a single token case over the strided view.
* vulkan: address review from jeffbolznv
The batch stride of an in place tensor is taken from nb[2] as
nb[2] / type_size * block_size, which holds when nb[2] is padded and
not a multiple of nb[1]. A test_mul_mat case with a padded batch stride
covers it.
* vulkan: keep the A and B ranges exact in mul_mm
The quantized A loads of mul_mm carry no row bound and rely on the
descriptor range to read zeros past the last row of a partial tile.
Binding A and B up to the end of the buffer let those tiles read the
leftovers of a previous node and hung the NVFP4 mul_mm on NVIDIA
without coopmat2. The range is the strided extent of a tensor read in
place and the staged size otherwise.
* server : allow splitting RANK pooling for causal LLM rerankers
Rerank models fall into two categories: bidirectional cross-encoders
(BERT, etc.) that require all tokens in a single physical batch, and
causal LLMs repurposed as rerankers (Qwen3, Qwen3-VL) that can use
chunked prefill like any other decoder.
Previously the server rejected all RANK-pooling inputs larger than
n_ubatch, and the graph builder hardcoded QWEN3/QWEN3VL arch checks to
determine last-token pooling. This broke long-document and multimodal
reranking for causal models.
Fix: expose llama_get_causal_attn(ctx) so the server can check the
effective runtime attention type (reflecting any --attention override
or set_causal_attn call). Also expose llama_model_is_causal(model)
for querying the static architectural property from GGUF metadata.
can_split() now permits chunked prefill for RANK pooling when the
context is causal. The graph builder's inline arch check is replaced
with the same cparams.causal_attn predicate, removing the duplication.
Assisted-by: Opencode/Qwen3.8-27B
* remove unused llama_model_is_causal, fix whitespace
Assisted-by: opencode
---------
Co-authored-by: timothywang21 <timothywang21@users.noreply.github.com>
- register --rpc unconditionally and call llama_supports_rpc() only from its handler
- print server "initialization ..." log after args are parsed
Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
Recent PLaMo-3 models use YaRN, while some earlier PLaMo-3 models do not.
The recent PLaMo-3 store their YaRN settings as flat config keys
(rope_scaling_factor, initial_context_length) and build the dict at runtime
in Plamo3Config.rope_parameters. The current converter misses these settings
and writes plain RoPE metadata to GGUF. Mirror the runtime settings into
rope_parameters so the corresponding rope.scaling.* is written to GGUF.
The SYCL FWHT covers 64 to 512 via the standard butterfly network, plus
384/640/768/1280 via the Kronecker/Paley construction added separately in
Hadamard hint can produce (1024, 2048, 4096, 8192); those still fall through
to the default case and run as a dense GEMM against the materialized
rotation tensor, correct but O(n^2) instead of O(n log n).
fwht_kernel_wide runs one row per work-group instead of per sub-group, so
each work-item keeps N/NT values rather than N/WARP_SIZE. Butterflies below
the sub-group width still shuffle; those up to the work-group width go
through work-group local memory; the rest stay in registers. Same butterfly
and sign convention as the existing narrow kernel.
ggml's SYCL backend registration (dpct::dev_mgr) unconditionally requires a
GPU-labeled platform to exist and throws before any op-level test can run,
so test-backend-ops could not be exercised on this box (a GPU-less pod) even
via the CPU device. Verified instead with a standalone harness: the same
kernel body run through a real SYCL CPU device (Intel oneAPI DPC++ 2026.1,
OpenCL CPU backend), checked against an independent recursive-doubling
Hadamard reference, cross-validated by first running the existing unmodified
narrow kernel through the identical harness and confirming it passes (rules
out a reference-convention bug before trusting a pass on the new code).
Random-input results for all four widths, single- and multi-row:
N=1024 NT=256 rows=1 max_abs_err=1.7e-07 max_rel_err=4.9e-04 PASS
N=2048 NT=256 rows=1 max_abs_err=1.9e-07 max_rel_err=2.0e-04 PASS
N=4096 NT=256 rows=1 max_abs_err=2.0e-07 max_rel_err=1.4e-04 PASS
N=8192 NT=256 rows=1 max_abs_err=2.5e-07 max_rel_err=3.8e-03 PASS
N=1024 NT=256 rows=7 max_abs_err=2.4e-07 max_rel_err=1.0e-03 PASS
N=2048 NT=256 rows=5 max_abs_err=3.0e-07 max_rel_err=9.4e-04 PASS
N=4096 NT=256 rows=3 max_abs_err=2.7e-07 max_rel_err=1.7e-03 PASS
N=8192 NT=256 rows=2 max_abs_err=2.5e-07 max_rel_err=1.9e-03 PASS
This covers the kernel algorithm itself; it does not exercise the ggml
dispatch/supports_op integration end to end, which needs a real GPU (or a
SYCL GPU plugin) to get past backend registration. test-backend-ops build
is verified: fwht.cpp recompiles with zero warnings as part of ggml-sycl.
* hex-topk: trying to improve/cleanup the pipeline
* hex-sampling: add STEP op
* hex-sampler: add SUM op
* hex-sampler: update CPY to support sampling cases
* hex-binary: add support for chunking to handle large logits
* hex-argmax: super basic version of ARGMAX
* hex-binary: support for scalars in extended buffers
* hex-binary: fix wrong indexing for dim 1 broadcasts across dim 2 slices
* hex-argsort: fix missing header
* hex-sampler: cleanup dma usage in the sampler related ops, and binary
* hex-build: disable autovectorizer, it is better to use explicit hints for critical loops
* hex-binary: fix perf regression due to is_1d fallback
* hex-ops: update supported ops
* cuda: add F16 input to the FWHT
The CUDA FWHT accepts F32 input only. This makes the source type a template
parameter, so the kernel reads an F16 source directly instead of requiring a
converted copy. The F32 path is unchanged.
supports_op accepts an F16 src1 against an F32 src0 for the Hadamard hint.
Every other F16 src1 against a non-F16 src0 is still refused.
ggml_cuda_op_mul_mat_use_fwht is the single predicate both supports_op and
the dispatch call now share, checking contiguity and same-shape(src1, dst)
in addition to the type/hint conditions above. Without a shared predicate,
supports_op could admit an op that ggml_cuda_op_fwht then rejects only after
the unconditional same-shape assert has already fired; that gap predates
this change (it applies to the existing F32 path too) but this PR is what
touches supports_op, so it closes it here.
test-backend-ops on an A10 (lambdalabs): MUL_MAT 1297/1297, including all
24 Hadamard cases (18 existing F32, 6 new F16).
* cuda: use ggml_cuda_cast in the FWHT load, drop the comment
* llama : add discard for deferred state writes
* llama : add tensor zeroing helper for backends without tensor memset
* llama : clear K/V data after failed sequence restore
* llama : clear recurrent state data after failed sequence restore
* llama : simplify discard and restore cleanup
* llama : report error when abnormal cell count is found in state_read_meta
* llama : clear attention state on hybrid restore failure
* tests : cover failed state restore cleanup
* llama : clear MLA state on dsa restore failure
* tests : update test for rebased test suite
* llama : clarify comment in llama_memory_recurrent::state_read
* Added tiled mul_mat.
For each mul_mat_one_chunk, quants are unpacked into (max) 256x256 tiles of int8,
one routine per quent. Then microkernel computes 16x16 tiles before writing out
256x256 float reults to main memory.
Tests/benches in tests/test-tiled-mulmat.cpp. 3-6x speed improvement
for large matmul, break even at 4096x64 * 64x4096, 80% performance (net
loss) for GEMV. Error rates trivial (order of 1-e04 max, 1-e05 rmse).
* Fixes for ARM/windows builds
* more windows fixes, ggml-cpu.h isn't visible in MSVC for some reason
* unified iqp + tiled on the Q5_K, IQ4_XS set for benchmarking, updated benchmark
* Fixed accidental removal of llama_build_and_test(test-backend-ops.cpp)
* First integration of iqp code
Co-authored-by Bartowski <3266127+bartowski1182@users.noreply.github.com>
* Cleaning up declaration of iq unpacking helpers to align with the bit unpackers
* Removed iqp path
* Fix cross-platform warnings
* Disabling benchmarks unless explicitly enabled
* Fix backend_init for DLL-based builds, add self and bartowski to CODEOWNERS for tiled
* Put benchmarks behind a flag
* kernel fix for AVX2, iq quants
* Fix for asan, leaking memory in test-tiled-mulmat and avoid stack use after return
* guarding env flags with std::call_once
* Simplified repacking for VNNI to a single call per macrotile
* No threadlocals anymore, aligned wdata access
* Doing aligned reads since we ensure alignment with padding in wdata
* Eliminated per-thread gather of Q8_K rows in mul_mat_id, we now gather/repack in a single pass. Repack method now takes pointer array to support both dense/normal and mmid paths. Interface with ggml-cpu.c simplified as a result
* Unified/simplified dispatch and support checks. Put details on wdata needed inside the kernel.h body, simplified interactions with ggml-cpu.c.
* Cleanup includes and whitespace, update src1_repack to return false if we don't need a special repack, so the common case is handled by driver
* Better detection of win32 and additional whitespace fixes
* Gating fuzz tests behind a parameter and some extra prints to try and fix slow CI hosts
* Optimized AVX2 kernel
* Changed interleave format and added ability to interleave in-place after dequant
* Repacks now happen in-place, 16x64 microtiles are independent of each other
* Only repack rows in groups of 16 as they're needed. Save work in low n_rows cases and optimize L1 usage in other cases
* Use long panels for memory-bound regime (M <= 16), reintroduce IQP path for benchmarks
* Fix unused warnings and cleanup. Improved IQ dequantization speed.
* Removed separate process benchmarks
* Revert "Removed separate process benchmarks"
This reverts commit 0688cf43d5.
* AVX2 optimizations and guards for tests on windows
* Removed temp perf harness
* Remove perf-mulmat from build
* Removed IQP path, simplified tests to not use sub processes
* Cleaning up alignment of wdata
* Whitespace fixes and aligning L2 workspace to clean 512kb boundaries
* Update ggml/src/ggml-cpu/tiled/tiled-kernel.cpp
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Cleanup merge-duplicated declaration of test-backend-ops target
* Undo accidental line deletion in ggml.c
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
The original function was broken on Windows for some unicode paths
Paths without a trailing separator now create the last directory too,
matching the function name. All current callers already include a
trailing separator, so this change does not affect them.
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
which resulted in different greedy transcripts for 4.5% of English and 6.5% of Japanese
test utterances. In Japanese, some differences changed entire words.
This change:
* uses `log(x + 2^-24)` instead of clamping to the log floor
* uses a symmetric Hann window, equivalent to `torch.hann_window(periodic=False)`
* adds the normalization epsilon to the standard deviation instead of inside the square root
Only the `lfm2a` preprocessor opts into these behaviors. Other audio preprocessors are unchanged.
Tested on top of 84e76d8 using `llama-server` with CUDA and `temperature=0`, compared against
http://github.com/Liquid4All/liquid-audio fp32.
Test set:
* 200 LibriSpeech `test-clean` utterances (EN)
* 200 Common Voice `ja` test utterances (JP)
* identical 16 kHz audio passed to both implementations
| Greedy transcript identical to `liquid-audio` | Without fix | With fix |
| --------------------------------------------- | ----------: | ----------: |
| EN F16 | 191/200 | 200/200 |
| JP F32 | 187/200 | 200/200 |
| JP F16 | 187/200 | 199/200 |
The remaining JP F16 difference is a comma and matches the reference implementation's own bf16
output.
Mel relative L2 error versus `liquid-audio`:
* EN: 3.2% -> ~2e-6 median
* JP: 3.9% -> ~2e-6 median
* opencl: add A8 Q5_K non-MoE non dp4a + dp4a binary kernel
* opencl: fix s transpose - s only transposed for bin kernels
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* vulkan : fix build issue of legacy glslc version by adding GGML_VULKAN_COOPMAT_GLSLC_SUPPORT macro check for Intel FA shader compiling
* vulkan : add preprocess condition to filter out unsupported FA 2 phases kernels before creation.
* vulkan : move lock_guard for Intel FA shader pointer creation under CM1 compiling preprocessor
* metal: FWHT kernels for block widths above 512
The Metal FWHT covers widths 64 to 512, one row per simdgroup with N/32 values
per lane. Wider blocks need more registers per lane than that layout allows.
kernel_fwht_tg runs one row per threadgroup with 256 threads, so each thread
keeps N/256 values. Butterflies below the simdgroup width still shuffle, those
up to the threadgroup width go through threadgroup memory, and the rest stay in
registers. Same butterfly and sign convention as the simdgroup kernel.
Widths 64 to 512 keep the simdgroup kernel. 1024 through 8192 use the new one,
for both F32 and F16 sources.
The wide kernels allocate float[N] of threadgroup memory, 32 KB at 8192, so the
size check takes the device limit and reports those widths as unsupported where
they would not fit. Without that a device with less threadgroup memory would
accept the op and then abort on a nil pipeline.
test-backend-ops on M5 Pro: MUL_MAT_HADAMARD 26/26, MUL_MAT 1265/1265.
* cont : add TODOs
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* rpc : include nb in the get_alloc_size cache key and floor the result at ggml_nbytes
* cont : remove redundant comment
* cont : add TODO
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* musa: use 16-byte copies for MUSA like sm_70+
ggml_cuda_get_max_cpy_bytes() derives the copy width from __CUDA_ARCH__. mcc
never defines it, so MUSA fell into the generic branch and returned 8 bytes
instead of the 16 bytes that every sm_70+ target gets. The value sizes the
per-thread copy unit of the FlashAttention K/V staging code (fattn-common,
fattn-vec, fattn-tile, fattn-mma-f16 shared-memory loads) and of mmq-vec-dot,
so every MUSA FlashAttention kernel moved half as many bytes per instruction.
On an MTT S5000 (mp_31, MUSA SDK 5.2.0) with Qwen3.8-27B-UD-Q4_K_M, -ngl 999,
-p 512 -n 64, -fa on: 751.15 -> 794.73 t/s prefill and 15.59 -> 15.69 t/s
decode. -fa off is unchanged (1050.05 -> 1052.86 t/s prefill), FLASH_ATTN_EXT
is unchanged (3984 ok / 0 fail / 1323 unsupported) and perplexity is
unchanged.
* musa: enable the CUB paths on MUSA
GGML_CUDA_USE_CUB and USE_CUB are selected by "CUDART_VERSION >= 11070", which
the MUSA SDK never satisfies: CUDART_VERSION is not defined anywhere under
/usr/local/musa/include, so the condition is always false and every CUB-based
path stayed compiled out on MUSA even though the SDK ships CUB and the kernels
build for mp_31. Select them from GGML_USE_MUSA as well. The device-wide
algorithms are usable too: cub::DeviceSegmentedSort compiles and produces
correct results on mp_31.
This lifts the ne[0] <= 1024 limit that ggml_backend_cuda_device_supports_op
applied to ARGSORT and TOP_K on MUSA. On an MTT S5000 (S5000, mcc 5.2.0):
ARGSORT 48 ok / 52 not supported -> 100 ok / 0 (CUDA parity), TOP_K 0 ok /
354 not supported -> 527 ok / 0. The other 20 per-op suites are unchanged, the
Qwen3-0.6B f16 (14.4679) and Qwen3.8-27B iq4_nl (5.1724) perplexities are
unchanged, and the 0.6B graph keeps the same nodes and splits (18 CPU + 18
MUSA0, SET_ROWS 1008) as before.
* musa: take the upstream code path where the toolkit supports it
Several guards were written for an older MUSA toolkit. Verified against MUSA SDK
5.2.0 and on an MTT S5000 (mp_31):
- device init: query cudaDevAttrCooperativeLaunch instead of hardcoding false.
The device reports cooperativeLaunch=1 and musaLaunchCooperativeKernel works
(verified with a kernel whose result was checked).
- device init: keep prop.warpSize instead of overriding it with 32. The device
reports 32 anyway, so this only removes the divergence.
- CUDA_SET_SHARED_MEMORY_LIMIT and the FA shared-memory raise: musaFuncSetAttribute
returns success and sharedMemPerBlockOptin is 192 KiB, so the kernels can use
more than the default 48 KiB.
- vendors/musa.h: add the cudaDeviceGetAttribute and cudaDevAttrCooperativeLaunch
mappings the device-init change needs.
Measured on one S5000 with Qwen3.8-27B Q4_K_M (-ngl 999, -r 3): pp512 968.27 ->
957.09 t/s, tg64 10.09 -> 10.23 t/s, FLASH_ATTN_EXT sweep identical (3975/3982
both), perplexity identical (80.2841 +/- 7.26772 both).
* musa: drop compile-time guards that MUSA's runtime gates already cover
mcc never defines __CUDA_ARCH__, so the arch-gated fallbacks in this group
were already taken on MUSA and the GGML_USE_MUSA guards on top of them only
kept the upstream text from being compiled:
- wkv.cu: the "#pragma unroll" suppression has no effect on the generated
code that is not already covered by the surrounding guards
- common.cuh: the MUSA-only __builtin_unreachable() in no_device_code() is
not needed to silence the compiler
- ssm-scan.cu: the SSD (Mamba-2 prefill) block and its dispatch are gated at
runtime by GGML_CUDA_CC_IS_NVIDIA(cc) and turing_mma_available(cc), which
are both false for PH1 (cc 0x100310), so compiling them changes nothing
- common.cuh: warp_reduce_max(half2) is guarded the same way as
warp_reduce_sum(half2) (FP16_AVAILABLE); the MUSA-only guard left the
function with no return statement. It has no caller today.
MTT S5000 (mp_31, MUSA SDK 5.2.0), MUSA_ARCHITECTURES=31: build rc=0. Against
an unmodified build of the same tree on the same card, FLASH_ATTN_EXT
(3984 ok / 0 fail / 1323 unsupported), SSM_SCAN (15/0), RWKV_WKV6 (6/0),
GATED_DELTA_NET (38/0) and MUL_MAT (1299/0/385 unsupported) are identical, and
perplexity with -fa on is bit-identical (5.1639 +/- 0.36673, 4 chunks).
* musa: do not use MMQ on PH1
test-backend-ops on an MTT S5000 (mp_31, MUSA SDK 5.2.0) fails 260 cases and every
one of them goes through the MMQ path:
- MUL_MAT with a batched src1 (any bs/nr != [1,1]): 109 cases across all
quantized types, e.g. 12 of 13 cases at n=16, while the plain [1,1] layout
passes
- every quantized MUL_MAT_ID: 147 cases, while the f16/f32 variants of the same
shapes pass
- MUL_MAT with more than ~512 tokens: 4 cases (n=509..4096); the small-n cases pass
The cuBLAS/dequant path is correct for all of them and the MMVQ path used for
small batches is unaffected, so quantized matmuls now take that path on PH1
instead of returning wrong values. 27B perplexity with default flags goes from
nan to finite, and the full suite reports 0 failures out of 22237 cases.
The MMQ defect itself (fastdiv, __umulhi, uint3 kernel parameters and
__CUDA_ARCH__-based MMA availability were all checked and are correct on this
part) is not addressed here.
* musa: keep the block barrier of the fused TOPK_MOE kernel reachable
topk_moe_cuda returns early for the rows past the end of the graph, but one block
covers TOPK_MOE_ROWS_PER_BLOCK (8) rows, so the last block is only partially filled
whenever n_rows is not a multiple of 8. On MUSA a warp that has already returned
blocks the block wide __syncthreads() below, which makes the kernel hang and the
launch time out. CUDA tolerates the exited warps, which is why the CUDA numbers
never showed it.
For MUSA, clamp the row index of those warps to the last row so that every warp of
the block reaches the barrier; they recompute the last row and write the same
values. The CUDA code path is unchanged.
On an MTT S5000 (mp_31) the fused TOPK_MOE cases change from a launch timeout with
no completed case to 418 ok / 0 not supported / 0 failed, i.e. the CUDA result, and
the other 101 per op suites are unchanged (0 failed, no count changes).
* musa: enable GATED_DELTA_NET
The op was turned off for every MUSA target because mcc could not build the kernel
at the time. The current toolkit builds it: with mp_31 and MUSA SDK 5.2.0 the file
compiles with zero errors and all 36 test-backend-ops GATED_DELTA_NET cases pass
against the CPU reference. 27B perplexity is unchanged.
While the op is refused, the scheduler has no choice but to run it on the CPU: 48
GATED_DELTA_NET nodes per forward pass. On an MTT S5000 (Qwen3.8-27B Q4_K_M, -ngl
999, one container, -r 3):
pp512 (FA off) 964.51 -> 2119.26 t/s
tg64 (FA off) 10.15 -> 15.50 t/s
* musa: name the stream capture query API for the graph aware kernels
argsort.cu and mean.cu call cudaStreamCaptureStatus, cudaStreamIsCapturing and
cudaStreamCaptureStatusNone inside their USE_CUDA_GRAPH blocks, but the MUSA
compatibility headers do not alias those names, so building with the experimental
GGML_MUSA_GRAPHS option fails with 7 errors in those two files. Map the three
names to their musa* counterparts, under the same guard that enables the graph
code, so the default build is untouched.
The option stays off by default: on an MTT S5000 the captured path measured
slower (pp512 693 vs 772 t/s, tg128 15.20 vs 15.39 t/s over two sessions) and the
borderline MUL_MAT cases are not reproducible between runs.
* musa: build the CI and docs for PH1 (MTT S5000)
The MUSA CI job and the documented default still targeted the first generation
(MTT S80, MUSA_ARCHITECTURES=21) while the current MUSA SDK targets PH1
(MTT S5000, 31). Move the job, ci/run.sh's default and the build docs to 31,
and run the job in the PH1 MUSA SDK devel image:
registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64
That image needs two things the previous one did not: python3-venv for the
ccache-buckets step, which builds a virtual environment for the Hugging Face
CLI, and no time prefix on the build command, because container jobs run their
steps with sh and the image ships no time binary.
- #28068 builds the GDN q/k l2norm as ggml_scale(ggml_rms_norm(x, eps/n), 1/sqrt(n)). This adds 2 SCALE nodes per GDN layer, 96 extra kernel launches per ubatch on Qwen3.8-27B (48 GDN layers).
- The extra kernels take no measurable GPU time, but each launch has a host/driver cost. It is small with plain batch processing and about 10x larger with draft-mtp speculative decoding.
- rms_norm_f32 gets a do_scale flag, the same pattern as do_multiply/do_add, so the fused path shares the kernel, the reduction and the launcher. It computes scale * (rsqrt(mean + eps) * x), which matches the unfused rms_norm + scale bit for bit, so #28068 numerics are kept.
- Fusion only fires when SCALE has no bias and the rms_norm output has a single consumer (ggml_can_fuse).
- Metal (#28948) and SYCL (#28931) already fuse the same pattern.
Measured on 2x GTX 1080 Ti (sm_61, PCIe 3.0 x16 + x4), i7-13700KF, Windows 11, driver 582.66, CUDA 12.9.
Qwen3.8-27B-UD-Q4_K_XL, -ngl 99 -ts 53,47 -ot token_embd=CPU, master fee39dd92.
llama-bench -ub 128,512 -p 512,2048 -n 128 -r 5, tok/s:
build pp512@128 pp2048@128 pp2048@512 tg128
master 367.7 419.1 385.4 12.90
master + fix 372.6 420.6 388.6 12.98
+1.3% +0.4% +0.8% +0.6%
llama-server cold prefill, -c 56000 -ub 128 -b 2048, draft-mtp n-max 3 p-min 0.5, mean of 2 rounds x 3 reps:
build pp 8000 pp 20000
master 356.5 322.0
master + fix 371.4 (+4.2%) 337.4 (+4.8%)
- Launches per ubatch go from 1032.9 + 841.7 back to 978.9 + 799.7 (CUDA0 + CUDA1), the b10828 count. The GPU op sum is unchanged.
- test-backend-ops RMS_NORM_SCALE, NORM_SCALE, RMS_NORM_MUL_ADD, RMS_NORM_MUL_ROPE, RMS_NORM, RMS_NORM_BACK, NORM, L2_NORM and SCALE all pass on both GPUs.
- Perplexity is identical to the unfused build: 3.2030 +/- 0.0559 at -c 2048, 16 chunks.
- Draft acceptance counts per request match the unfused build.
Assisted-by: Claude Opus 5.5
- return early when the graph has no nodes
- drop the redundant reset of capture_compute: the decrement at the top
of the function already transitions the counter from 0 to -1, so a
capture happens exactly once
- hint at METAL_CAPTURE_ENABLED=1 in the capture error message
- pass capture_compute == 0 (not the raw counter) as use_capture to
ggml_metal_op_init, so GPU debug-group markers are only emitted on the
captured compute
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : cache sparse FA indices in shared memory
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : simplify shared memory size calculation
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* pi : update general
* metal : unroll sparse index load
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* hexagon: fix accuracy issue in Q8_0 N=1 MUL_MAT
* hex-quant: fix register spills
* hex-mm: use dma for all dyn.quant paths
Co-authored-by: Aparna M P <aparmp@qti.qualcomm.com>
* hex-mm: remove obsolete run_quant_task
* hex-mm: update tracing to properly wrap the events
* hex-mm: use act for activation data in all paths
* hex-mm: use act_ instead of src1_ to avoid confusion in fused kernels
* hex-mm: remove/reroute the rest of the non-DMA act (aka src1) logic
* hex-dma64: yet another pass at cleaning up the dma_addr_t casts
* Update ggml/src/ggml-hexagon/htp/matmul-ops.h
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update ggml/src/ggml-hexagon/htp/matmul-ops.c
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update ggml/src/ggml-hexagon/htp/matmul-ops.c
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Update ggml/src/ggml-hexagon/htp/matmul-ops.c
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Aparna M P <aparmp@qti.qualcomm.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* vulkan: add int8 coopmat quantized matmul shader
* apply scales inline
* use scalar sums
* probe and directly access coopmat values instead of going through shmem
* add q8_0 support
* add BK_STEP to shader, default to 2
* use larger workgroups
* double buffering
* preload scales
* coopmat load first, then wmma
* use float for scales
* add faster RDNA int->float conversion
* workgroup scheduling for cache proximity
* clean up
* use wave32
* restructure for vgpr use
* skip computation for inactive tiles
* only force subgroup size 32 on AMD RDNA
* use BK_STEP 4
* fix compilation
* move quant-specific prefetch function out of main file
* add q4_1, q5_0, q5_1 support
* restructure mmq cm1 functions
* enable mul_mat_id support
* fix segfault
* fix mul_mat_id bug
* support iq4_nl and mxfp4
* remove elem row/col fast path, invalid for RDNA4
* use shmem arrays for LUTs
* use 4-byte loads where possible
* add q3_k, q4_k, q5_k, q6_k and nvfp4 support
* fix l warptile
* improve performance
* improve performance
* improvements
* dedup b scales
* merge shmem arrays
* undo uint8_t, gate to RDNA3/4
* add RDNA4 architecture, use for hardcoded coopmat elem thread access, set BK_STEP back to 4
* improve offset application
* clean up
* fix iq4_nl and nvfp4 performance
* rdna4 tuning
* use BK_STEP 2 on MUL_MAT_ID
* adapt to upstream changes
* fix shmem support function, clean up comments
* fix warptile logic
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
* vulkan: add IQ4_XS support to the coopmat1 integer matmul shader (#28440)
Adds IQ4_XS to mul_mmq_cm1: dedicated block_a_load/block_a_to_shmem that
expand both nibbles of each packed32 word through cm1_kvalues, LOAD_VEC_A 8
and an IQ4_XS-sized a_panel_bytes estimate for the L2-friendly scheduling.
Assisted-by: OpenAI Codex
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* avoid compiling f16 acc shader variants
---------
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* 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
make-release-desc.sh now emits "Changelog since [vX.Y.Z](<repo>/releases/tag/vX.Y.Z)"
instead of a plain version string, so the release notes link back to the previous release.
The repo URL is derived from the origin remote (SSH or HTTPS); if it cannot be
resolved (local run without origin), the title falls back to plain text.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
Resolve the target arch with get_model_architecture so vision targets
(e.g. Lfm2VlForConditionalGeneration) map to their text model for the vocab.
Fix double rope reorder for LFM2/LFM2.5 DSpark drafters
* 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>
Restore get_cache_directory() as fs::path as string() can be lossy on Windows
Partially reverts #29125
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
ggml_conv_1d_dw builds its im2col as f32 when the kernel is bf16, then
multiplies the two, so a depthwise convolution over bf16 weights asks
for kernel_mul_mv_f32_bf16, which was never instantiated. The base, the
_4 and the _short families are filled in next to their bf16 neighbours,
inside the same runtime guard, so a device without bf16 support is
unaffected.
* CUDA: enable sparse-fa for dsv4 prefill (again)
* CUDA: unroll the query loop of the sparse mask scan
The query loop of flash_attn_mask_to_sparse_indices has a runtime trip
count, which keeps the unrolled scan over the values of a lane from
issuing its loads together. Template the kernel on ncols1 so the loop
is bounded at compile time: batch one decodes compile to straight line
code and the scan drops from 46 to 17 us at 49k columns on sparse
decode shapes.
* CUDA: pick the out of bounds check of the sparse mask scan in host code
The query loop of the ncols1 == 8 scan keeps a runtime bound and an
early exit, so it does not unroll past its first iteration. Template the
kernel on whether the last group of queries is partial, decided on the
host from n_queries, and hoist the column bound out of the loop: the
loop becomes straight line code and the batched sparse op at 49k
context drops from 586 to 244 us.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* 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>
The OpenAI chat completions API specifies content part type "video_url"
with a {"url": ...} object, and clients typically send data: URIs
(e.g. data:video/mp4;base64,...). The llama-server only accepted the
non-standard "input_video" type and rejected data: URIs for video
(accept_base64_uri=false), so any OpenAI-conformant client failed with
"unsupported content[].type" or "Invalid uri format".
- accept "video_url" as an alias of "input_video"
- read the media object from whichever key was used
- allow data: URIs for video (data:video/*), as already done for images
This commit adds a new recipe/target to the Makefile which allows the
logits verification to be run on pre-existing model outputs.
The motivation for this is that for large models it can take a long time
to run them models, and especially for the original model which seldom
changes this is very time consuming. With this change we can run the
original model one which will store the tokens and logits, and then
manually run the converted model and the run use this recipe to verify
them against the orignal model.
* dspark: add Gemma 4 draft support
Add GGUF conversion and runtime support for full-attention and SWA Gemma 4
DSpark drafts, including tied output weights and boolean backbone metadata.
Assisted-by: Codex
* dflash: infer Gemma draft features from metadata
* 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>
Since #28732 our internal symbols are exported. A duplicate copy dlopened and
dlclosed by ggml_backend_load_all() then interposes them, so its destructors
destroy the live vk_instance and later device queries hit the GGML_ASSERT on
vk_instance.device_indices. Hidden visibility exports only GGML_BACKEND_API,
as before #28732.
Fixes#29138
Assisted-by: henk:claude-fable-5
Avoid returning references through lambdas that hold a local cast pointer, which triggers -Werror=dangling-reference in some CI compilers. Reuse the precomputed select_expr pointer directly.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* server: route every model load through the queue
A model loaded by the fast path has no queue entry, so tick() evicts
it at its LOADED transition before its own request is proxied. Every
load now joins the queue, whose entry protects the model until its
waiters leave.
* server: do not admit requests into a stopping model
A request for a model that is being stopped still sees it LOADED and
is proxied into the dying child. Such a request now joins the queue
and is served by the next instance. The stopping mark is cleared
under the same lock that sets UNLOADED, so no request can see a
model that is neither stopping nor unloaded while its child is gone.
* vulkan : Intel FA kernel optimization for split k path
* vulkan : Host code update for Intel split k FA kernel path selection, fix A770 Linux op test failures
* vulkan : use symmetric coopMatMulAdd() in flash_attn_decode_phase_1 shader to resolve test op failre on A770 Linux with 26.2.3 mesa driver
* vulkan : fix editorconfig issue in flash_attn_decode_phase_2.comp
---------
Co-authored-by: Liu, Russell <russell.liu@intel.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>
* model : add DFlash layer-input taps for HunyuanVL
DFlash speculative decoding needs the target graph to expose the residual
stream entering each layer (res->t_layer_inp[il]) - the draft model reads
those tensors to build its cross-context. Qwen3 and the other DFlash-capable
targets register them, but the Hunyuan graphs do not, so serving a DFlash
draft against a HunyuanOCR target aborts during the first graph build:
GGML_ASSERT(t_layer_inp[il] != nullptr && "layer input tensor is null")
Register the tensor at the top of the layer loop, mirroring qwen3. The
layer input is the residual stream entering layer il, i.e. the output of
layer il-1, which is what the draft's target_layers metadata refers to
(the converter writes target_layer_ids+1). hunyuan-dense.cpp reuses this
graph, so it is covered as well; hunyuan-moe has a separate graph and is
untouched.
The vector is only read when a speculative implementation enables those
layer ids, so there is no behaviour change without a draft model.
Tested with tencent/HunyuanOCR 1.5 and its DFlash draft: image requests now
run, draft acceptance is ~0.5 and the OCR output is byte-identical to the
non-speculative run.
Co-authored-by: wendadawen <wendadawen@qq.com>
* convert : fix DFlash draft conversion against HunYuan targets
Converting a DFlash draft with a HunYuan target failed in two ways.
1. DFlashModel.set_vocab() reuses the target class' vocab handling by
calling it unbound with the draft instance, but HunYuanModel.set_vocab()
called self._fix_special_tokens(), a method that only exists on
HunYuanModel, so the conversion always aborted with
AttributeError: 'DFlashModel' object has no attribute '_fix_special_tokens'
Make the vocab helpers module-level functions taking the model
explicitly, so they do not depend on the instance being a HunYuanModel.
They have no other callers, so the two id lookups are folded into
_fix_special_tokens().
2. The delegated call runs with self.dir_model pointed at the target but
keeps the draft's self.hparams, so config lookups inside the target's
vocab code (the pad_token_id < 0 guard, eod_token_id) read the draft's
config instead of the target's. That aborts on targets with
pad_token_id = -1 (e.g. the HunyuanOCR v1.0 checkpoint) and otherwise
writes special token ids that disagree with the target.
Add _vocab_hparams(): it returns the target's config (with text_config
merged to the root, as TextModel does) when the model is a draft
converted with --target-model-dir, and the model's own hparams
otherwise, so a normal conversion is unaffected.
Tested: converting tencent/HunyuanOCR/dflash succeeds with both the 1.5 and
the v1.0 target; converting the base model without --target-model-dir
produces a byte-identical GGUF to before.
Co-authored-by: wendadawen <wendadawen@qq.com>
* convert : fix DFlash draft vocab against HunYuan targets
Switch hparams to the target config for the duration of the borrowed
set_vocab(), matching the existing dir_model swap, instead of teaching
HunYuanModel::set_vocab about draft models.
* convert : fix HunYuan special token ids for DFlash drafts
* convert : use load_hparams for HunYuan special token ids
* convert: add MiMo-V2.6 support
Hoist the K3 mxfp4 conversion repack into base.py so it can be reused
Remove decoder from mmproj convert
* Update conversion/mimo.py
* fix: use autoparser
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Piotr Wilkin <piotr.wilkin@syndatis.com>
The snapdragon CI builds packages only to feed the QDC device tests, so
Hexagon NPU binaries never reached the releases page. Build both targets
in release.yml and attach them as release assets.
This commit tweaks the Toaster element to include a close button.
These toasts often cover other UI elements like the model selector, and
this change avoids having to wait for them to disappear on their own
(e.g. after a load failure).
* hex-gdn: start putting together HMX support for GDN
* hex-gdn: working hmx but not-pipelined and slow for now
* hex-gdn: re-write vtcm layout handling and prep for pipelining
* hex-gdn: starting to pipeline hmx and dmas
* hex-gdn: add hvx threading for most pipeline stages
* hex-gdb: add detailed trace events
* hex-gdn: vectorize expfs and use aligned hvx reads/writes
* hex-gnd: vectorize the rest of expf
* hex-gdn: optimize tail processing (pad partial chunks)
* hex-gdb: avoid float up/down casts in hot loops
* hex-fa: remove float up/down casts from inner loops
* hex-gdn: do exp() in f16 to improve HVX utilization
* hex-gdn: optimize tiler
* hex-hmx: bump hmx-queue to 128 and dispatch all GDN gemms at once
* hex-gdn: further pipeline improvements
* hex-gdn: optimize gdn prep stage
* hex-gdn: yet more tweaks to optimize GND_SOLVE task and pipeline
* hex-gdn: improve accuracy and optmize gdn-prep further
* hex-gdn: fix rebase conflict
* hex-bufs: revert max_bufsize enforcement, it is enough to just enforce max_vmem
* hex-scripts: improved inspect script to avoid false alarms in reg spill detector
* hex-fa: improve inline softmax with in-reg VKQ32 accum
* hex-fa: minor improvement for dma pipeline in hvx kernel
* hex-fa: reduce ddr reads by 20-30% during token gen
* hex-gdn: proper alignment for hvx vtcm spads
* llama-context : report graph inputs and input tensors during sched reserve
- fix the tg (token generation) graph bs label to use n_seqs instead of a hardcoded 1
- report the number of graph inputs from llm_graph_result::inputs for both the pp and tg graphs
- report the number of input tensors (nodes and their src tensors flagged with GGML_TENSOR_FLAG_INPUT)
- log a warning when an input tensor has an op other than GGML_OP_NONE
- log a trace line for each input tensor and the nodes (name and op) that use it
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* cont : count input tensors before reserving the sched
* wip
* llama-graph : name the unnamed graph input tensors
- name the kv-cache idxs input tensors (attn_inp_k_idxs, attn_inp_v_idxs)
- name the recurrent state copy idxs input tensor (rs_s_copy)
- report the input tensor shape in the sched_reserve trace
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : rename "graph inputs" to "graph input objects"
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : report the sched reserve graph stats on a single line
- print nodes, splits, input objects and input tensors in one line
- when the pp and tg graphs differ, print each value as 'pp / tg'
and annotate the line with the batch sizes used for each graph
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : pad logs
The 5-argument load_ldmatrix added in 1884824fd only defines tile<16,8>, so the Volta tile<8,4> does not match. See https://github.com/ggml-org/llama.cpp/issues/29222 for details. Building on 1884824fd, generalize the tile shape of the 5-argument load_ldmatrix from <16,8> to <I,J>, so the non-swizzle branch forwards to the 3-argument loader for any shape. Local compilation and testing passed.
Assisted-by: DeepSeek V4.1 Flash (OpenCode)
convert_unary handles the contiguous case through the general strided kernel,
one element per thread: each lane reads 4 bytes and writes 2. Converting the
activations for a bf16 matrix multiplication that way moves 126 MB in 1021 us
on gfx1151, about 65% of what the memory system can do.
Give the contiguous path its own kernel that takes four elements per thread
through a vector type, so a warp loads 512 bytes at a time instead of 128. It
is used only when the element count is a multiple of four and both pointers
carry the alignment the vector type needs, and falls back to the strided
kernel otherwise.
Model level, Qwen3.8-Next-Flash IQ3_XXS on gfx1151, llama-bench -ub 2048 -r 6,
mean of the last 3 reps, ABBA counterbalanced:
pp2048 688.0 680.0 -> 694.3 691.1 +1.26%
tg128 24.8 24.8 -> 24.8 24.8 +0.14%
Every conversion in a prefill takes the new kernel (kernel trace: 1146
convert_unary_cont_vec4, no convert_unary). Output is bit identical; MUL_MAT,
MUL_MAT_ID, CPY, CONT, GET_ROWS and SET_ROWS pass.
Assisted-by: Claude Opus 5
* 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
Allow configuring --temp, --top-p, --min-p, --repeat-penalty,
--presence-penalty and --frequency-penalty via LLAMA_ARG_* so
llama-server can be fully controlled from an EnvironmentFile
(e.g. systemd on Debian).
Use `llama-gen-docs` to regenerate the readme files.
In router mode, authentication belongs to the router. unset_reserved_args()
already unset LLAMA_API_KEY, but did not unset LLAMA_ARG_API_KEY_FILE.
When --api-key-file was passed, children re-validated against file keys only,
causing clients using --api-key to 401 on chat completions (#28820).
In addition, router internal calls without auth headers (such as
POST /v1/streams/lookup and DELETE /v1/stream) were silently rejected with 401.
Unset LLAMA_ARG_API_KEY_FILE in unset_reserved_args() so no API keys reach
child instances. This keeps keys out of child argv, ensures all keys the router
accepts work end-to-end, and prevents router internal stream calls from 401ing.
Fixes#28820
* Fixed json enum handling
Added common_json_value handling for enum values.
Added tests/test-json.cpp to cover testing of some aspects of common_json.
* Removed tests as requested.
* Applied recommended style and simplification
Simplified by delegating enum constructor to the constructor of the underlying type
Matched style of surrounding templating code
* Implemented ARM NEON DP q1 4x4 repack
* Hoisted out scaling by b_d in gemm
* Added 4x8 NEON I8MM repack kernels
* Cleanup for q1 arm repack
* Added missing aliases for arch fallback
* Corrected unused var statements
* Extended table guard condition to account for i8mm w/o dp build
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
* Moved new declarations and references to groups' top
* Moved declarations for uniformity
---------
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
* 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
* ui : let the chat column shrink below its content width
The chat column is a flex item, so its automatic minimum size kept it as wide as the widest row inside it. Message rows cap at max-w-3xl plus padding, so a narrower window pushed a page-level horizontal scrollbar.
Set min-w-0 on the column so the inner scroll containers take over.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : wrap markdown tables in a scroll container
Markdown tables render as a bare <table>, which keeps its content-driven minimum width and can stretch the chat column past the window. The table-wrapper CSS already existed, but nothing produced the wrapper.
Add a rehype plugin that wraps each table in div.table-wrapper, following the existing enhance-* plugins.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : scroll long inline content inside markdown blocks
Long unbreakable content (inline code, paths, hashes) widened the message row and spilled over the neighbour elements. Give each markdown block a horizontal scroll container, and the content root one as well, since the trailing block renders with display: contents and has no box of its own.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : use exact transition properties for markdown images
transition: all repainted every property and 300ms felt sluggish. Name transform and box-shadow at 200ms ease-out, and gate the hover scale behind (hover: hover) and (pointer: fine) so touch taps do not trigger it.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : fit wide image attachments to the message width
Attachment thumbnails used a fixed height with w-auto, so a wide image kept its aspect-driven width and, being flex-shrink-0 in a right-aligned bubble, overflowed to the left of the message row.
Cap the thumbnail with max-height and max-width instead of a fixed height so it scales down proportionally, and let it shrink outside the single-row carousel.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : keep long tool call titles inside the message row
A tool title could not shrink below its content, so a long path escaped the message row. Let the title span shrink and scroll, and for the file tools put the value on its own line only when it does not fit, with the value as the only scroll container.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : render get info as a collapsible block with a table
get_info rendered its own always-open row with the values trailing the label. Use the shared ToolCallBlock chrome so it collapses like the other tools, and list os and cwd as table rows with the key as a row header.
The error and pending states now show inside the body, including the plain-string errors the server tools path produces.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* test : pin the server mode in the add menu a11y story
The story asserts the add menu's first enabled item is the reasoning submenu, which is mounted only outside router mode. The vitest dev server proxies /props to whichever server is running, so the assertion depended on the machine's server mode and failed whenever a router was up.
Pin the mode in the story, including props.role so a re-detection cannot flip it back.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui: wrap long markdown tokens instead of scrolling every block
Making each markdown block and the content root a horizontal scroll
container turns any hover transform into a scrollbar: the blockquote
translate and the image zoom overflow their block and flash a scrollbar
under it. Each block also becomes a block formatting context, so the
paragraph margins stop collapsing across blocks and the spacing doubles.
Drop both overflow-x rules and let long unbreakable tokens wrap with
overflow-wrap: break-word on the content root. break-word leaves the
min-content width untouched, so wide tables and code blocks keep
scrolling inside their own containers.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* 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>
* server-models : show source per model in log
- Show [source] tag (preset/models_dir/cache) per model instead of cryptic * marker
- Show HF hub cache path in the 'Loaded cached model presets' log
- Add hf_cache::get_cache_dir() public accessor
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : pad log
* 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
Add support for the new DSV4 HC op variants used by qwen4exp:
- hc_pre with per-element sigmoid gate (gated variant)
- hc_post with identity mixing (comb == nullptr)
Assisted-by: pi:llama.cpp/Qwen3.8-27B
argsort_f32_i32_cuda_cub called the one-shot DeviceRadixSort::SortPairs
API with d_keys_in == d_keys_out (temp_keys, temp_keys). CUB's internal
double-buffer ping-pong requires distinct key buffers: with aliased
buffers the sort partially overwrites its own input mid-pass and emits a
corrupted permutation, surfacing as intermittent garbage indices (e.g.
backend top_k over a 248k-column vocab on Maxwell/CUDA 12.5/CCCL 2.x,
which then triggered out-of-bounds gathers in downstream get_rows).
Use a distinct keys-out buffer for all six call sites (plain and
segmented, ascending and descending, size-query and execute).
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: Oliver Simons <osimons@nvidia.com>
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.
* ggml-cpu: add F16 input to the FWHT
The CPU FWHT accepts F32 input only. This change makes the source type a
template parameter. The CPU path now accepts F16 input and F32 input.
The CPU MUL_MAT reference now converts an F16 src1 to F32. It does this when
the caller sets the Hadamard hint.
No backend has an F16 FWHT kernel yet. The test cases come with the backend
changes that add one.
* ggml-cpu: assert the F16 FWHT input path, and use the bulk converter
Address review feedback.
The F16 branch writes plain floats into wdata, which is only correct when
vec_dot_type is F32. That invariant held because supports_op only accepts an
F16 src1 for the Hadamard hint with F32 src0 and dst, but nothing enforced it.
Assert it next to the existing src1 type check so widening supports_op cannot
silently break the write.
Replace the hand-rolled conversion loop with ggml_cpu_fp16_to_fp32.
* llama: read the SWA pattern as a period or a per-layer array
Add llama_model_base::load_swa_pattern(), which reads
sliding_window_pattern either as one flag per layer or as a period
expanded by set_swa_pattern(), and use it in every loader that reads
the key as a period.
These loaders silently ignored an array and applied their default
period, although the converters of olmo2, gemma3n and exaone4 write
arrays. The published GGUFs match the defaults, so their outputs do
not change. The loaders that already accepted both forms lose their
duplicated scalar-then-array block, and use their declared default
period when the key is absent.
* model-saver: write the SWA pattern and the MLA SWA geometry
Write sliding_window_pattern as one flag per layer, nextn layers
included, for every model using SWA. The array is never collapsed to
a scalar, since the loaders read a scalar as a period.
Also write the MLA key/value lengths and KV LoRA rank of the SWA
layers, required by dots3note.
This enables the saver for plamo3, gemma3, cohere2, cohere2moe,
olmo2, exaone-moe, afmoe, mimo2, spark2_5, muse-glimmer, mellum,
laguna, granite_swa, dots3note and maple, all passing the bit-exact
roundtrip of test-llama-archs.
* vulkan: add IQ3_S MMQ matmul kernels
* Make block_a_to_shmem do 2-byte loads (110 bytes is divisible by 2)
* Align the check, IQ3_S is also using K tile size
* 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.
* Update to openvino-2026.4
* Update OV docs
* ggml-openvino : fix clangd and MSVC warnings
* fix int to ptr cast, more internal linkage enforcement, and avoiding duplicate switch case
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
* ui: fix accidentally removed reasoning menu in single model mode on desktop
* ui: formatting task run to fix storybook test
* ui: mount the add menu reasoning submenu outside router mode only
The models selector already owns the reasoning submenu in router mode,
so the add menu only mounts it in single model mode. The first enabled
item of the add menu is now the reasoning submenu, the accessibility
story expects it.
---------
Co-authored-by: Ben Babik <work@benjaminbabik.com>
Co-authored-by: Pascal <admin@serveurperso.com>
* ci : add API/ABI check to make-release workflow [no ci]
This commit adds an API/ABI compatibility check to the make-release
workflow.
The motivation for this to allow us to detect any potential breaking
changes in API/ABI compatibility between releases and fail the the
release if there are any.
The workflow can be triggered manually as before and this check can be
skipped if needed as it does take some time which might be useful when
doing a dry-run and not specifically interested in the API/ABI check.
By default this will check the current release against the latest
release, but this can also be configured in the workflow, or in the
script run on the command line, to check a different tag.
* add check for minor version bumps [no ci]
This commit also changes the build type to be RelWithDebInfo so that the
reported information is more useful.
* 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>
Both im2col.comp and im2col_3d.comp declare D_ptr without an explicit
buffer_reference_align, so glslang emits writes through it as Aligned
16. The shaders advance the pointer by D_SIZE, a per-variant define
set to 4 for float and 2 for float16_t, so most write addresses are
not 16-byte aligned. This triggers
VUID-RuntimeSpirv-PhysicalStorageBuffer64-06315 under GPU-AV.
Declaring buffer_reference_align = D_SIZE matches the alignment to the
actual write stride and takes validation hits from 20 to 0 for both
IM2COL and IM2COL_3D.
Fixes#28960
* model: calculate split states for attn_qkv from n_head * n_embd_head_k
required for gemma4 with --fuse-qkv, where n_embd is 5376 but Q is 8192.
* model: handle fused full attention layers for qwen35/qwen35moe
* model: add TODO: [TAG_SPLIT_QGATE_QWEN]
llama probes weight placement with a rope where all params are 0, so rejecting
n_dims == 0 or freq_base == 0 puts rope_freqs on the CPU. That splits the decode
graph at every full-attention layer (gemma-4-E2B: 5 splits instead of 2).
Assisted-by: Claude Opus 5
* 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>
The `sizeof(int16_t)` branch in `permute_transpose_impl` calls
`rvv_transposed_s32_mn_to_nm` instead of `rvv_transposed_s16_mn_to_nm`.
This is a copy-paste bug from the `sizeof(int32_t)` branch above it.
The s32 function uses 32-bit segment load/stores (`vssseg8e32.v`) on 16-bit
data, reading 2x bytes per element and producing completely wrong
transposition results -- 14 out of 16 positions are corrupted for a 4x4
int16 matrix.
The correct function `rvv_transposed_s16_mn_to_nm` already exists (line 390)
and is used elsewhere in flash attention (line 1488).
The server caches the most recent compute graph per device so that
GRAPH_RECOMPUTE can re-execute it without resending tensor data. The
cached graph nodes hold direct pointers to backend buffers that were
live at graph_compute() time. If any of those buffers is later
released via FREE_BUFFER, the next GRAPH_RECOMPUTE re-executes the
cached graph through the dangling pointers (use-after-free).
The bug is reachable by an unauthenticated remote client. The
dangling pointers point into chunks an attacker can reshape via
subsequent ALLOC_BUFFER/SET_TENSOR commands, and the resulting
read/write through the cached graph is sufficient to leak libc
addresses and hijack the buffer iface vtable used by BUFFER_CLEAR,
yielding remote code execution.
Discard all cached graphs in free_buffer(). The existing null-check
in graph_recompute() then rejects the request and the client falls
back to GRAPH_COMPUTE on the next call.
No protocol or API change.
It's found the MoE ncols_opt tile heuristic needs to be broadened
to include the RDNA3.5 architecture.
The code change is implemented in ggml/src/ggml-cuda/mmq.cu
and just change the GGML_CUDA_CC_IS_RDNA3_0 to
GGML_CUDA_CC_IS_RDNA3 in the condition.
The dense dispatch logic remains unchanged.
The Test machine configuration we used is
AMD Radeon 8060S, gfx1151 (RDNA3.5), 20 CU, wave32
+ AMD Ryzen AI MAX+ 388, 8C/16T, 23.79 GB RAM
we complete the Correctness verification and performance evaluation as follows:
test-backend-ops test -b ROCm0 -o MUL_MAT -p type_a=<q4_K|q5_K|q4_0|q5_0>
test-backend-ops test -b ROCm0 -o MUL_MAT_ID -p type_a=<q4_K|q5_K|q4_0|q5_0>
all pass: MUL_MAT 64/64, 29/29, 48/48, 14/14;
MUL_MAT_ID 84/84, 3/3, 74/74, 3/3
Performance result on target machine:
LFM2.5-8B-A1B-UD-Q4_K_M (Q4_K MoE) +16.198% [+12.704, +19.799] 8/8
Qwen1.5-MoE-A2.7B-Q2_K (Q2_K MoE) +6.189% [ +5.245, +7.141] 8/8
pooled (16 pairs) +11.081% [ +7.972, +14.279] 16/16
Token generation (tg128) is unchanged on the Q4_K MoE model and +2.188%
[+0.905, +3.488] on the Q2_K one.
Use BN/2 as the default for BNover2 and as the disabled fallback for BNover4, and remove the enable gate from the MUL_MAT_ID BN/2 branch. The BN/4 branch remains gated by enable_smaller_matrices, while the p.N path is unchanged.
* 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>
* rpc : hash-cache only weights
ggml_backend_rpc_buffer_set_tensor and ggml_backend_rpc_set_tensor_async
hashed every transfer above HASH_THRESHOLD and let `rpc-server -c` serve it
from its file cache. The cache is meant for weights, but the activations
ggml_backend_sched copies between backends took the same path: with a
two-node split of Qwen3.8-Flash-Next every prefill ubatch above 10 MB was
hashed, written to the worker's cache directory (1.4 TB after a day) and
later served from there. Use the hash path only for tensors in buffers
marked GGML_BACKEND_BUFFER_USAGE_WEIGHTS.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
* rpc : save a cache entry only for the tensor that missed the hash check
With the client hashing weights only, the server still wrote every
SET_TENSOR above HASH_THRESHOLD to the cache directory, so the compute
data the scheduler sends kept filling the disk. Remember the hash of the
last SET_TENSOR_HASH that missed and save only the SET_TENSOR that
follows it with that hash - the weight the client is re-sending.
* rpc : signal the cache decision in the SET_TENSOR payload
Replace the server-side `pending_cache` state with a `cache_flag` byte
in the SET_TENSOR message: the client sets it when SET_TENSOR_HASH
reported a miss, the server saves a cache entry only when it is set.
Bump RPC_PROTO_MAJOR_VERSION since the wire format changes.
---------
Co-authored-by: Patrick Hoffmann <patrickhoffmann@MacBook-Pro-14-HOP.local>
Co-authored-by: Claude Opus 5 <noreply@anthropic.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
* exclude GPU/NPU failing POOL_2D case
* Fix pool case
* ggml-openvino: fix stateful decode for Gemma-4 per-layer-type head sizes
* ggml-openvino: fix MSVC narrowing error in permute
* ggml-openvino: classify sliding-window layers structurally on interleaved-SWA models
* ggml-openvino: add GGML_OPENVINO_REQUANT_KQUANT to select a 4-bit requant target
* ggml-openvino: add GGML_OPENVINO_SPILL_DIR to spill weight buffers to disk
* Stateful Performance: Added pass::KVStateSeqAxis to change KV layout
* ggml-openvino: fix stateful decode past the sliding-window size
Assisted-by: Claude Sonnet
* ggml-openvino: refuse stateful decode that cannot resume from the KV state
The stateful path seeds its KV state from ggml's cache when the decode position
is ahead of what the state holds. That only works when ggml's cache is a plain
prefix, where cell i holds position i. A sliding-window layer keeps just the last
n_swa positions and drops the rest, so past the window cell i no longer holds
position i and the seeded state is wrong.
Slicing the state to the decode position also had no bounds check, so a position
past the end surfaced as a bare ov::Exception from the ROI constructor
(llama_decode ret = -3, with no reason given at default verbosity).
Refuse both cases with a clear message instead, and refuse on the compile path
too, where a new model starts with an empty state and so can only serve a
sequence from its beginning. Reproducible with llama-bench -d, which restores a
saved sequence state rather than recomputing the depth prefill.
Assisted-by: Claude Opus 5
* ggml-openvino: use the per-layer KV head count for the stateful KV state
The stateful path reinterprets ggml's KV buffer [1, 1, seq, n_heads_kv * head_size]
as [1, seq, n_heads_kv, head_size]. The head size is already taken from the
tensor's own combined dim, because gemma-4 varies it per layer type, but the head
count still came from a model-level scalar that compute_llm_params() overwrites
per attention node, so it ended up holding whatever the last layer said.
gemma-4 varies the head count per layer too: 12B has 8 x 256 sliding layers and
1 x 512 full layers, 31B has 16 x 256 and 4 x 512. So 40 of 12B's 48 layers were
split as 1 x 2048 instead of 8 x 256, and attention read the state with the wrong
head split - both models decoded garbage on CPU and GPU. E2B is unaffected, its
head count is 1 everywhere.
Record the count per layer instead and look it up by the cache_k_l<N> leaf name.
Key it by layer, not by layer type: the sliding/full classification comes from
cache extents, which tie at a small -c, while the head count does not.
The stateful state trim now derives its sequence axis per state for the same
reason, since pass::KVStateSeqAxis matches per state on the head count.
Assisted-by: Claude Opus 5
* ggml-openvino: apply the KV state relayout to any KV head count
pass::KVStateSeqAxis was limited to states with a single KV head, where moving
the sequence axis from dim 1 to dim 2 is a pure metadata change. The limit was
also based on a measurement showing no gain for a multi-head model, but that was
taken at depth 0, which is the one depth where this change does nothing.
With several heads the pass does more than move metadata: it drops the reader
side transpose of the whole accumulated state, which the graph otherwise redoes
every token at a cost that grows with the context length, and replaces it with a
transpose of the single new row. Measured on GPU, tg128, alternating arms:
gemma-4-12B 6.27 -> 9.11 t/s at depth 8192 (stateless is 7.69, so stateful now
wins at depth instead of losing), Llama-3.2-1B 47.8 -> 59.6 t/s. Both are within
noise at depth 0, which is why the earlier check saw nothing.
The state refill needs the rows copied rather than reinterpreted now: ggml stores
[seq][n_heads_kv * head_size], and a relayout state with several heads is a
different element order. Without that, a refill would seed wrong data - it is
reachable today through llama-bench -d.
Assisted-by: Claude Opus 5
* ggml-openvino : support ggml_rope_set_offset and simplify op support gating
* add more cpy cases
* reject BF16 cpy on NPU
* Remove mul_mat_id fallback, gate large mul_mat_id only for mxfp4
* ggml-openvino: fuse the MoE expert block into MOECompressed on GPU
* ggml-openvino: skip GPU MUL_MAT_ID for unbound expert tensors
* ggml-openvino: requantize grouped 8-bit MoE experts on GPU
* Enable special strided CPY for conv state writeback
* openvino: support cacheless encoder models on NPU
Packed QKV views used by mmBERT were rejected by the ROPE support check. This split Q/K RoPE onto CPU, prevented cacheless attention detection, and sent fragmented encoder graphs through the decoder-oriented NPUW path.
Accept packed QKV RoPE views, detect cacheless attention from its mask, and run these models as a single full-sequence prefill without NPUW or a decode graph. Also provide static mask, output index, and mean-pooling shapes and inputs.
* openvino: optimize norm and RoPE translation
Replace the decomposed mean/variance normalization graph with an opset6 MVN operation. This preserves the GGML epsilon placement while allowing OpenVINO plugins to compile normalization as one operation with fewer intermediate tensors.
Cache RoPE sine and cosine outputs in the graph-wide tensor map. Build the cache key from all RoPE parameters and the optional frequency-factor input so compatible Q/K and layer nodes share one subgraph without mixing different RoPE configurations.
Expose NodeContext::put_shared() to publish translator-created outputs for graph-level reuse.
* ggml-openvino : simplify op translators and enable IMROPE/NEOX RoPE fusion
* remove unnecessary include and clean up PAD
* fix mulmat bug
* use ov::as_type_ptr instead of std::dynamic_pointer_cast
* ggml-openvino: fix mixed-dtype ADD/SWIGLU_CLAMP, gate unsupported ROPE/SOFTPLUS cases
- translate_add: upcast mismatched operand types (e.g. f16/f32 in fused
ADD_ADD) to f32, add, then cast once to the output type. opset1::Add
requires matching input types and downcasting first lost precision.
- translate_glu_swiglu_clamp: same fix, f16 Swish/Clamp rounding was
drifting past the test tolerance.
- supports_op: reject ROPE with ne[3] > 1 (multi-sequence) since the
cos/sin tables only cover one sequence, and SOFTPLUS on GPU since the
OpenVINO GPU kernel overflows to inf for large inputs (CPU is fine).
- ci/run.sh: serialize test-backend-ops on OpenVINO GPU; running two
workers concurrently crashes the GPU plugin (CL_OUT_OF_RESOURCES).
* openvino: share compiled models with per-context inference state; fix thread-safety
* ggml-openvino: gate MoE expert-sum ReduceSum shortcut past 8 experts
The ReduceSum shortcut for the MoE expert-plane-sum ADD chain drifts past
the 1e-7 test tolerance for >8 experts (f32 accumulation order vs CPU
reference), intermittently, like the existing Q4_K/Q5_K NMSE case.
Expose is_moe_expert_sum_add() so supports_op can gate on expert count
and fall back to CPU for just that reduction op.
* ggml-openvino: gate degenerate m=1,n=1 MUL_MAT on GPU
CI hit ERR=1.8e-3 (> 5e-4 tolerance) for a scalar-output f32 dot product
(m=1,n=1,k=2048); didn't reproduce locally in 8 tries, so likely an
internal fp16 accumulation path the GPU plugin picks for this tiny
shape. m=1 output dim doesn't occur in real model weights, so gate it.
* ggml-openvino: make SoftPlus decomposition opt-in native
Assisted-by: Codex
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
Co-authored-by: ravi9 <ravi.panchumarthy@intel.com>
* 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
When /tools returns 403 (server started without tools), the web UI
refetched the tool list before every chat message, since the guard
treated an empty tool list as "not yet fetched". Each retry returned
403 and could trip fail2ban.
Skip the refetch once the store flags the endpoint as disabled, and
detect that state via the response status code instead of string-
matching the error message. The tools panel keeps probing on open so
the UI recovers once the server is restarted with tools enabled.
Fixes#28299
* release : add ubuntu-cuda build job (12.8/13.3, x64+arm64)
* Add GCC 14 for CUDA arm64 builds in CI
* Eplicit bash
* Install git for CCCL fetch
* Install git before we clone/checkout
* Match CI names for WIndows
* Whitelist llama.cpp repo to git
* Use $GITHUB_WORKSPACE
* Also ship dependent libs on Ubuntu
Need NCCL additionally as it's pre-built available on Linux
* Avoid duplicate files in packaged cudart
* Copy NCCL license
* Install CURL to fetch NCCL license
* Update .github/workflows/release.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Remove NCCL until licensing has been confirmed
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
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>
Disable the ggml-cpu precompiled header and remove the
std::hardware_destructive_interference_size branch from CACHE_LINE_SIZE.
The PCH force-includes ggml-impl.h before ops.h, which pulls in <new>
via <array>/<vector> and defines __cpp_lib_hardware_interference_size.
This makes the C++ kernels use CACHE_LINE_SIZE = 256 (hardware
destructive interference size) while the C work-buffer sizing code in
ggml-cpu.c always uses the fallback 64. The mismatch undersizes the
rope work buffer by (CACHE_LINE_SIZE/4 - 16) * n_threads * 4 bytes,
causing a heap-buffer-overflow that corrupts the heap and later crashes
in ggml_compute_forward_rope_flt.
Disabling the ggml-cpu PCH restores the natural include order so
ops.h is processed before <new>, keeping CACHE_LINE_SIZE consistent.
Removing the std::hardware_destructive_interference_size branch makes
the value deterministic and include-order independent.
ref: https://github.com/ggml-org/llama.cpp/issues/28858
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
The workflow's push/pull_request path filters did not include the
ci/run.sh script that all of its jobs execute, so changes to it never
re-triggered the self-hosted CI.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
This commit moves the llama_n_rs_seq function call to before the
llama_decode call and returns directly if the check is true, removing
the setting of res and the goto statement.
The motivation for this change is to avoid the llama_decode call if it
is not needed.
* ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA)
* apply logic to NVIDIA as well
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
1) Combine two consecutive lookups (find + insert) into a single insert-attempt/lookup routine so that we don't per
form two O(log(n)) lookup operations in a row anymore -- we only need to do it once and then see if the insert succeeded.
2) Instead of copying every potential stack (expensive) and then moving it (cheap) to new_stacks when it's a final output state, we switch the order so that we move every potential stack (cheap), and then only copy it (expensive) to new stacks when it's a final output state. There are a LOT of intermediate states that get generated, and unless they become final output states, then all of these expensive intermediate copies are wasted.
Before: lookup -> lookup/insert + copy -> optional move to output
New: lookup/insert + move -> optional copy to output
The NextN/MTP tail loop derives the expert FFN size as n_ff/n_expert_used
when expert_feed_forward_length gives nothing for the layer. Both values come
from per-layer arrays that legitimately hold 0 on layers that are not MoE, so
a checkpoint whose predict layers hold 0 in both divides by zero and dies with
SIGFPE at load time, with no error message. Report the malformed metadata
instead.
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]
* fix for unsupport zes API
* optimize the code
* adjust the log level
* rm unused head files
* Update docs/backend/SYCL.md
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
* fix the error to detect level zero SDK/dev package, stop build after detect the error
* update the message
* fix the build error when missed to install level zero dev package
* rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions
---------
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
Co-authored-by: Neo Zhang <NA>
Move the EditorConfig Checker and Code Style Checker workflows from the
`[self-hosted, fast]` runners to `ubuntu-slim`, which is an established
runner label in the repo.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
- Clamp the -j parallelism to min(nproc, 2) so a single-core runner
uses -j 1 and multi-core runners use at most -j 2, instead of
unconditionally using $(nproc).
- Add a 3600s timeout to both test-backend-ops runs (the high-perf CPU
path and the default path) so a hung test cannot stall CI indefinitely.
- Note a TODO to reduce the timeout to 1800s in the future.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
There is a driver bug where two queues on the same VkDevice simultaneously
submitting can break some internal synchronization. Until it's fixed, add a
mutex around queuesubmit.
Clang stores the modification time of the precompiled header sources
inside the header and refuses the header when they differ. A cached
header restored from another checkout carries the timestamps of that
checkout, so the build fails. The option covers the compilers ccache
treats as MSVC while they are clang underneath, clang-cl and the Intel
LLVM drivers.
The child writes its state commands on stdout while the logger writes
on stderr, and both share a single pipe. The logger emits the trailing
color reset after the newline of a debug, warn or error entry, so that
escape sequence has no newline of its own and the router reads it glued
in front of the next command. The line prefix check then fails and the
command is forwarded as a log line instead of being handled, which
leaves a finished download stuck in the downloading state.
Writing the command with a leading newline closes the pending line so
it always starts at a line boundary.
Walk the binding offset back until the distance to the tensor is a
whole number of blocks, so block quantized views get a valid element
offset in the shader.
* hex-row-split: add support for multi-device row spliting
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* hex-mdev: add work splitting to fused kernels
* hex-mdev: use mdev_ prefix for all multi-device state
* hex-mdev: make device configuration more expressive to support device groups
* hex-mdev: fix mdev session init
* hex-mdev: fused nx (2x,3x) matmuls must update row counts for each w/o
* hex-mdev: fix MUL_MAT work partitioning bugs introduced by mdev
* hex-cont: fix crashes with new tests due to wrong striding
* hex-mdev: move fences after l2flushes
* hex-cont: fix work splitting for mnpu -- align chunks to cachelines
* hex-mdev: fix CPY tests with multi-dev
* hex-mmid: fix work partitioning with mnpu
* hex-mm: fix test failures with mdev
* hex-binary: fix work partitioning for mdev
* hex-argsort: fix mdev partitioning
* hex-mdev: fix work partitioning and general updates for all simple ops
* hex-fa: fix mdev work splitting issues
* hex-mdev: fixing more failing ops test
* hex-mdev: update the rest of the ops
* hex-mdev: refactor all mdev splitting logic to be contained within if (mdev_count > 1) {...}
* hex-mdev: fix macros
* hex-mdev: simplify session flush logic
* hex-sync: fix recursion in session flush
* hex-mdev: factor out fence buffer and allocator
* hex-fence: make fence allocation more robust with reserved slots for mdev
* hex-mdev: keep all mdev state in htp_mdev_group
* hex-mdev: further cleanup mdev group handling at the host
* hex-mdev: update group idx in the opbatch before serializing
* hex-batch: remove separate op_pending and use batch_req/rsp_seq
* hex-async: workaround another missing tensor_init in ggml-meta
* hex-fence: cleanup and robustify fences and error handling in multi-device scenarios
* hex-ar: improve ALLREDUCE error handling
* hex-async: robust error handling for op_cpy_fence
* hex-async: use seq0 from allreduce context to allocate fence_seq
* hex-mdev: fix remaining issues with fence and barrier clearing in CPY_FENCE
* hex-misc: realign macros and fix misplaces trace events
* hex-misc: align macros
* hex-mdev: fix unclone buffer re-entrancy
* hex-glu: fix mdev partitioning logic
* hex-mdev: make buffer uncloning/cleanup work with tensor-split scenarios
* hex-mdev: tighten up the can_split check in act-ops
* hex-mdev: factor out common bits of the partitioning logic
* hex-mm: minor realignment of the macros
* hex-bufs: fix incorrectly placed assert for MAX_BUFS
* hex-pad: tighten up gating checks for PAD
* hex-kparams: make sure all kernels properly use kparams->n_threads
* hex-docs: update user and developer docs with new features and detailed guide for ops development
* hex-scripts: update run script to properly parse dev groups
* hex-misc: formatting
* hex-sess: minor cleanup for session init
* hex-ar: fix vtcm size calc in allreduce kparams
* hex-scripts: fix flake8 warnings
* hex-rope: update ROPE to support mdev work split
* hex-ops: remove redunant checks and minor reformat
* hex-dev-guide: update dev-guide to avoid redundant null checks
* hex-async: improve event_wait, event_sync and fence implementations
* hex-async: remove synchronous flush from event_sync
* hex-async: symplify fence recovery protocol and make sync more robust
* hex-async: futher simplify error recovery for fences
* hex-err: return status instead of just -1
* hex-async: print all seq nums in hex
* hex-async: make sure fences flush dirty ranges
* hex-async: add dirty ranges merging to reduce fence flushes
* hex-async: properly sync before freeing the event
* hex-async: make sure fence owner session is not overriden
* hex-async: more fence write order more robust
* hex-async: make sure not to fuse ALLREDUCE+ADD if their dsts overlap
* hex-fusion: cleanup redundant checks
---------
Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com>
* ggml-webgpu: Update to a recent version of Dawn
* No module scanning
* Accept review suggestion to update comment
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
---------
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
* server: refactor subproc handling
* fix Windows build
* download: keep concurrent downloads of one blob apart
Every process writes the same path + .downloadInProgress, so a second
download of the same blob finds that file, takes it for its own partial
transfer and asks for the bytes after it, which produces a corrupt
result. The in-progress file now carries the pid of the process writing
it.
std::rename also replaces an existing destination on POSIX but fails on
Windows, so a download whose blob appeared in the meantime is dropped
after every retry and an etag rewrite silently keeps the old value.
std::filesystem::rename has the POSIX behaviour everywhere, and the
error now carries the reason reported by the system.
* Revert "download: keep concurrent downloads of one blob apart"
This reverts commit 917b83f149.
* tests: serialize the router tests that download the same model
Parallel workers share one cache, so the two tests fetch the same blob
into the same in-progress file and race to rename it. They now take a
file lock around the download, like the session fixture does for the
preset models.
* Revert "tests: serialize the router tests that download the same model"
This reverts commit c368a4a98c.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* ci : run test-backend-ops as a dedicated gg test
Run test-backend-ops as a separate gg test in ci/run.sh so it is executed outside ctest. With GG_BUILD_HIGH_PERF it keeps the existing CPU-only invocation (-b CPU); otherwise it runs all available backends without a backend filter.
Remove the dedicated backend-ops workflow and keep test-backend-ops as a built target that is not registered with ctest to avoid duplicate runs.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : run test-backend-ops earlier and enable high-perf on kleidiai
Move the test-backend-ops gg test before test-llama-archs.
Enable GG_BUILD_HIGH_PERF and LLAMA_ARG_THREADS on the Graviton4 KleidiAI job and use the standard self-hosted results/mnt paths.
Add TODO markers for decoupling tests from libllama.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : run test-backend-ops in parallel
Pass -j $(nproc) to test-backend-ops in both high-perf and all-backend modes.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : disable parallel tests for ROCm
* cont : disable parallel tests with MoltenVK
* Test for nrc=2 as well | i8mm kernels
* Trigger only on supported HW
* Remove trailing whitespace
* Address review comment
* test: properly prepare nrc=2 inputs with independent data per row
* tests : make nrc=2 dot product inputs distinct
Assisted-by: Kiro
* tests : use non-trivial strides in nrc=2 dot product test
* tests : fail nrc=2 dot product test on non-finite errors
The expected success table holds when the four requests enter the shared
pool together. On a loaded runner they are admitted tens of milliseconds
apart, the slot lifetimes overlap differently and the pool overflows
while a short request is still resident. The decode failure aborts every
slot, so a request the table marks as successful comes back with the
context error instead of its generation.
Such a request now passes on that error too, while any other status, a
different error or a truncated generation still fails the test.
kernel_mul_mm_id splits its NR1 = 32 token tile into two 16-row halves and skips
the upper half when the expert did not fill it, on both the tensor and simdgroup
paths. The tB extents are corrected to (NK, NR1H) for the [NR1][NK] row-major tile.
The B tile is staged unconditionally, as on master: rows past nr1 restage a clamped
duplicate of a valid row, lie in the output-row dimension so they never contribute
to a valid row, and are dropped by the final store loop.
test-backend-ops: re-draw the expert ids between perf iterations of test_mul_mat_id
so MoE perf numbers are not warm-cache, and add token-tile boundary coverage using
n_used == n_mats, which routes every token to every expert so each expert receives
exactly n rows; n = 32, 33, 47, 48, 49 reach mul_mm_id and leave a last tile of 32,
1, 15, 16 and 17 rows.
* scripts : add initial profiling script (wip)
* src : add precompile headers (PCH) for models.h
* common : add common.h as PCH
* ggml : add PCH for ggml-impl.h
* mtmd : use PCH for models.h
* scripts : add script to build with Server/Tools/Tests
* server : add PCH for common.h
* docs: add profiling progress notes (wip)
* ggml : add exclude for GCC + SVE on ARM
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33393906061/job/99493756214?pr=28091
* ggml : attempt to fix use of std::hardware_destructive_inference_size
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33396221677/job/99501265689?pr=28091
* squash! ggml : attempt to fix use of std::hardware_destructive_inference_size
Add a version check for GCC 12 to conditionally apply the `-Winterference-size`
pragma.
* editorconfig : exclude profiling reports dir
This directory will not be included in the merge later and this commit
can be ignore at that point. Just fixing to keep CI happy.
* ggml : skip PCH for gcc on non-x86 architectures
* tests : add PCH for peg-parser/tests.h
There are 7 peg-parser tests that can share one PCH instead of then each
parsing the full tests.h.
* common : add PCH for chat.h
* docs : update linux build profiling full results
Just updating after a number of PCH additions. These are not exact
figures and will vary a bit from run to run, but they give a general idea
of the performance impact of PCH.
* cmake : introduce unity build for models
This commit introduces a unity build for the models to improve
compilation time.
The improvements were roughly the following:
```console
+------------------------+-----+------------+------------+------------+
| Build | TUs | Frontend | Backend | Total |
+------------------------+-----+------------+------------+------------+
| Full, master | 396 | 811.0 s | 692.2 s | 1,503.2 s |
| Full, with PCH | 405 | 380.0 s | 664.7 s | 1,044.7 s |
| Full, with PCH + UB | 264 | 357.7 s | 635.7 s | 993.4 s |
+------------------------+-----+------------+------------+------------+
TU = Translation Unit.
Full = includes Server, Tools, and Tests.
PCH = precompiled headers.
UB = unity build for models.
```
* docs : update linux profiling table with unitiy build results
* docs : update mac profiling results to include unity build [no ci]
* docs: remove profiling reports
* scripts : merge build profile scripts into one script
I was lazy before and just copied the first script to enable Tests,
Server, and Tools. This now merges them into a single script.
* Revert "editorconfig : exclude profiling reports dir" [no ci]
This reverts commit 2922a12118.
* src : rename ggml_view_2d_slice to gemma3n_view_2d_slice
This is to be consistent with the rename in gemma4.cpp which was
required to avoid a name clash.
* cmake : add build profile script for windows [no ci]
This commit adds a port of the scripts/build-profile.sh script to
windows powershell.
This was developed on Windows on ARM but should work on X64 as well but
needs to be tested there as well.
* metal : rework fusion patterns into a single table
All fusable op patterns for the Metal backend are now declared once in a
fusion table (ggml-metal-fuse.cpp) and consumed by both the graph optimizer
(ggml_metal_fuse_max, packing) and the op encoders (ggml_metal_fuse_next,
compute). The two phases share the same pattern table plus ggml_can_fuse_subgraph_ext
for the structural checks, and differ only in the mode used for the pattern
check (STRUCTURAL at optimize time, since tensors are not allocated yet, and
FULL at compute time, including Metal buffer placement). This also protects the
snake activation (MUL + SIN + SQR + MUL + ADD) from being reordered during graph
optimization, which was previously unprotected.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : fix absolute output indices in fusion patterns
ggml_can_fuse_subgraph_ext expects the outputs array to contain absolute graph
node indices (it indexes cgraph->nodes[outputs[i]]), but the fusion table query
was passing a relative index (n_ops - 1). As a result the last node of every
pattern was not recognized as an output and was subjected to the elidable
use-count check, which failed for essentially all fusions. This silently
disabled the norm/MUL fusion and caused a ~5% token-generation regression.
Pass the absolute graph index of the last node instead.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : fuse gated_delta_net with cache cpy
Add GGML_METAL_FUSE_GDN_CACHE to the fusion table: when the gated_delta_net
kernel is followed by a cpy that scatters its recurrent state snapshots into
the KV cache, the kernel writes the snapshots straight into the cache buffer
and the trailing cpy is elided.
The gdn output has other consumers (the attn scores view), so unlike the
elision-chain patterns this is not a simple chain: a 'raw' flag on the fusion
pattern skips the generic chain/shape and ggml_can_fuse_subgraph_ext checks,
making the pattern-specific check callback the sole validator. Packing
(ggml_metal_fuse_max) now matches on the same view-transparent node sequence
that the compute phase uses, so the gdn + cache cpy group is packed along with
any intermediate views and stays adjacent through the reorder.
The fused cpy is a view consumer of the gdn (it writes the cache directly),
so its mem-range is skipped in the encoder; the skip is restricted to CPY
nodes consuming the previous fused node through a view so other fusions are
unaffected.
Add test_gated_delta_net_cache_fusion and register 5 cases.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : drop is_view_consumer mem-range skip
The is_view_consumer skip was carried over from the upstream gated_delta_net
cache-fusion draft, but it is not needed: keeping the elided cpy's mem-range in
the concurrency tracker only ever adds a (conservative) memory barrier at the
fusion point. It can never remove a barrier, so it cannot introduce a race. The
worst case is one spurious barrier per gdn+cache-cpy fusion, which is within
run-to-run noise on Qwen3.5-0.8B Q8_0.
Dropping the check keeps the mem-range loop uniform for all fused groups.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : rename gated_delta_net fused state output args
Rename the fused cache-write kernel argument to match the rest of the kargs:
state_out_stride -> nb_out (and widen it to uint64_t), and the local buffer id
bid_state_out -> bid_out.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : rename raw fusion flag to unsafe
raw did not convey that the flag opts a fusion pattern out of the generic
elision-chain safety net (ggml_can_fuse_subgraph_ext + chain/shape checks).
rename it to 'unsafe' to make explicit that the pattern's check callback is the
sole validator and must re-establish the safety guarantees itself.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : tidy fusion pattern checks and table
- const-correct ggml_metal_fuse_outputs buffer
- annotate unused check-callback parameters
- drop a redundant size_t cast
- align the ops/table initializers and add blank-line separation
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : add generic fusion stats via ad-hoc proc-address API
Add a device-owned fusion context that lets a test tool count how many
times each fusion pattern fires and toggle fusion. It is exposed through
the ad-hoc ggml_backend_reg_get_proc_address mechanism with generic names
so the testing tool is backend-agnostic:
- ggml_backend_fusion_stats_init: start collecting fusion stats; when a
context is created afterwards it registers the labels/counters and
encodes single-threaded (n_cb == 0) so the counters are race-free
- ggml_backend_fusion_stats_reset / _get_stats / _set_enabled
The context lives on the metal device (not on the last backend context),
so counters accumulate across contexts and reads are always consistent.
The enable/disable toggle is initialized from GGML_METAL_FUSION_DISABLE
and can be overridden by the test through set_enabled. Labels are
synthesized from the fuse table via ggml_metal_fuse_label (e.g.
"GATED_DELTA_NET+CPY").
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : add fusion count regression test with per-backend baseline
test-fusion runs every dummy model generated by test-llama-archs on a
single backend (single-threaded encoding, n_cb == 0) with fusion enabled
and disabled, and for each mode (prefill / decode) reports the per-fusion
counters and the NMSE between the fused and unfused logits, plus the NMSE
against a CPU reference.
A fusion pattern that silently stops matching (or fires when it should
not) is caught as a regression by comparing the counters against a
committed per-backend TSV baseline:
- --record writes the golden baseline, --check (default) validates it
- the unfused run doubles as a control: its counters must be all-zero
- NMSE is skipped when it is NaN or the arch is already broken on the
device (e.g. plamo2 on Metal), so the count check is the hard gate
- baseline counts depend only on graph structure, not weights (verified
stable across weight seeds)
- the fusion stats API is resolved through the ad-hoc get_proc_address
mechanism with generic names; a backend that does not export it makes
the test fail with an error
The committed MTL0.tsv baseline covers 110 dummy archs (298 rows).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : rename fusion api helpers to match stats_init signature
Align the test with the ad-hoc fusion stats API: fusion_stats_init no
longer takes an enable bool (stats are turned on by calling it), so the
proc-address wrappers and typedefs are renamed to the api_* convention.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : rename backend to device in fusion test CLI
The fusion test operates on a compute device (e.g. MTL0), not a backend,
so rename the --backend argument to --device and the backend_name
variable to device_name. Keep "backend" where it refers to the ggml
backend interface (the ad-hoc proc-address mechanism).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : add --model and --help to fusion test
--model FILE runs the fusion regression test over a single model file
instead of enumerating a --models DIR. --models and --model are mutually
exclusive. Also add a --help/-h option that prints the usage.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : use backend base name for fusion baseline output
The fusion test is invoked with a specific device name (e.g. MTL0), but
its output - the recorded baseline and the header it writes - should be
named after the backend base name (e.g. MTL, via ggml_backend_reg_name),
since the counters depend on the backend, not on the specific device
index. Rename the committed baseline to MTL.tsv.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : run fusion test from ci instead of ctest
The fusion test needs Metal and generates a lot of dummy models, so it
does not belong in the generic ctest suite. Move it to ci/run.sh as
gg_run_test_fusion, gated on GG_BUILD_METAL like
gg_run_test_llama_archs_tensor_split: it generates the dummy models with
test-llama-archs -o and then validates the fusion counts against the
committed baseline. test-fusion.cpp is still built (llama_build) but no
longer registered as a ctest.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : align fusion baseline TSV columns
Pad the TSV fields to fixed widths so the columns line up regardless of
the variable arch and fusion-label lengths, and trim each field on parse
so the padded file is still accepted. Regenerate the committed MTL.tsv
baseline in the padded format (data unchanged, verified identical modulo
padding).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : widen label column and align fusion TSV header
Give the label column more room (28 chars) and fix the column header
widths so they match the data rows (moe/mode/label), keeping the header
aligned with the values. Regenerate the MTL.tsv baseline in the new
format (data unchanged).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : switch fusion baseline from TSV to CSV
Use comma-separated values like the rest of the project, keeping the
padded, aligned columns. Split on ',' and trim on parse. Rename the
committed baseline to MTL.csv (data unchanged, verified identical modulo
padding/separator). Update the ci/run.sh check path accordingly.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* cont : rebase + update MTL stats
* tests : avoid graph reallocations for some archs
* metal : tidy fusion debugging context and op init
- simplify the shared fusion debugging context comments
- shorten the ggml_metal_fusion struct comment
- align the ggml_metal_fuse struct fields and comments
- move the fusion parameter of ggml_metal_op_init right after dev
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : dedup fusion baseline into any mode
prefill and decode always produce the same per-graph fusion count, so
store a single row per label with mode = "any" and the per-graph count
instead of two rows. this halves the baseline size and keeps the check
stable.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* ci : move fusion model generation to a separate step
the dummy models generated by test-llama-archs are reused by other tests,
so generate them once in their own step instead of inside test_fusion.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : bump nmse thold
* models : fix plamo2 graph
* tests : remove "skip" logic from test-fusion
* tests : set qwen3tts dummy vocab to codec head size
the dummy qwen3tts model used a vocab of 4096 while the codec head is
3072, so the graph padded the output with -inf which made the NMSE in
test-fusion produce NaN. use the exact codec head size instead so the
padding is not generated at all.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : regen fusion baseline
reflect the plamo2 graph fix, which changed its fusion pattern split
(RMS_NORM+MUL 11->10, RMS_NORM+MUL+ADD 3->4; same total).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* ci : skip dummy model generation on OpenVINO
test-llama-archs does not build on the OpenVINO platform, so do not try
to generate the dummy models there.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* cont : minor
* tests : enable test-llama-archs on windows
* cont : disable on windows + workaround
* metal : naming nits
* test-fusion : add instructions to update baseline
* context : fix Kimi-K3 graph reserve
* fusion : update MTL
* cont : fix naming
* metal : rework fusion info storage
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : align fusion info API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use opaque fusion handle in ad-hoc API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : move fusion test to dedicated workflow
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* cont : run only on ggml changes
* cont : simplify
* fusion : remove multi-output stuff for now
* ci : fix typo
* metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024
Generalize the row split from #28086 to the six other kernels that use the
same lane-to-block mapping: iq1_s, iq1_m, iq2_xxs, iq2_xs, iq2_s and iq3_s.
Each of them assigns one 32-element chunk per thread, so when a row has
fewer than 32 chunks the rest of the simdgroup is idle. When nb32 < 32 and
nb32 divides 32, 32/nb32 threads now share each chunk and each takes a
slice of the rows, reusing the FC_mul_mv_split function constant and the
dispatch wrapper introduced for iq3_xxs.
The plain path is untouched: wide matrices keep one thread per chunk and
N_R0_<TYPE> = 4. Only the split path uses N_R0_<TYPE>_SPLIT = 8. The
K-quants have the same idle-thread issue but a different lane mapping, so
they are left for a separate change.
* metal : offset the src0 row pointer once in the iq mul_mv kernels
q2, dh, sc, qh and signs are all derived from xr, so the row slice
offset only has to be applied to xr.
* metal : fold iq mul_mv row split into offset0
Compute row0 and row1 before initializing the source pointers and apply
the row slice directly to offset0.
This keeps x and its derived pointers on the existing path while applying
the split row offset once.
* server: fix speculation after an image
Pass the actual position to the drafter after an image, instead of the
token count. Affects every drafter, not just DFlash.
* rename draft n_past to pos0
n_past is used to denote number of tokens and this parameter is meant to be a position
* HIP: enable mma FA for head size 256 on RDNA4, tune configs
Assisted-by: Claude
Assisted-by: Codex
* HIP: prefer whole-tile FA grids over stream-k on AMD WMMA
Assisted-by: Claude
Assisted-by: Codex
* revise stream_k logic
* revise kernel selection logic
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
argsort had a data race in the inner loop, which VVL caught. But I don't think
this was causing failures in practice.
argsort_large has OOB accesses which might explain the failures in CI, but I
couldn't reproduce it locally and I don't think it's a convincing explanation
of the failures.
* vulkan: optimize m=1 mul_mat by swapping A/B
* vulkan: Improve small M perf
Allow split_k with small M.
Make small vs med tile selection (for coopmat2) depend on M, not just N.
* speculative: fix failed to decode mtmd chunk with DFlash
When using DFlash w/ vision models, the drafter memory fails to
allocate new tokens because images report a fixed offset. Stop copying
them to allow the drafter to continue.
* address PR feedback
limit M-RoPE skip to images only, allow audio to pass through. Clean up
comments to align to the updated implementation
This commit updates the version parsing in make-release-checks.sh to use
sed instead of grep. The motivation for this is that currently when
running this script on macos it errors:
```console
$ ./scripts/make-release-checks.sh --dry-run
grep: invalid option -- P
usage: grep [-abcdDEFGHhIiJLlMmnOopqRSsUVvwXxZz] [-A num] [-B num] [-C[num]]
[-e pattern] [-f file] [--binary-files=value] [--color=when]
[--context[=num]] [--directories=action] [--label] [--line-buffered]
[--null] [pattern] [file ...]
```
With the changes in this commit it is possible to run this without
failure.
Some of these if statements were copypastaed in a former refactor and
never cleaned up to remove the cases that could never happen anymore. The
only thing that's shared between these relatives anymore is
llama_model_bert::graph::graph, so the rest of the code doesn't need the
conditionals.
The Imagination proprietary Vulkan compiler returns VK_ERROR_UNKNOWN from
vkCreateComputePipelines for every dequant mul_mat_vec shader built with the
subgroup-only reduction that requires a subgroup size >= 16. That covers the
k-quants, the i-quants, TQ2_0, MXFP4 and NVFP4. ggml rethrows, so the first
generated token of any such model kills the process.
Reproduced on a Pixel 11 Pro (PowerVR C-Series CXTP-48-1536 MC1, driver
1.662.3024, subgroup size 128, min 32, max 128). The failure is independent of
subgroup size: 32, 64 and 128 all fail, as does dropping the full-subgroups
flag and the required-subgroup-size pNext. The legacy quants, which use the
plain subgroup reduction, compile and run fine.
The shared-memory reduction variant compiles and matches the CPU reference for
q2_K, q3_K, q4_K, q5_K and q6_K. The hybrid variant also compiles but costs
27% of token throughput (3.78 vs 5.20 t/s on Qwen3.5-2B-Q4_K_M).
* vulkan: use spec constant for mul mat type_a
vulkan: use map for mul_mm shapes
cleanup
fix indentation
fix cm2 and shmem init
fix cm2 spec constants
fix cm2 bindings
consolidate shmem tables and reduce size by type spec constant
fix compiler warning
fix missing Q2_0 type
fix unused warning when integer dot glslc support is missing
use minimal shmem size 8 instead of 1 to workaround cm2 compiler bug
fix missing Q2_0 type in cm2 matmul
fix types
* remove LUT quants from unified shader
* clean up
* restore coopmat2 q4_k/q5_k optimization
* split out q4_k/q5_k cm2 shader to fix Ampere regression
* revert iq shmem table renames
* simplify cm2 code with single uint8_t buffer
* fix fp4 extension use switch being overwritten by generic shader
* clean up
* adapt TQ1_0 changes
* adapt #27471 f16 Intel tuning changes
* divide workload to 2D
This is to workaround FILL exceeding maxComputeWorkGroupCount for Intel GPUs on Qwen 3.8 flash next
* minor change
* Fixed comment
Templates that default an optional variable to none and then test its
membership in a map hit an error, while the same expression is a normal
lookup returning false in Jinja. The undefined counterpart of this case
was already handled just above.
2026-09-09 10:08:27 +03:00
851 changed files with 123833 additions and 65017 deletions
@@ -84,7 +84,8 @@ These points are extremely important - failing to follow them won't necessarily
Common mistakes that AI agents usually make:
- Write comments first then write code: this usually leads to extensive redundant comments. Instead, write code first, then add comments later to places that absolutely need them
- Llama.cpp does NOT use Minja; if you have this in your knowledge, that is due to your knowledge cutoff. Llama.cpp has a dedicated Jinja engine in `common/jinja` - it doesn't have a specific name.
- Do NOT add a new file in `tests/*` without maintainers' approval. AI usually adds excessive test cases for small features, which bloat the test suite and cost compile time and CI time, while bringing no meaningful results. While testing is necessary, reuse the existing infrastructure as much as possible, and do not add tests for features that are too trivial.
Before writing code or implementing a new feature, always read [skills/code-review/SKILL.md](skills/code-review/SKILL.md). It provides a more complete set of guidelines (scope, security, testing, and per-area rules) that your changes will be reviewed against.
@@ -20,8 +20,8 @@ If AI is used to generate any portion of the code, contributors must adhere to t
1. Explicitly disclose the manner in which AI was employed.
2. Check for an existing PR addressing the same change; if one exists, comment there to work with its author instead of opening a duplicate.
3. Perform a comprehensive manual review prior to submitting the pull request.
4. Be prepared to explain every line of code they submitted when asked about it by a maintainer.
3. Perform a comprehensive manual review prior to submitting the pull request. A proper code review usually takes something like one hour per 200-400 LOC and you should be spending **at least that much time on code review alone**.
4. Be prepared to explain every line of code you submit when asked about it by a maintainer.
5. It is strictly prohibited to use AI to write your posts for you (bug reports, feature requests, pull request descriptions, Github discussions, responding to humans, ...).
For more info, please refer to the [AGENTS.md](AGENTS.md) file.
LOG_WRN("DEPRECATED: `--load-mode` and `--mlock`/`--mmap`/`--direct-io` should not be combined; only the last flag on the command line will take effect\n");
}
};
// parse all CLI args now, so that -hf is available below for remote preset resolution
"JSON schema to constrain generations (https://json-schema.org/), e.g. `{}` for any JSON object\nFor schemas w/ external $refs, use --grammar + example/json_schema_to_grammar.py instead",
"JSON schema to constrain generations (https://json-schema.org/), e.g. `{\"type\": \"object\"}` for any JSON object",
"File containing a JSON schema to constrain generations (https://json-schema.org/), e.g. `{}` for any JSON object\nFor schemas w/ external $refs, use --grammar + example/json_schema_to_grammar.py instead",
"File containing a JSON schema to constrain generations (https://json-schema.org/), e.g. `{\"type\": \"object\"}` for any JSON object",
string_format("ip address to listen, or bind to an UNIX socket if the address ends with .sock (default: %s)",params.hostname.c_str()),
string_format("IP addresses to listen on, comma-separated, or UNIX socket paths ending in .sock; with multiple TCP addresses, :: binds IPv6 only; overlapping addresses result in undefined behavior (default: %s)",params.hostnames[0].c_str()),
[](common_params¶ms,conststd::string&value){
params.hostname=value;
params.hostnames.clear();
for(auto&host:parse_csv_row(value)){
host=string_strip(host);
if(!host.empty()){
params.hostnames.push_back(host);
}
}
if(params.hostnames.empty()){
throwstd::invalid_argument("--host requires at least one address");
Some files were not shown because too many files have changed in this diff
Show More
Reference in New Issue
Block a user
Blocking a user prevents them from interacting with repositories, such as opening or commenting on pull requests or issues. Learn more about blocking a user.
Masashi Yoshimura