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)
* hex-fusion: support for GELU+MUL, REPEAT+GET_ROWS and BINARY_NX (back-to-back ADD/MUL/...)
* hex-fusion: factor out fusion logic into separate .h and .cpp and update all related code
* hex-fusion: allow for RMS_NORM+MUL fusions for non-contig tensors
* hex-fusion: streamline fusion eligibility checks and cover missing corner cases
* hex-fusion: streamline fusion flow
* hex-opbatch: make n_tensors limit more scalable
* hex-bufs: adjust the trace buff sizing
* hex-glu: cleanup glu-ops interfaces and add kernel_params
* hex-kparams: clean up kparams handling in unary, ssm-conv, and other ops
* hex-kparams: update the rest of the ops PAD,DIAG,... to use kparams
* hex-unary: add support for collapsing dims
* hex-ops: cleanup redundant runtime checks
* hex-ops: consistent use of kernel_type in all kparams
* hex-fusion: remove redundant check in gdn+cpy
* hex-fusion: further streamlining of fusion rules and code
* hex-fusion: cleanup naming
* hex-opfmt: cleanup op formatting
* hex-fusion: bump fusion lookahead to 64
* hex-fusion: fuse SSM_CONV+SILU
* hex-fusion: fix hvx register spills in fused ssm_conv+silu
* hex-fusion: add more GLU fusions
* hex-fusion: remove duplicate checks and use ggml-common helpers
* hex-kparams: refactor kparams into htp-kparams.h/cpp
* hex-fusion: fix binary-nx fusion regression
* hex-kparams: another pass at cleaning up kparams (removing unused, etc)
* hex-kparams: streamling nameing and move _check_ wrappers into .h
* hex-kparams: templetize htp_kparams_check_<op> to reduce code bloat
* hex-ops: de-duplicate op dispatch switches
* hex-kparams: trim down exposed kparams calls
* hex-kparams: simplify type casting of kparams structs
* hex-fence: de-dup fence tensor construction
* hex-opbatch: simplify tensor shape init/checks
* hex-mm: cleanup leftovers of the original QK defines
* hex-fa: fix stack overflow on older npu firmware
* hex-ctx: move dma_caches into ops_ctx
* server: realtime API: send complete after committed audio drains
* server: realtime API: complete a commit at its end row
Completion is triggered when the stream runs dry, which is a timing
proxy: audio appended while the tail drains leaks into the committed
item, a final commit before the drain merges both items, and back to
back commits collapse into one.
mtmd_audio_stream_flush() records the row count at the end of the
flushed audio (mtmd_audio_stream_get_n_rows_flushed), next to the rows
produced so far (mtmd_audio_stream_get_n_rows). The slot keeps a FIFO
of these end rows and completes each commit once every row up to it is
decoded.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* opencl: add Adreno depthwise conv2d
Assisted-by: OpenCode
* opencl: use local program for depthwise conv2d
Assisted-by: OpenCode
* opencl: use singed type for args
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
qwen4exp published the hc-wide residual [n_embd, hc, n_tokens] as
t_layer_inp, but the drafter reads layer inputs as [n_embd, n_tokens]
rows: llama_context sizes each embd_layer_inp buffer as n_embd * n_batch
and the DFlash encoder input is n_extract * n_embd wide. Evaluating a
DFlash draft over a qwen4exp target therefore aborts on the first request
in llama_context::extract_layer_inputs.
Collapse the streams with a plain mean before exposing them, as deepseek4
already does for its hyper-connection residual, and build the mean only
for the layers a drafter actually requested.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* Multi-column on Intel XMX matrix engines refactored
Multi-column (2 to 80 activation columns) mul_mat_vec_q for reordered K-quant and Q8_0 weights on the Intel XMX matrix engines (int8 DPAS, ESIMD), for the SYCL backend. It targets speculative decoding verification (2 to 8 tokens per step) and short prompt batches (9 to 80 tokens). Three types also use it for a single column (Q2_K, Q3_K, Q8_0).
* Fixed regression bug
real bug found: the first version produced NaN whenever the row count wasn't a multiple of 4
* Removed GGML_TYPE_Q8_0 activation
* rewritten xmx_mul_mat
64-byte alignment and 16-byte pitch check on the activations
* Implemented fused gate/up kernel
One launch computes up * act(gate) for both FFN weights, reading the activations once and applying SWIGLU or GEGLU in the epilogue. It covers up to 16 columns and works for all six XMX types.
* Removed the temporary XMX_NOGLU environment switch
* Q2_K min-sum optimization
Q2_K weight is d·scl·q − dmin·mc, where scl (the 4-bit scale) and mc (the 4-bit min code) apply to each half of a 32-weight group
* Added dispatch check and GRF sweep
tg128 on Q4_K_M goes from 26.67 to 28.23 tok/s on the B70 (+5.9%).
Q6_K is unchanged.
MUL_MAT passes 1727/1727 and MUL_MAT_VEC_FUSION 1361/1361.
Perplexity at batch 1 is 5.2064 with XMX and 5.2058 without.
* Renamed variables
* Added 2 environment variables
GGML_SYCL_XMX_MAX_COLS (default 80, any value 0 or higher) and GGML_SYCL_XMX_GLU_MAX_COLS (default 16, range 0 to 16): the largest batch for the fused gate/up kernel. It is clamped to 16 because that is the most the fused kernel is compiled for, so the variable can only lower it.
* Update esimd.hpp
1. All five quant structs were repeating the same boilerplate, so I moved it into shared pieces.
2. A new helper xmx_term does one DPAS, converts it to float and scales it per row. The four branches are now one or two lines each, and the shared Cb += z * d8r is written once.
The refactor passes all three op tests. MUL_MAT_VEC_FUSION is 1369/1369, MUL_MAT is 2281/2281 (the count includes your uncommitted f32 tests) and MUL_MAT_ID is 937/937.
* nemotron : support non-standard GQA, BPE tokenizer and domyn-small pre-type
* gguf-py : write Metaspace prepend_scheme as GGUF metadata
* vocab : parse metaspace prepend_scheme into enum, rename key
* vocab : use metaspace prepend_scheme in detokenize, add TODOs for other tokenizers
* vocab : use constant for metaspace escaped space
---------
Co-authored-by: Martin Cimmino <martin.cimmino@domyn.com>
* server : support typed content input for rerank endpoint
Assisted-by: Opencode/Qwen3.8 27B
* server : address rerank multimodal review findings
- reject multimodal rerank input when the model has no rerank chat template (images were silently dropped)
- translate handle_media download/base64 failures to 400 instead of 500
- reject negative top_n with 400 instead of letting resize throw
- fix comment typos, redundant comments, indentation and trailing whitespace
- README: document that TEI input accepts the same shape as Jina; drop stray blank line
Assisted-by: opencode
* Trimmed README
---------
Co-authored-by: timothywang21 <timothywang21@users.noreply.github.com>
* sycl: fix CONV_2D wrong results for non-contiguous input
The kernel indexes the `input` as a contiguous [W, H, C, N] tensor, but
supports_op claimed the op for any layout. A non-contiguous (e.g. `CWHN`
permuted) input was read with the wrong element offsets, silently
producing wrong results.
Gate the CONV_2D op on `input` and `kernel` contiguity, matching the
CUDA backend contract, and assert input contiguity in the op.
`test-backend-ops -b SYCL0 -o CONV_2D`: goes from 16 failures to 0,
none.
```text
2006/2006 tests passed
Backend SYCL0: OK
Backend 2/2: CPU
Skipping
2/2 backends passed
OK
```
Assisted-by: pi (unsloth/Qwen3.8-27B-GGUF:Q4_K_M via llama.cpp)
* sycl : CONV_2D cwhn=1 support
Make the SYCL backend compute `GGML_OP_CONV_2D` correctly for `cwhn=1` (CWHN
channel-order) inputs and kernels, so the previously failing test-
backend-ops CONV_2D cases pass on SYCL instead of being skipped or
miscomputed.
- Implement CWHN layout support in `ggml/src/ggml-sycl/conv2d.cpp`
(`ggml_sycl_op_conv2d`).
- Update SYCL `supports_op` for `GGML_OP_CONV_2D` in
`ggml/src/ggml-sycl/ggml-sycl.cpp` to accept both layouts.
Assisted-by: pi (Greatjedi/Qwen3.8-27B-Intel-Arc-Tuned-GGUF via llama.cpp)
* sycl docs : update ops for SYCL
Satyam Thakurandorigin-labs-ai
Alexey Kopytko
Georgi Gerganov
Xuan-Son Nguyenandyomaytk
Daniel Bevenius
Ruben Ortlam
Arseniy Obolenskiy
Yuri KhrustalevandXuan Son Nguyen
Hans
kurquharandMax Krasnyansky
Adrien Gallouët
WalterCool
thelittlefireman
SreeramandJohannes Gäßler
Christian Loehnert
Aman Gupta
Aleksander Grygier
Kartik Gulia
Sigbjørn Skjæret
Raul Torres
Titaniumtown
Pranesh GonegandlaandPranesh Gonegandla
Todd Malsbary
pzhu
tc-mb
Martin Kleinberger
Gaurav Garg
budiony
Bill Sideris
martin-cimminoandMartin Cimmino
Tim Wangandtimothywang21
Michael Taylor
tkran2
Gabe Goodhart