* llama-quant : fix quantizing 4D tensors
nrows_total was ne[1]*ne[2], so for a 4D tensor only the first ne[2]
slices were quantized and the rest of the rows were dropped. Use
ggml_nrows(tensor) to count all rows, including ne[3].
Fixes#29916
* llama-quant : detect expert tensors by name and fix imatrix size for 4D tensors
The MXFP4_MOE path used ne[2] > 1 to detect MoE expert tensors. This is
wrong for non-expert tensors with 3 or more dims, e.g. MLA projections
(BailingMoE3) and 4D conv kernels. Use the tensor name instead: all
tensors used by GGML_OP_MUL_MAT_ID are named *_exps or *_chexps. This
also removes the BailingMoE3 special case.
The imatrix holds one vector of ne[0] values per 2D slice, so the
expected size is ne[0]*ne[2]*ne[3], not ne[0]*ne[2]. This matches how
the imatrix is indexed per row for 4D tensors in llama_tensor_quantize_impl.
Assisted-by: Claude Opus 5.5
* cuda: fuse FFN gate/up and GLU in MMQ with interleaved gate/up rows
For prompt processing, run the gate and up matmuls of a dense FFN as one
mul_mat_q launch and apply the GLU in the write-back.
- The loaders take the gate weights as a second pointer and fill the weight
tile with up and gate rows interleaved in chunks of 8, so each thread holds
an up value and the matching gate value (rows r and r + 8 of an MMA tile).
- The MMA loop is unchanged: same tile, warp layout, vec_dot and one sum[]
per thread. The fused write-back computes the GLU from registers and writes
half the rows.
- The fused path uses the regular tiling (no stream-k) with the column tiles
innermost so that neighbouring blocks share weight tiles.
- Enabled for Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS on NVIDIA Turing and newer.
* cuda: assert the fused MMQ gate argument and trim redundant comments
* cuda: move the fused MMQ hint check to should_fuse and add TODOs
* cuda: pass prec_src1 to the fused MMQ write-back
* cuda: fuse only with efficient tiling and no fallback, merge the GLU write-back
* cuda: require has_fusion explicitly, drop the fusion-only src1 precision check
* cuda: launch the fused MMQ kernel through the stream-k kernel with one block per tile
* cuda: inline the MMQ tiling-efficiency check
* tests: one fused gate/up MMQ case per type and GLU op
* cuda: always launch the fused MMQ kernel with one block per tile
* tests: shorten the fused gate/up MMQ test comment
* cuda: select the fused MMQ loader and write-back in ggml_cuda_mmq_get_util_funcs
* cuda: run the fused MMQ kernel on the plain tiling path
* cuda: put the column tiles innermost in the MMQ plain tiling path
* vulkan: Add a larger tile size for conv3d (coopmat2 only)
This adds a larger tile size, in practice it is only for coopmat2 where
we don't need shared memory to store the output (would use more than 64KB).
Add some additional representative backend tests. At least, this is what codex
came up with from looking at sd.cpp Wan VAE.
* only apply large tile size for newer driver versions
WM of the small tiles was 32, or 8 for 8-wide subgroups, whatever the
tile's BLOCK_SIZE and WARP. With 16-wide subgroups (Mali) the second
subgroup of a workgroup wrote into the next tile's columns. With 8-wide
subgroups some _id tiles left rows uncomputed. Derive WM so there is one
warp tile per subgroup.
Tested on my Pixel 9 Android phone. Before the change there were 20
MUL_MAT and 6 MUL_MAT_ID failures. After the change they all pass.
Fixes#28637
Assisted-by: Claude Opus 5.5
* hex-tools: support for --compare option in the profile.py
* hex-unary: improve collapse logic to also do splits if nrows is small
* hex-colapse: handle all colapsing at the host and make it generic
* hex-colapse: remove leftover unary helper
* chat : reject JSON nested more than 128 levels before walking it
The schema builder, the jinja engine and the decision code walk request
JSON with recursion. A request nested deep enough overflows the stack of
the HTTP thread and llama-server dies (SIGBUS on macOS, SIGSEGV on Linux):
for a tool schema from a few hundred levels on macOS and about 10000 on
Linux.
Check the depth first and answer 400 past 128 levels: the chat inputs in
common_chat_templates_apply (tools, tool-call arguments,
chat_template_kwargs, json_schema), the json_schema of /completion, and the
whole /v1/systemone body. A gemma4 tool result that is too deep stays text.
* fix the root, not the call side
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* model : add DeepSeek-V4.1-Flash (deepseek41)
* deepseek41 : remove MXFP8, convert the Engram table to BF16
* deepseek41 : compute the compressed-KV read layers once in the cache
* deepseek41 : use get_cells() for the shared-history check, drop llama_kv_cache::seq_pos_tok_le
* deepseek41 : gather the Engram rows in-graph like qwen4exp, drop the host gather
* deepseek41 : keep the Engram hash constants in hparams.ple_*, one module per Engram layer
* deepseek41 : drop the LLAMA_DSV41_KV_MIRROR toggle
* deepseek41 : keep il last in build_hc_mixes
* deepseek41 : sort the top-k rows before the sparse gather
* tests : add deepseek41 to the Metal fusion baseline
---------
Co-authored-by: Daniel Han <danielhanchen@gmail.com>
* hexagon: KV-head team split for HMX flash attention in multicore
* hexagon: merge FA_HEAD_SPLIT and FA_SPLIT into single GGML_HEXAGON_FA_SPLIT flag
0 = off (token-split only); 1 = head-split, plus team-split when
n_kv_heads < N_cores so stranded cores still participate (e.g. gemma-4
with 2 KV heads on 4 cores). Where heads divide evenly across cores,
plain head-split is used.
The src1 power-of-two rescale in kernel_mul_mm_id is superseded by the
per expert-token amax scaling done in the compute graph
(ref: https://github.com/ggml-org/llama.cpp/pull/30861), which keeps the
MoE down-projection activations within f16 range before the op runs.
Drop the two-stage reduction kernels, their scratch buffer and the
GGML_PREC_F32 dispatch added in https://github.com/ggml-org/llama.cpp/pull/26223,
as well as the test cases that drive MUL_MAT_ID src1 out of f16 range.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* ui : move models manager to its own route, keep Discover as a dialog
Models now lives at #/models instead of a dialog. Discover opens as a
standalone dialog over it, and leaving the page returns to the chat it
was opened from.
Assisted by: Claude (Opus 5.5 Medium)
* ui : polish the models page layout and filters
Give the models page a page header with the Discover action, calmer
column labels, tinted section and family headers, and an empty state.
Load the model list on a direct visit, slide the detail pane over the
table below xl, and draw the filter toggles as one box matching the
context select, with room for their shadow.
Assisted by: Claude (Opus 5.5 Medium)
* ui : round the models table family headers
Assisted by: Claude (Opus 5.5 Medium)
* ui : fix models page navigation and Discover gating
Record the models page's way back after every navigation, so Back,
a typed URL or a reload no longer return to a stale chat. Escape
closes the detail pane, then leaves the page. Gate every Discover
entry point on one shared check, recheck the API key on /models,
skip the duplicate router fetch on a direct load, and share the
route load between pages.
Assisted by: Claude (Opus 5.5 Medium)
* ui : show model details in a drawer
Open the models manager's detail pane in a drawer, from the right on
desktop and from the bottom on a phone, instead of an inline side
panel and a full-screen cover. Pad the pane on both sides, drop its
close button on a phone, and assert that filtered rows are gone from
the page rather than merely hidden.
Assisted by: Claude (Opus 5.5 Medium)
* ui : make whole models table rows clickable
A click anywhere on a model or quant row opens it, unless it lands on
one of the row's own controls (load toggle, actions menu). The name
button stays the keyboard path to the same action.
Assisted by: Claude (Opus 5.5 Medium)
Use the dropdown menu surface (border, radius, background, shadow,
padding) for the hover card and the context gauge popup, and fix the
hover card open/close animation to use bits-ui data-state attributes.
Assisted-by: Claude (Opus 5.5 Medium)
The two T4 runners in the ggml-ci group (ggml-6-x86-vulkan-t4,
ggml-8-x86-nvidia-a10) share the same labels and can both pick up the
same jobs, so extend the per-runner results dir from #31046 to them.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
The generic scale pass in llama-model.cpp loads the per-expert
NVFP4 scale2 tensors (ffn_{up,gate,down}_exps_s) and the per-tensor
MXFP8 scales for the dense FFN and the shared expert, but
minimax-m3.cpp never forwarded any of them to the graph.
build_moe_ffn() defaults these to nullptr, so build_lora_mm_id()
skipped its w_s branch entirely and the routed-expert weights came
out global_scale times too large, swamping the residual (reported
G ~= 15123 for M3.1 Flash) and producing token soup.
Forward all scale tensors at the three call sites, matching the
pattern used by flint and the other NVFP4-capable MoE models.
Co-authored-by: origin-labs-ai <origin-labs-ai@users.noreply.github.com>
Add test-speculative-rejection-sampling to the ctest -E exclusion list of the macos-latest-arm64 CI job.
Assisted-by: dsh:llama.cpp/ggml-org/Flint-GGUF (RTX 5090 32GB, DDR5 128GB)
The amax POOL_1D reduction spans the whole matmul row, so a device cannot
compute it from an axis-0 split. The Meta split-state case reported
MIRRORED, which gives each device a partial max of its own slice. Drop
that case and deny -sm tensor for the arches which build the amax graph,
see [TAG_MOE_AMAX_SCALE].
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* ggml : allow MUL_MAT_ID -1 index to skip the computation
* nits
* update vulkan code
* ggml-webgpu: ensure all threads reach workgroup barriers in mul_mat_id_vec
(cherry picked from commit d8a4b9efb6e2b9abf3a866e847756b7e6bbc6b69)
* fix cuda
* openvino
* exclude CUDA from this PR
* add TODO tag
---------
Co-authored-by: yomaytk <yoshimura.masashi.frbs@gmail.com>
* tests : add speculative rejection draft sampling test
This commit adds a test for speculative rejection sampling test that
excercises the speculative rejection sampling mechanism in
common/speculative.cpp.
The motivation for this is have a test that verifies that the
speculative draft models emitts tokens that match the target models
distribution. So when rejection sampling occurs it should still
produce the same probability distribution as if we directly sampled
from the target model.
Running the test:
```console
$ ctest --test-dir build -R "^test-speculative-rejection-sampling$" \
--n_trials 5000 -V
```
This will generate two models, one target and one draft model.
For subsequent runs use the "fixture exclude any" option to avoid
recreating the models:
```console
$ ctest --test-dir build-cpu-debug -R \
"^test-speculative-rejection-sampling$" \
-FA "^speculative-models$" --n_trails 5000 -V
```
* remove draft model
* ci : fix data race between runners sharing the model cache
Two runners per mac mini share HOME, so concurrent jobs corrupt each
other downloads, converted models and venv under the shared MNT cache.
- gg_wget: download in a temp dir, publish with mv, so readers only
ever see complete files (-N cannot be combined with -O)
- convert, quantize and wikitext unzip now produce under a per-runner
scratch dir in OUT, published to MNT via hardlink when missing
- gguf-split and quantize tests get a per-runner TMP_DIR
- python venv moved from MNT to OUT
- metal, vulkan and webgpu mac jobs pass a per-runner OUT path
Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)
* ci : clean up scratch models after tests
Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)
* ci : fix q4_0 publish when a cached model exists
cp -n skips an existing file with exit code 1, so the ln fallback
aborted the qwen3 job under set -e. Publish through a unique temp
file with an atomic mv: readers never see a partial file and a
stale cached q4_0 gets replaced.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)
* ci : cache-first gg_wget to avoid re-downloads
HF sends no Last-Modified, so wget -N re-downloads the full files on
every run. Skip the download when a non-empty cache file exists: the
HF CI artifacts are static. Downloads still go through a temp dir and
an atomic mv, so concurrent runners never read partial files.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)
* HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable
amdgcnspirv rejects target-specific builtins unless they are guarded by `__builtin_amdgcn_is_invocable`, so 66 of 145 HIP files failed to build
Wrap the perm calls in a helper that keeps v_perm_b32 when the target has it and falls back to a portable byte select otherwise
* move perm helper to vendors/hip.h
* mtmd : lfm2 bilinear resize matches torchvision
Lfm2VlImageProcessor and d1-omni vision.py resize uint8 images with
torchvision (antialias=True). That is Pillow's algorithm, except that the
weights are int16 with as many fractional bits as the largest weight allows,
where Pillow uses 22 bits. The Pillow port was 1-2 levels off on about 1% of
the values of the d1-omni and LFM2.5-VL crops.
RESIZE_ALGO_BILINEAR_TORCH rounds the weights of the Pillow port to that
precision and scales them back to 22 bits, which gives the output of
torchvision with the same passes, and "bilinear" maps to it for LFM2 and
d1-omni. Other models keep the Pillow algorithms.
* mtmd : lfm2 bicubic resize matches torchvision
Lfm2VlImageProcessor resizes with torchvision for the bicubic checkpoints too
(d1-3B, LFM2-VL), so "bicubic" maps to RESIZE_ALGO_BICUBIC_TORCH for LFM2: the
Pillow bicubic filter (a = -0.5, as torchvision with antialias) with the int16
weights of torchvision. The Pillow port was 1-2 levels off on about 0.25% of
the values of the d1-3B crops.
* use bitfield, add hparams setter
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Add Soprano-80M and Soprano-1.1-80M (https://huggingface.co/ekwek/Soprano-1.1-80M)
as an audio generation model: a Qwen3 backbone plus a Vocos decoder that runs
as a gen-only mmproj.
- convert: detect Soprano checkpoints and export the backbone as qwen3 and the
decoder.pth Vocos weights as the mmproj
- vocab: add the soprano pre-tokenizer (lowercase, whitespace collapse,
per-digit split, [UNK] fallback)
- mtmd: add the Vocos graph and ISTFT, allow loading a gen-only mmproj for
Soprano and include the gen-audio context in mtmd_get_memory_usage
- mtmd-helper: add the Soprano pipeline; the decoder runs over the whole
utterance, so stream mode sends the audio in one chunk after flush()
- tts: document Soprano usage
Assisted-by: Codex
Assisted-by: Claude Opus 5.5
* hexagon : add SGN support
Assisted-by: OpenCode
* hexagon : optimize SGN with sign bit
Assisted-by: OpenCode
* hex-sgn: add f16
* hex-unary: update SGN to use generic unary_colapse and relax colapse checks to apply to more conditions
* hex-unary: improve dims colapsing and run it in the host-side solver
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
The __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 (and _gfx12) builtins
take their A/B operands as vectors of short (_ExtVector<16, short> /
_ExtVector<8, short>), not __bf16. Declaring the fragments as
ext_vector_type(N) __bf16 fails to compile:
error: cannot initialize a parameter of type 'short
__attribute__((ext_vector_type(16)))' with an lvalue of type 'const
bf16x16_t' (vector of 16 '__bf16' values)
short and __bf16 are both 16-bit and the value is only ever passed to
the builtin via reinterpret_cast over raw bytes, so changing the element
type is bit-identical. The operand type has been short since the builtin
was introduced, so this works across LLVM/ROCm versions without a
version gate. Fixes HIP builds for RDNA3 (gfx1150/1151) and RDNA4
(gfx1200/1201).
Signed-off-by: Pablo Cholaky <waltercool@slash.cl>
* CUDA: fix build after chunked CUDA graphs (#30804)
* Apply suggestion from @JohannesGaessler
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Fix newlines in common.cuh
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* qwen4exp : build the QSA selection without the kq_mask
* clean
* lang
* qwen4exp : reshape the fused QSA selection to the kq_mask streams
ggml_top_k_mask with a single keep row returns [n_kv, n_tokens], while
build_attn_qsa expects the kq_mask layout [n_kv, n_tokens/n_stream, 1, n_stream].
With several streams the shape assert failed at graph reserve.
* test the QSA optimization
* reword
* update build_fused_top_k_mask instead of handwiring
* simplify
* keep il at the end for consistency
Run npm audit fix, then bump the remaining vulnerable packages within
their current majors: storybook 10.6.2, vitest 4.1.12, vite 7.3.7,
@modelcontextprotocol/sdk 1.32.1, dompurify 3.4.17, mermaid 11.17.2,
@vite-pwa/assets-generator 1.0.4 and @tailwindcss/typography 0.5.20.
Override katex to 0.18.11 and postcss-selector-parser to 7.1.6 since
their dependents still pin vulnerable ranges, bump the sharp override
to 0.35.5, and pin @storybook/svelte so npm does not resolve the 11.0
alpha through the addon-svelte-csf peer range.
Storybook stays on 10.6.2 as 10.6.3 is newer than min-release-age.
Assisted by: Claude (Opus 5.5 Medium)
Anthony Wang