* metal : faster MXFP4 mat-vec kernel
Decode the fp4 values with the fp16 bit trick instead of a threadgroup lookup table, and load the packed nibbles as aligned 32-bit words. The kernel no longer uses threadgroup memory.
* metal : launch contiguous unary and binary ops with full threadgroups
The contiguous unary path and the row-broadcast binary path used one thread per threadgroup. Use up to 256 threads per threadgroup and bound-check the tail.
* metal : split GLU rows across threadgroups
A GLU row ran on one threadgroup, so single-row decode used a single threadgroup for the whole op. Split each row over threadgroups of up to 256 threads.
Requested in the #31125 review. The cases duplicate coverage that is
already in the file and add CI time. Test-only, nothing changes in the
kernel, the docs or the CMake option.
Assisted-by: Claude
* hex-upscale: basic upscale ops (not fully vectorized)
* hex-upscale: optimize nearest 2x, 4x and bilinear 4x
* hex-upscale: add support for non-contig rows and dims collapsing
* hex-upscale: use vgather instead of scalar staging
* hex-upscale: move all kernel select logic to the host, review comments and cleanups
* tests: add permuted UPSCALE test to cover qwen3-vl image->text scenarios
* sycl : add DPAS flash attention kernel for quantized KV cache
Decode and small batches (1 to 31 query tokens) now run flash attention
straight on q4_0 and q8_0 KV blocks, instead of scalar dequant or
converting the whole cache to f16 per call. Q.K uses s8 x s4 and s8 x s8
DPAS on XMX, V.P stays f16. dpas.s4.s8 is inline vISA since joint_matrix
has no s4.
Covers head sizes 64 to 512, GQA up to 32 and attention sinks. Prompts
stay on the existing kernels. GGML_SYCL_FA_DPAS=0 turns it off.
B580, Llama 3.1 8B Q4_K_M, q4_0 KV, tg128: 31.3 to 54.2 t/s at 16K
context. FLASH_ATTN_EXT 4582/4584, same 2 failures as master.
Assisted-by: Claude
* sycl : 192 head dropped, Wildcat Lake added to docs
Assisted-by: Claude
* sycl : stricter checks from review
Assisted-by: Claude
* [SYCL] : rename XMX macro, add GGML_SYCL_SUPPORT_XMX option, document DPAS
Assisted-by: Claude
* sycl : widen reordered q4_0 dequantization stores
The reordered-layout Q4_0 conversion kernel processed one qs byte per
work-item with two scalar stores, strided 16 elements apart. Give it
the same treatment the Q8_0 kernel got in #30775: a flat
sycl::range<1> launch with four qs bytes (8 outputs) per work-item
inside one 32-element block and two vec<dst_t,4> stores.
Same float math and rounding as the scalar kernel; output is
bit-identical. The existing k % QK4_0 == 0 assert guarantees whole
32-element blocks, so the k/8 launch covers every row exactly. The
now-unreferenced scalar reorder kernel is removed.
Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +6.4% to +7.3% pp512/pp2048 on Q4_0
models (tested up to 27B dense); fresh-layout prefill and decode
unchanged; PPL within run noise.
* sycl : vectorize reordered q4_K dequantization stores
The reordered-layout Q4_K conversion kernel already loads its
quantized inputs through sycl::vec but writes the dequantized
results with scalar stores. Build the results in a vec<dst_t,n>
register and emit two vector stores instead. Same per-element
math; output is bit-identical.
Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +14.7% to +15.3% pp512/pp2048 on
Q4_K-dominant models (tested up to 27B dense); fresh-layout
prefill, decode and PPL unchanged.
* sycl : remap reordered q5_K dequantization to vec4 loads/stores
The reordered-layout Q5_K conversion kernel gave each thread only
four outputs fed by scalar ql/qh byte loads plus a local-memory
scale broadcast, so the stores were never its bottleneck: widening
the stores alone measured within noise on mixed-quant models and
+0.5% on a q5_K-dominant model. Re-decompose to 32 threads per
block, each consuming four consecutive ql and qh bytes and emitting
two aligned vec<dst_t,4> stores; a 4-element run shares one scale
pair, so the per-element math and the output are identical.
Arc Pro B70, primed prefill on a q5_K-dominant 27B model: +3.6%
pp512/pp2048 over the scalar kernel; fresh-layout prefill, decode
and PPL unchanged. The GET_ROWS reorder path moves to the same
32-thread decomposition.
* sycl : remap reordered q6_K dequantization to vec4 stores
The reordered-layout Q6_K conversion kernel wrote four scalars per
thread, strided 32 elements apart, each pair of neighbors sharing
neither instruction stream nor store transaction. Re-decompose the
same 64-thread work-group so each thread emits one aligned
vec<dst_t,4>: a 4-element run shares one 16-element scale group,
one ql nibble source and one qh bit pair, so the per-element math
and the output are identical to the scalar kernel.
Arc Pro B70, primed prefill: +28.7% pp512/pp2048 on a Q6_K 27B
model; fresh-layout prefill, decode and PPL unchanged.
* sycl : vectorize reordered q3_K dequantization stores
Each thread of the reordered-layout Q3_K conversion kernel already
computes four consecutive outputs; build them in a vec<dst_t,4>
register and emit one vector store instead of four scalars. Same
per-element math; output is bit-identical.
Arc Pro B70, primed prefill: +16.8% pp512/pp2048 on a Q3_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.
* sycl : vectorize reordered iq3_xxs/iq3_s dequantization stores
Both reordered-layout IQ3 conversion kernels write eight contiguous
outputs per thread as scalar stores (four from each grid lookup).
Build two vec<dst_t,4> registers and emit two vector stores. Same
per-element math; output is bit-identical.
Arc Pro B70, primed prefill: +20.3% to +21.1% pp512/pp2048 on
IQ3_XXS- and IQ3_S-dominant 27B models; fresh-layout prefill,
decode and PPL unchanged.
* sycl : remap reordered q2_K dequantization to vec4 stores
Re-decompose the reordered-layout Q2_K conversion kernel the same
way as Q6_K: one aligned vec<dst_t,4> per thread instead of four
scalars strided 32 apart. A 4-element run shares one scale byte and
one 2-bit lane, so the per-element math and the output are
identical to the scalar kernel.
Arc Pro B70, primed prefill: +13.1% pp512/pp2048 on a Q2_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.
* sycl : widen reordered mxfp4 dequantization stores
Give the reordered-layout MXFP4 conversion kernel the same flat
sycl::range<1> launch as Q4_0: four qs bytes (8 outputs) per
work-item inside one 32-element block and two vec<dst_t,4> stores.
Same LUT lookups and per-block scale; output is bit-identical. The
k % QK_MXFP4 == 0 assert guarantees whole blocks, so the k/8 launch
covers every row exactly. The old scalar kernel is removed.
The reordered layout for MXFP4 is installed on the MoE expert path
(reorder_qw_mxfp4_moe); dense MXFP4 tensors keep the normal layout.
Arc Pro B70, MXFP4 MoE 20B model: +57% pp2048 / +61% pp512 prefill
(primed and fresh alike), decode unchanged, PPL within the model's
run-to-run noise.
SCALE now accepts a src0 view with arbitrary row strides. translate_scale
used the raw input, and on the dynamic path a VIEW translates to its whole
source tensor, so a strided view input scaled the entire source. The model
output then had the source's size, and copying it back overran the result
buffer, which crashes test-backend-ops on the SCALE view cases.
Slice the input with process_view_input_new, as the other translators do.
Reject an in-place SCALE on a non-contiguous view: its result is copied
back densely and would land in the wrong elements.
* hexagon: IM2COL over multiple threads with output-col blocks
Assisted-by: Claude Opus 5.5
* hex-im2col: fix padding handling
* hex-im2col: loop unroll common cases
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
test_cancel_request enables context shift so the generation only ends
when the request is cancelled, then polls /slots with a deadline.
Without context shift the generation fills the context and frees the
slot on its own on fast machines, so the test passes even with
cancellation disabled.
test_server_sleep uses the existing wait_for_sleep helper instead of a
fixed sleep.
* common: bound control vector layer indices by the model layer count
A direction tensor layer index is parsed from the tensor name with no
upper bound, and the buffer size and write offset are computed in int.
A large index overflows both, so the accumulation writes before the
start of the buffer.
Pass the model layer count to the loader, reject indices above it like
the existing zero index check, and compute the sizes in size_t.
* mtmd: fill M-RoPE positions by whole merge tiles
The position loops of the Qwen-VL, GLM, PaddleOCR and Qwen2.5-VL
encoders hardcode a 2x2 merge inside a stride of n_merge, so a patch
grid that is not a multiple of the merge writes past the positions
buffer. Iterate over whole merge tiles with n_merge inner steps, which
bounds the writes by the patch count and gives the same positions on
aligned grids.
PaddleOCR also requires a merge aligned image at graph build, as the
Qwen-VL family already does.
* opencl: run the tiled q4_K lm_head through a tiled GEMM at batch > 1
* opencl: split tiled dp4a q4_k gemm into its own file and test for dp4a
Assisted-by: Codex
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* hexagon: add conv2d op
* hex-conv2d: add double-buffered staging
* hex-conv2d: full 64-wide activation block transpose straight from staged buffer
* hex-conv2d: implicit gemm
* hex-conv2d: remove scalar path and kernel
* hex-conv2d: unify conv2d kernel params for all paths
* hex-conv2d: scale HMX result before f16 storage and then restore during unpacking to prevent overflow
* hex-conv2d: fix f16 overflow in HMX paths
* hex-conv2d: refactor to rebase to new op fusion machinery
* hex-conv2d: restructure DMA pipelines and trace scopes
* hex-conv2d: removed the HVX_WIDTH and HVX_TILED paths
* hex-conv2d: remove in-loop vector spills
* hex-conv2d: refactor kparams, vtcm-layout and dma-pipeline to match the other ops
* hex-conv2d: add basic support for CONV_2D_DW
Co-Authored-By: Pranav Dhinakar <pdhinaka@qti.qualcomm.com>
* hex-conv_2d_dw: use DMA in all paths
* hex-conv_2d_dw: support for chunking larger inputs
* hex-conv_2d_dw: remove scalar/hvx access for w tensor (replaced with sync dma copy)
* hex-conv-2d: fix warnings
* hex-conv2d-dw: vectorize all paths
* hex-conv-2d-dw: tree-based reductions and unrolling to optimize
* hex-conv-2d-dw: uninline compute loops from dma loops
* hex-conv-2d-dw: move all kernel selection into the host solver and simplify/specialize kernel dispatch
* hex-softmax: update obsolete gates to enable SOFT_MAX for mmproj
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Pranav Dhinakar <pdhinaka@qti.qualcomm.com>
* mxfp4 sve 256 impl completed
* working with pretty good perf improvement
* review comments addressed
* review comments for SVE 256 check
* trailing white spaces removed
---------
Co-authored-by: Anubhav Srivastava <anubhav.srivastava@fujitsu.com>
* llama : fix DSV4 compressor padding for simple batches
Count actual tokens per sequence when padding compressor writes. Add a graph reuse regression test for odd-sized DeepSeek V4 and V4.1 batches.
Assisted-by: Codex
* tests : cover multiple sequences in DSV4 padding regression
Exercise one and two sequences at different positions while keeping a fixed ubatch shape.
Assisted-by: Codex
* llama : remove simple batch splitting for DSV4
Use the existing sequence splitting to keep n_seq_tokens consistent with
the compressor padding calculation. Extend the graph reuse regression to
6-token batches.
Assisted-by: Codex
* codeowners : refresh teams and maintainer PRs link
Update the CODEOWNERS team legend with the current `ggml-org` team
memberships, and extend the maintainer PRs link in the README to cover
all current members of the `maintainers` team.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* codeowners : list remaining backend sub-teams
Add the `ggml-*` sub-teams that were missing from the team legend at the
top of the file: base, cpu, et, hip, openvino and riscv. The legend now
covers every sub-team of the `ggml` and `llama` maintainer groups.
No ownership rules change - only the comment block is extended.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* codeowners : use ggml-hip team for HIP backend
Route `/ggml/src/ggml-hip/` to `ggml-org/ggml-hip` instead of a single
collaborator, so review requests reach the whole team.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* codeowners : add llama-examples team
List the newly created `llama-examples` team in the team legend at the
top of the file.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* codeowners : use llama-examples team for examples
Collapse the per-directory rules that only listed a single collaborator
into a single `/examples/` rule owned by `ggml-org/llama-examples`.
Directories owned by other collaborators keep dedicated rules, since the
last matching pattern takes precedence in CODEOWNERS.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* cont : update zdnn
* codeowners : use llama-base team for src
Route `/src/` to `ggml-org/llama-base` and list the new team in the
legend at the top of the file. Rules for individual `src` files are kept,
since the last matching pattern takes precedence; `/src/models/` drops
its individual owner, which is already covered by `llama-models`.
Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* mtmd : add LFM3-VL and LFM3-Omni video input
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Assisted-by: Claude Opus 5.5
* mtmd : reserve the soundtrack command to avoid a GCC 11 warning
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* ui: build the reasoning menu from the template effort levels
When the chat template reports effort levels, the menu offers Default,
Off and those levels, and a pick is sent as reasoning_effort. A
template that reports none keeps the token budget menu.
Picking a model other than the one that produced the last answer resets
the conversation effort to Default, and an effort the current menu
lacks resolves to Default.
* ui: wait for a cold router model before building the request
A router serves the props of a loaded model only, so a request sent to
a cold model was built without its template effort levels and fell
back to a token budget. The request now waits for the model to load,
joining a load already in flight, and a stop during the wait sends
nothing.
Status waiters hold a list per model so a load and the requests that
wait on it each get theirs, and a props request in flight is shared
with a second caller instead of answering null.
* fix: clamp the A prefetch to end_k in the Vulkan coopmat1 int8 matmul, like the B prefetch
* fix: zero the A scale for k steps past end_k in the Vulkan coopmat1 int8 matmul instead of relying on the clamped block
* chat : strip trailing whitespace from parsed reasoning and content
Assisted-by: Claude Fable 5.1
* chat : stop qwen3-coder reasoning and content before template separators
Assisted-by: Claude
* tests : remove the Qwen parser case from the chat session tests
Assisted-by: Claude
* common : handle partial prefix until match at end of generation
---------
Co-authored-by: Zeshawn Zahid <zzshawnn@gmail.com>
Co-authored-by: Alde Rojas <hello@alde.dev>
* ui : show save and reset settings buttons only when relevant
Assisted-by: Claude (Opus 5.5 Medium)
* ui : move reset settings button next to reload app, pad settings caption
Assisted-by: Claude (Opus 5.5 Medium)
* ui: support OGG audio attachments by converting them to WAV
Fixes#24164
Assisted-by: Claude (Opus 5.5 Medium)
* ui: convert all non WAV/MP3 audio uploads to WAV
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep new chat draft when adding a system prompt
Fixes#24338
Assisted-by: Claude (Opus 5.5 Medium)
* ui: remove unused initialMessage from chat form
Assisted-by: Claude (Opus 5.5 Medium)
* ui : replace the sidebar search mode with a search dialog
Search opens a spotlight dialog (Cmd+K or the Search entry under New chat in the sidebar) with keyboard navigation over recent chats and sidebar actions. Phones keep the search page.
Sidebar actions show their shortcuts, with new ones for Settings (Cmd+,) and Models (Shift+Cmd+M).
Assisted-by: Claude (Opus 5.5 Medium)
* ui : ignore search dialog keys while an IME is composing
Enter confirms an IME candidate and the arrow keys move through the candidates, so the search field leaves them to the IME instead of opening the highlighted result or moving the highlight.
Assisted-by: Claude (Opus 5.5 Medium)
* server : add realtime conversation and transcription on /v1/realtime
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Assisted-by: Claude Opus 5.5
* examples : default the realtime client to the server port 9931
Assisted-by: Claude Opus 5.5
* server : sample realtime responses right after their own decode
Slots respond in turns on one context: sampling at the start of the next step read the logits of whichever slot decoded last, so concurrent sessions mixed their outputs.
Assisted-by: Claude Opus 5.5
* move more stuff to server-realtime, making it self-contained
* examples : fix the type check of the realtime client
Assisted-by: Claude Opus 5.5
* examples : use the model card prompts in the realtime client
Assisted-by: Claude Opus 5.5
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* model : add LiquidAI/d1-2B and LiquidAI/d1-omni-2B
Assisted-by: Claude Opus 5.5
* convert : keep the speech task templates out of the d1 GGUFs
Assisted-by: Claude Opus 5.5
* small refactoring: no_batching
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* server: expose the reasoning effort levels of the chat template
The capability probe renders a user turn for each candidate level
(minimal, low, medium, high, xhigh, max) and for an unknown level. A
level is honored when it renders without raising and differently from
the unknown level, which shows the template fallback. Levels sharing a
render are aliases and the highest one is kept.
The caps serialize to JSON, and chat_template_caps in /props reports
the list as reasoning_effort_levels, from lowest to highest.
* jinja: keep the caps free of the json dependency
The caps keep reporting their flags as a map, and the effort levels
stay a plain member. common_chat_templates_get_caps() builds the JSON
the server reports from both, so jinja depends on json only for the
probe inputs.
In --models mode the thread count per model is computed and printed,
but never set in the params, so every context runs on the common
default, the number of physical cores. With several models in
parallel this oversubscribes the CPU, and the busy-wait barriers of
the CPU backend make it collapse.
Set the count in cpuparams and cpuparams_batch.
* model : add LiquidAI/LFM3-Audio-2B and LiquidAI/LFM3-Omni-2B
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Assisted-by: Claude Opus 5.5
* convert : read the LFM3-Audio config.json as released
A transformers that knows lfm3_audio rewrites the config into another layout (no lfm/encoder keys).
Assisted-by: Claude Opus 5.5
* mtmd : disable flash attention for the LFM3-Audio depthformer
Assisted-by: Claude Opus 5.5
* tts : add a voice option for LFM3-Audio and LFM3-Omni
The model card offers a female and a male voice; --tts-voice for llama-tts, "voice" for the server /tts.
Assisted-by: Claude Opus 5.5
* clean up template handling
* convert : detect LFM3-Audio and LFM3-Omni by model_type
Assisted-by: Claude Opus 5.5
* move banned tokens --> logits bias --> mtmd_helper
* leave mtmd-cli unchanged
* no more template_asr_system
* left over
* llama : keep the LFM3-Audio encoder convs and codec unquantized
Assisted-by: Claude Opus 5.5
* common, mtmd : shorten comments
Assisted-by: Claude Opus 5.5
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* common : support deferred tools and tool_reference content parts
Assisted-by: Claude Opus 5.5
* common : credit the deferred tools co-author
Co-authored-by: Aldehir Rojas <765227+aldehir@users.noreply.github.com>
Assisted-by: Claude Opus 5.5
* tests : fix the type check of the deferred tools test
Assisted-by: Claude Opus 5.5
* model : add LiquidAI/LFM3-1.7B and LiquidAI/LFM3-VL-2B
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Xuan-Son Nguyen <son@huggingface.co>
Assisted-by: Claude Opus 5.5
* tests : set the SWA rope base of the LFM3 test model
Assisted-by: Claude Opus 5.5
* llama : enable recurrent state rollback for LFM3
Assisted-by: Claude Opus 5.5
* Apply batched suggestions from code review
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Tarek Dakhran <github@dakhran.com>
* gguf-py : align the lfm3vl mapping comments
Assisted-by: Claude Opus 5.5
* llama : fix CRLF line endings from the review suggestions
Assisted-by: Claude Opus 5.5
* tests : add the LFM3 Metal fusion baseline
Assisted-by: Claude Opus 5.5
* llama : keep the LFM3-VL depthwise conv kernels in F32 when quantizing
Assisted-by: Claude Opus 5.5
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Xuan-Son Nguyen <son@huggingface.co>
Keyu Zhu
lhezandHongqiang Wang
Aman Gupta
Jesse Symons
Georgi Gerganov
Max Krasnyansky
Trivikram ReddyandMax Krasnyansky
Anant Shrivastava
Yonghye Kwon
Ravi Panchumarthy
Marco ColomboandMax Krasnyansky
Pascal
Nikhil Sharma
Abhijith Jajam
Bartowski
anubhavfujitsu
abhijain1204fujitsuandAnubhav Srivastava
Aaron Teo
Frosty40
anujj
Xuan-Son NguyenandSigbjørn Skjæret
Tarek DakhranandXuan Son Nguyen
Kris
Aleksander Grygier
uvos
Mengsheng Wu
