Commit Graph
12013 Commits
Author SHA1 Message Date
Tarek Dakhran 889f2ec96d quant : allow quantization for LFM3-Audio codec decoder (#31386)
* quantize : quantize the LFM3-Audio codec decoder, it is Q4_0 QAT

* quantize : drop the QAT note from the LFM3-Audio comment
2026-10-30 11:52:59 +01:00
Alessandro de Oliveira Faria (A.K.A.CABELO) 1ab23c33b7 vendor : update cpp-httplib to 0.64.0 (#31400) 2026-10-30 11:49:58 +01:00
Sigbjørn Skjæret 27b8464e82 ci : add missing ccache-clear to windows behavior check (#31413) 2026-10-30 11:34:25 +01:00
Adrien Gallouët d63fbfe794 cmake : embed an app manifest in all Windows exes (#31329)
* cmake : embed an app manifest in all Windows exes

Signed-off-by: Adrien Gallouët <angt@huggingface.co>

* Remove LLAMA_WIN_APP_CMAKE

Signed-off-by: Adrien Gallouët <angt@huggingface.co>

---------

Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-10-30 10:56:53 +01:00
Daniel Bevenius e216f4ec87 ci : extract most apple jobs into daily workflow (#31408)
This commit extracts all apple jobs except for macos-latest-arm64 into a
new daily workflow.

The motivation for this is to reducd the load on the macos runners.
2026-10-30 10:01:28 +01:00
Lorenzo Bona 344846b4fc gemma3n, gemma4 : create token_embd before the tied output (#31394)
The duplicated output looks up token_embd in its buffer context, so
token_embd must exist first. With the old order a load where input and
output share a buffer type (CPU-only, -ngl 0) held the table twice:
gemma-4-E2B-it Q4_K_M, no mmap, CPU model buffer 3163.73 -> 2947.73 MiB.
2026-10-30 09:36:27 +01:00
Anton IandPascal fb5997274e kv-cache : keep rotation constants on the GPU (#31073)
* kv-cache : keep rotation constants on the GPU

Assisted-by: Codex

* tests : verify resident KV rotation buffers

Assisted-by: Codex

* kv-cache : document remaining multi-device rotation work

Assisted-by: Codex

* kv-cache : support per-layer rotation buffers

Assisted-by: Codex

* kv-cache : fix hybrid SWA rotation inputs and mixed placement check

Assisted-by: Codex

* models : use per-layer rotation in mistral-large4 and dflash

* tests: drop the rotation buffer tests

---------

Co-authored-by: Pascal <admin@serveurperso.com>
2026-10-30 10:15:06 +02:00
Pascal ab3768cdc2 memory: refuse rollbacks past the snapshots of the last ubatch (#31368)
* memory: refuse rollbacks past the snapshots of the last ubatch

The recurrent memory and the DSV4 cache accepted any rollback up to
n_rs_seq, whether or not the last ubatch wrote that snapshot. Each now
records the deepest rollback its last ubatch guarantees and seq_rm
refuses anything past it:

- recurrent: n - 1 for a ubatch of n <= n_rs_seq + 1 tokens, 0 for a
  longer one, which only guarantees the final state
- DSV4: min(n, n_rs_seq), the snapshots are rebuilt from the state
  before the ubatch

The dead rollback block of the rollback test becomes the assertion it
asked for, and rollback_bound checks the exact bound of each memory
after ubatches of 1, 2, n_rs_seq + 1 and n_rs_seq + 2 tokens.

* tests: address review from ggerganov

Detect architectures with llama_arch_from_string in comp_padding and
arch_reserves_single_seq, as in rollback_bound.
2026-10-30 10:11:58 +02:00
shaofeiqiandLi He e0f142d5dd opencl: add initial q2_0 support (#31340)
* opencl: add q2_0 AOS mv

Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Assisted-by: Codex

* opencl: add q2_0 SOA mv and tiled mm

Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Assisted-by: Codex

* opencl: add q2_0 noshuffle gemm/gemv for adreno

Assisted-by: Codex

* opencl: add q2_0 dp4a for adreno

* opencl: optimize q2_0 adreno decode and narrow batches

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-10-29 21:30:05 -07:00
Tarek Dakhran 3c809f4438 common : end reasoning budget forcing on an end tag from the prompt (#31374) 2026-10-29 19:27:50 -05:00
Michael Taylor 36a0e2ffc0 sycl: FLASH_ATTN_EXT failures (#30761) 2026-10-29 17:37:28 -04:00
Georgi Gerganov 418f905a4e ggml : add batched tensor copy API (#30881)
* ggml : add generic batched tensor copy API

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* llama : use batched tensor copy API in context io

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* ggml : add batched tensor copy API

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* llama : batch device io tensor copies

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* ggml : add CUDA support for batched tensor APIs

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* ggml : remove async tensor_n APIs and add null buffer hooks

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* ggml : fix batched tensor 2D detection

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* ggml : group batched tensor operations by buffer type

Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* cont : minor indentantion

* cont : const correctness

* tests : add mock buffer type tests for batched tensor API

Restructure test-backend-tensor to run a unified test set on each backend buffer type and on an ad-hoc mock buffer type. The mock set exercises the dispatch logic in ggml_backend_tensor_{set,get,copy}_n: hook invocation, per-tensor fallback, cpy_tensor_n all-or-nothing, and buft grouping.

Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)

* cont : add TODO
2026-10-29 20:52:17 +02:00
Slobodan JosicandMario Limonciello eb4e244be7 ROCm/HIP: add RDNA4_M family for gfx1170/gfx1171 (#31373)
gfx117X is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4_M and use it where it makes sense.


Assisted-by: Codex
Assisted-by: GitHub Copilot

Co-authored-by: Mario Limonciello <mario.limonciello@amd.com>
2026-10-29 18:23:29 +01:00
Oliver SimonsandSigbjørn Skjæret 68d20b884c Quant: OCP FP8 E4M3 support (#27512)
* WIP OCP FP8 E4M3 support

* Fix python type-checker

* Materialize input_scales also for FP8 when converting

* Fix conversion of https://huggingface.co/RedHatAI/Qwen3.8-27B-NVFP4

1. Per row FP8 wrongly fell to NVFP4 conversion, as NVFP4 checked on
   dims alone. Also check on dtype for FP8
2. Need to reshape FP8 QKV projections scales in the same way that
   weights are reshaped

* Allow remaining FP8 tensors to be dequantized during conversion

* Allow tensor-strategy for FP8-only compressed-tensor ckpts

* Fix and test gguf-py behavior for FP8_E4M3

* Disallow QKV-fusion for FP8 temporarily

* Fix conversion of fused qkvz & FP8 for Qwen3-Next

* Avoid fusing FP8 for --fuse-gate-up-exps as we may have different ws

* Emit warnings when not fusing FP8 during conversion

* More FP8 conversion fixes

* Add supported for strided copies to/from FP8 in CPU

* remove unneeded scratch buffer from FP8 cpu cpy/dup path

* add explanatory comments and TODOs

* Add dedicated encoder and decoder tests for FP8

* Whitelist ggml_op_get_rows in metal instead of blacklisting

required to fix failing CI, and white-listing should be more
future-proof than adding FP8 to the black-list

https://github.com/ggml-org/llama.cpp/actions/runs/38433916498/job/115422880633?pr=27512

* Fix more conversion bugs

* Defer weight materialization during conversion

* Blacklist fp8 in metal backend instead of white-list

* Adjust indents

* Fix lint + execute tests

* Add fast CPU kernel for F8_E4M3

* Conversion fixes

* Fix remaining gaps

* Enable tied weights with shared scales

* Make mandatory-F32 rules apply to FP8 during conversion

* Fix python type checks

https://github.com/ggml-org/llama.cpp/actions/runs/39765670052/job/119780379595?pr=27512

* Fix kaleidai unit-test failures

https://github.com/ggml-org/llama.cpp/actions/runs/39765670125/job/119780383150

* Relax numeric precision constraints for pre-check

* Leave fast-FP8 CPU on RISC-V for follow-up

We already exclude arm from the fast-path

* Allow output scale tensors to be located in any gguf shard

* Move test_conversion_fp8.py test execution to behavior-check.yaml

* Apply batched suggestions from code review

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>

* Move mixed compressed-tensor logic into dequant_fp8 + add unit-tests for it

* Install all requirements also on Windows for behavior-check

Fixes https://github.com/ggml-org/llama.cpp/actions/runs/40535852590/job/122300758223?pr=27512

* Unmap lazy tensors before clean-up in Windows

Should fix https://github.com/ggml-org/llama.cpp/actions/runs/40539797824/job/122314782459?pr=27512

* LF line endings

https://github.com/ggml-org/llama.cpp/actions/runs/40539798168/job/122314791118?pr=27512

* Extract FP8 tests from llama-arch to its own test-fixture

See https://github.com/ggml-org/llama.cpp/pull/27512#discussion_r4400628221

* Update gguf-py/gguf/scripts/gguf_convert_endian.py

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>

---------

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
2026-10-29 18:14:37 +01:00
Xuan-Son NguyenandGeorgi Gerganov 0893ec98f5 skills: specify rule about ggml_build_forward_expand (#31375)
* skills: specify rule about ggml_build_forward_expand

* Update skills/code-review/SKILL.md

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-10-29 15:44:12 +01:00
Xuan-Son Nguyen 78feec062d model: mistral le chonk (#31376)
* llama: extend DSA and sdpark to new model

* fix ci
2026-10-29 15:43:33 +01:00
Aman Gupta 0ac1a7fd35 llama : reuse the tensor staging buffer and skip the direct-io bounce buffer (#31321) 2026-10-29 19:22:43 +08:00
7konigsbergandZeshawn Zahid 7df8b876b9 server : restore the spec checkpoint for drafts longer than n_rs_seq (#31350)
Assisted-by: Claude

Co-authored-by: Zeshawn Zahid <zzshawnn@gmail.com>
2026-10-29 11:28:03 +02:00
Aleksander Grygier 8ab29faec4 ui: keep chats at the bottom until the reader scrolls up (#31311)
* ui: land at the bottom of a conversation every time it opens

Opening a long conversation often stopped short of its end:

- the press grace counted from page load (pressedAt = 0), so a chat that
  landed within 400 ms of a refresh did not follow the rows realizing
  after the jump
- the click that opened a chat (sidebar row, tab) started the same grace,
  so navigating by click had the same problem
- any scroll event away from the bottom stopped following, including the
  ones the browser fires itself while the chat settles
- the chat screen outlives its routes, so going back to the last landed
  conversation through the new-chat screen skipped the landing entirely

scrollToBottom() now clears the grace, only an upward scroll away from the
bottom stops following, and every navigation lands anew. Navigations into
a chat also disable SvelteKit's scroll reset, which moved the previous
conversation to its top for a few frames.

Where overflow-anchor is supported, a 1px pin after the messages is the
only scroll anchor while following, so the browser keeps the bottom in
place itself; WebKit keeps the resize observer.

Assisted-by: Claude (Opus 5.5 Medium)

* ui: let nested agentic content scrollers chain to the chat

Reasoning, shell output and code blocks used overscroll-behavior:
contain, so the wheel stopped dead at their edges instead of
scrolling the conversation.

Assisted-by: Claude (Opus 5.5 Medium)

* ui: stop following the stream as soon as the reader scrolls up

A trackpad or smooth wheel scrolls in steps smaller than the at-bottom
threshold, so each step up was re-stuck and snapped back on the next
growth while a response streamed. Any upward move that is not content
shrinking now releases at once.

The attach effect also read the hook's own state, so every flip
re-attached and re-measured, which undid such a release.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-29 10:18:13 +02:00
Aleksander Grygier f5d4d75419 ui: show submit button when the chat input has content (#31324)
ChatFormActions gated the mic and stop buttons on a `canSubmit` prop that
ChatForm never passed, so it was always false and typing text did not swap
them for the submit button. Use `canSend` instead and also hide the
skip-reasoning button while there is content to send.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-29 10:15:54 +02:00
Aleksander Grygier 3ed915e309 ui: reveal the collapsed sidebar from its logo and toggle it there (#31307)
* ui: reveal the collapsed sidebar from its logo and toggle it there

The collapsed sidebar now shows only the logo; hovering it reveals the
icon strip until the pointer leaves. The expanded sidebar drops its
top-right collapse button and the logo collapses it on hover instead.

Assisted-by: Claude (Opus 5.5 Medium)

* ui: fade the peeked sidebar in at full width and keep the toggle slide

Hovering the logo of the collapsed sidebar now fades in the fully open
sidebar, floating over the page at its full width, instead of resizing
it. Opening or closing with the toggle still resizes and slides: the
sidebar docks to strip width while it slides, then floats out of sight,
and it stays hidden under the pointer after a toggle close. The unused
icon-only strip is removed from the sidebar actions.

Assisted-by: Claude (Opus 5.5 Medium)

* ui: stack the sidebar above the chat tabs on desktop

The peeked sidebar floats over the page, but the chat tab bar shared its
z-index and came later in the page, so the tabs were drawn on top of it.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-29 10:15:31 +02:00
Aleksander Grygier f8fcef2546 ui: add MCP group checkboxes in tools settings and fix indeterminate checkbox colors (#31303)
MCP server groups in the settings Tools tab now get a group checkbox like the
built-in, server and browser groups, toggling the server-scoped tools key.
Indeterminate checkboxes now use the checked state colors.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-29 10:15:09 +02:00
Aleksander Grygier b7cde589e4 ui: request persistent storage on startup (#31299)
Ask the browser to keep IndexedDB and localStorage data under storage
pressure, so conversations and settings are less likely to be evicted.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-29 10:14:53 +02:00
lhez 3dd0f60bb4 doc: update ops for opencl (#31347)
Assisted-by: Codex
2026-10-29 10:14:29 +02:00
uvos 78febbad9d HIP/CUDA: enable sparse attn on CDNA (#31259)
* HIP/CUDA: enable sparse attn on CDNA

* Use __ballot for __ballot_sync on old ROCm versions

* Ignore new large head fattn mma variant
2026-10-29 10:14:11 +02:00
Giuseppe Scrivano abb76ab380 model : Support Limite 1B - Violetto (#29433)
* gguf-py : add Limite arch

Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>

* convert : add Limite support

Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>

* llama : add Limite support

Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>

---------

Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
2026-10-29 10:13:41 +02:00
Keyu Zhu 41be0bdc98 common: select A100 cores for SpaceMiT default threads (#30284)
* common: select A100 cores for SpaceMiT default threads

# Conflicts:
#	tests/CMakeLists.txt

* ggml: export SpaceMiT build marker

* common: scope SpacemiT thread detection macro to llama-common

* common: detect A100 cores in all Linux RISC-V builds

* tests: link barrier test against llama-common

* tests: drop SpacemiT default thread test

* common: follow platform check order
2026-10-29 10:12:42 +02:00
lhezandHongqiang Wang 22183e51f9 opencl: initial support for iq2 (#31267)
* opencl: add iq2_s AOS mv

* opencl: add iq2_s SOA mv and tiled mm

* opencl: add iq2_s noshuffle gemm/gemv for adreno

* opencl: add iq2_s dp4a for adreno

* opencl: optimize iq2_s adreno decode and narrow batches

Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>

* opencl: add iq2_xs AOS mv

* opencl: add iq2_xs SOA mv and tiled mm

* opencl: add iq2_xs noshuffle gemm/gemv for adreno

* opencl: add iq2_xs dp4a for adreno

* opencl: optimize iq2_xs adreno decode and narrow batches

Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>

* opencl: add iq2_xxs AOS mv

* opencl: add iq2_xxs SOA mv and tiled mm

* opencl: add iq2_xxs noshuffle gemm/gemv for adreno

* opencl: add iq2_xxs dp4a for adreno

* opencl: optimize iq2_xxs adreno decode and narrow batches

Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>

---------

Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
2026-10-29 10:12:11 +02:00
Aman Gupta 6d194531e1 metal : faster MXFP4 mat-vec, full threadgroups for unary/binary/GLU (#31075)
* 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.
2026-10-29 10:11:29 +02:00
Jesse Symons fade717ed8 tests : remove duplicate FA test cases added in #31125 (#31357)
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
2026-10-29 10:10:25 +02:00
Georgi Gerganov c9daee149d ggml-metal : fix unary kernels compile on older Metal toolchains (#31360)
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
2026-10-29 10:09:51 +02:00
Max Krasnyansky 8a199ad31a hexagon: support for UPSCALE op (#31331)
* 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
2026-10-29 00:22:34 -07:00
Trivikram ReddyandMax Krasnyansky da6c30ade8 hexagon: improve DDR bw utilization in transposed CONCAT (#31277)
* hex-concat: double row block size for transposed CONCAT

* hex-concat: double buffer scratchpads

* hex-concat: chunked src1 rows in transposed concat

* hex-concat: pad src1 VTCM stride to spread gather lanes

* hex-vgather: apply similar vgather opts to SSM_CONV and GDN

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-10-28 23:57:18 -07:00
Jesse Symons f7c01c5226 [SYCL] add DPAS flash attention kernel for quantized KV cache (#31125)
* 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
2026-10-29 14:01:21 +08:00
Anant Shrivastava bf48103ad1 sycl : vectorize reordered-layout dequantization stores (Q4_0, K-quants, IQ3, MXFP4) (#30976)
* 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.
2026-10-29 13:57:53 +08:00
Yonghye Kwon 6d5ac3c8d5 tests : fix comp_padding CUDA graph reservation (#31344)
Assisted-by: Codex
2026-10-29 07:04:28 +02:00
Ravi Panchumarthy 7ac5727d60 ggml-openvino: slice view inputs of SCALE (#31334)
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.
2026-10-28 21:11:13 -07:00
Marco ColomboandMax Krasnyansky 2e0d86bebc Hexagon: IM2COL improvement - block based HVX parallelization (#31050)
* 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>
2026-10-28 19:04:18 -07:00
Pascal e09a78a25d tests: fix CI flakes in server cancel and sleep tests (#31330)
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.
2026-10-29 01:35:01 +01:00
Nikhil Sharma 369dcf309c ggml-cpu: add SVE implementation of ggml_gemv_q8_0_4x8_q8_0 (#30585) 2026-10-29 08:21:05 +08:00
Abhijith Jajam d0ff6317cc docs : fix DSpark confidence threshold flag name (--spec-draft-p-min) (#31178) 2026-10-29 08:17:38 +08:00
Pascal cf24317bea common, mtmd: bound control vector layer indices and M-RoPE position loops (#31336)
* 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.
2026-10-28 23:15:01 +01:00
Hongqiang WangandLi He 74a4fb0037 opencl: add tiled q4_k GEMM for adreno (#30412)
* 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>
2026-10-28 14:55:28 -07:00
5220213255 hexagon: add conv2d support (#30995)
* 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>
2026-10-28 13:43:40 -07:00
Bartowski c3df4b7e22 Speed up IQ4 quant creation with pre-computed halfway points (#31201) 2026-10-28 14:35:56 -04:00
anubhavfujitsu ac8e1694d8 ggml-cpu: SVE 256 bit implementation of gemm_q6_K_8x8_q8_K kernel (#27472)
* q6_k GEMM working

* cleanup

* review comments

* review comment addressed - removed pg16_8

* incorporated the review comments - replaced switch with if and other points

* marking q8_k_blocklen UNUSEDto disable warning for non-i8mm cases
2026-10-29 00:01:37 +08:00
abhijain1204fujitsuandAnubhav Srivastava c02b5c641e ggml-cpu: SVE 256-bit Implementation of mxfp4 gemv kernel (#28141)
* 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>
2026-10-28 23:58:20 +08:00
Aaron Teo ff0d396e3d ci(s390x): switch to Ubuntu 26.04 runners (#31317) 2026-10-28 23:32:38 +08:00
Yonghye Kwon 838093f192 llama : remove simple batch splitting for DSV4 (#31300)
* 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
2026-10-28 23:25:47 +08:00
Georgi Gerganov 604cc43023 codeowners : refresh teams and maintainer PRs link (#31316)
* 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)
2026-10-28 17:00:54 +02:00