10908 Commits
Author SHA1 Message Date
Foad Abo DahoodandGitHub aac810230f metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024 (#28692)
* 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.
b10908
2026-09-11 12:30:20 +03:00
Logan ChuandGitHub 5cdd3d1dad model : fix MTP context kv cache allocation for deepseek2, glm4moe, … (#28630)
* model : fix MTP context kv cache allocation for deepseek2, glm4moe, cohere2moe architectures (#28626)

* model: add inverse architecture gating and comprehensive architecture testing for mtp layer filtering

* model : slim NextN filter comment, drop test-llama-archs changes
b10907
2026-09-11 12:02:31 +03:00
Jesus GulfoandGitHub b0dcb8192b server: fix speculation after an image (#28715)
* 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
b10906
2026-09-11 11:33:26 +03:00
16378d93f9 CUDA/HIP: Flash Attention tuning (gfx1201) (#28102)
* 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>
b10905
2026-09-11 09:58:20 +02:00
Sigbjørn SkjæretandGitHub 451b89bae0 ci : key cache to sanitizer matrix (#28708) 2026-09-11 07:56:18 +02:00
Jeff BolzandGitHub 481c65f091 vulkan: fix data race and OOB access in argsort(large) (#28705)
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.
b10903
2026-09-11 08:44:13 +03:00
shaofeiqiandGitHub df03399b88 opencl: add A8 Q4_0 mm binary kernel support (#28268) b10902 2026-09-10 11:25:40 -07:00
Jeff BolzandGitHub 28ff095829 vulkan: use CPU writes in ggml_backend_vk_cpy_tensor_async if the context is idle (#28618) b10901 2026-09-10 20:22:46 +03:00
Jeff BolzandGitHub 50182a53fa vulkan: use add_alloc_dep to enable topk_moe fusion for prefill (#28422) b10900 2026-09-10 20:21:29 +03:00
Jeff BolzandGitHub 6788edb4f3 vulkan: small M matrix optimizations for qwen (#28457)
* 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.
b10899
2026-09-10 20:20:18 +03:00
Sigbjørn SkjæretandGitHub 52d4268656 ci : add self-hosted-gpu-cuda and server-sanitize to hf-jobs (#28693) 2026-09-10 18:11:32 +02:00
shivamkumard-ctrlandGitHub 18c17b4d66 ci : Update WoA CUDA 13.4 release to use 13.4.1 GA redistributables (#28687)
- Move Windows ARM64 CUDA 13.4 builds from the Developer Preview archives to the 13.4.1 GA redistributables
b10897
2026-09-10 18:10:55 +02:00
Jesus GulfoandGitHub fa67698187 spec: fix failed to decode mtmd chunk with DFlash (#28587)
* 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
b10896
2026-09-10 17:10:55 +02:00
Daniel BeveniusandGitHub 41fc7584f0 scripts : use sed instead of grep for version parsing [no ci] (#28700)
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.
2026-09-10 15:44:40 +02:00
Iggy JacksonandGitHub d344123fe2 models: clean up some dead switch branches in old models (#28669)
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.
b10894
2026-09-10 16:09:35 +03:00
Gaurav GargandGitHub c32d1dabe8 tests : increase tolerance for Add fusion tests (#28691) b10893 2026-09-10 15:12:40 +03:00
Georgi GerganovandGitHub e5a8d439ce tests : drop SYCL special-casing in test-backend-ops.cpp (#28688) b10892 2026-09-10 13:43:43 +03:00
Julian PscheidandGitHub 3ff67eb43d vulkan: fall back to shared-memory reduction for dmmv on PowerVR (#28341)
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).
b10891
2026-09-10 13:42:23 +03:00
8c322d5bc4 convert : expand Nemotron H conversion fix (#28689)
* override function for n_h_l

* narrow change for extracting nested attribute

* simpler change; combines has_moe_params

* Apply suggestion from @CISC

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

---------

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
2026-09-10 13:41:57 +03:00
311d4211bf memory : avoid allocating V cache for indexer (it's not used) (#28330)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
b10889
2026-09-10 10:55:46 +02:00
72797e8919 vulkan : add command-buffer debug labels for GPU profilers (#28101)
* vulkan : add command-buffer debug labels for GPU profilers

Co-authored-by: gabby-zy <z2262718160@gmail.com>
Assisted-by: Claude Code

* vulkan : close the queue debug label with the label struct

---------

Co-authored-by: gabby-zy <z2262718160@gmail.com>
b10888
2026-09-10 10:06:07 +03:00
Aaron TeoandGitHub 4ea6d1bb6d ggml-cpu(s390x): add repack support for q4_0 (#28667)
ggml-cpu: clean comments

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
b10887
2026-09-10 09:51:17 +03:00
Aaron TeoandGitHub f1b6fbf35c ggml-cpu(s390x): add Q1_0 vector intrinsic support (#28606)
* ggml-cpu: add `ggml_vec_dot_q1_0_q8_0` support

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: clean up variable naming for understanding

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* docs: update support for Q1_0

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
b10886
2026-09-10 09:50:28 +03:00
Aaron TeoandGitHub d7e86430a7 model: fix all granite family parameter counts (#28643)
* model: fix all granite family parameter counts

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* model: fix additional include, add missing `A` prefix for active experts

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* model: fix code alignment, rm unused 40 block case

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
b10885
2026-09-10 09:49:56 +03:00
EveandGitHub 434ddbbc0e ci: fix sanitizer tests (#28583) b10884 2026-09-09 19:46:28 +00:00
Ruben OrtlamandGitHub 91f6a6cf36 vulkan: use spec constant for matrix matrix multiplication A-type (#25773)
* 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
b10883
2026-09-09 18:53:39 +02:00
6d9c82ea2b hexagon: rope updates (#28628)
* hexagon: vectorize RoPE theta cache on v75

* hexagon: vectorize MROPE/IMROPE theta pick

* hexagon: tighten NEOX RoPE rotate and aligned tail copy

* hex-rope: use inplace rope for all scenarios

* hex-rope: remove ctx->spad usage and legacy timers

* hex-rope: add kernel params and enforce vtcm reqs at the host

* hex-rope: cleanup unused params and tighten the mode checks

* hex-rope: add  missing ops header

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-09-09 08:40:24 -07:00
Masato NakasakaandGitHub 22397c31a0 vulkan: Convert FILL to distribute workgroups in 2D to avoid exceeding maxComputeWorkGroupCount (#28592)
* divide workload to 2D

This is to workaround FILL exceeding maxComputeWorkGroupCount for Intel GPUs on Qwen 3.8 flash next

* minor change

* Fixed comment
b10881
2026-09-09 16:54:15 +02:00
Sigbjørn SkjæretandGitHub 4b98ab805a py : lower numpy to 2.2.6 (#28654)
* Revert "py : bump numpy to 2.4.6 (#28649)"

This reverts commit 9cf3bf256b.

* bump numpy to 2.2.6
2026-09-09 15:56:27 +02:00
Sigbjørn SkjæretandGitHub 9cf3bf256b py : bump numpy to 2.4.6 (#28649) 2026-09-09 14:59:37 +02:00
4850c7727f llama : use int32_t for llama_sampler_chain_n return type (#28631)
Contributes to #4574

Co-authored-by: linsen <linsen@insta360.com>
b10878
2026-09-09 15:27:25 +03:00
d4abd573f6 CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#28552)
Recreated from #24546

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>

* CUDA: pick MMQ tile size against ncols_opt set on the host side

Assisted-by: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_011SYPfRhKoUpU3gMsGxq6go

---------

Co-authored-by: ravel7524 <58877666+ravel7524@users.noreply.github.com>
Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
b10877
2026-09-09 13:25:54 +02:00
5a4d0fecae CUDA: replace GGML_FA_ALL_QUANTS with GGML_FA_QUANTS, more control over what is compiled (#28079)
* CUDA: add configurable FA quant combinations

Assisted-by: Codex

* remove all flags but , add runtime fallback with warning for uncompiled combination

* Update docs/build.md

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>

* apply code review comments

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
b10876
2026-09-09 12:50:08 +02:00
Aaron TeoandGitHub 14a9d09f75 args: officially deprecate --mmap|mlock|dio (#28334)
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
b10875
2026-09-09 18:36:27 +08:00
Aaron TeoandGitHub e2d2c0d6aa model: fix granite3 moe unknown parameter count (#28632)
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
b10874
2026-09-09 16:46:22 +08:00
Xuan-Son NguyenandGitHub 6de9cdb26b mtmd: propagate video ID to bitmap (#28601) b10873 2026-09-09 10:37:36 +02:00
PascalandGitHub b31b71f3a0 jinja: treat a null left operand of in as a plain lookup (#28620)
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.
b10872
2026-09-09 10:08:27 +03:00
WakeUpMortyandGitHub df750f76bb vulkan: add dedicated iq4_xs mat-vec shader (#28426)
* vulkan: add dedicated iq4_xs mat-vec shader

Dedicated mul_mat_vec_iq4_xs for the dmmv path, replacing the generic fallback. ~+6-17% token generation on RDNA4 depending on model.

Assisted-by: Pi agent with Qwen3.8 27B

* vulkan iq4_xs: remove dead n_it unroll branch

Remove the n_it <= 8 experimental branch that attempted to fully unroll
the block loop. Since n_it is a runtime value, [[unroll]] is ignored by
the compiler, making both branches equivalent. Kept the simple loop
matching mul_mat_vec_iq3_s.comp.
b10871
2026-09-09 09:26:36 +03:00
Jiang, FishandGitHub 1945e09203 vulkan: add f16 B-type matmul pipelines and warp tile size tuning for Intel coopmat1 (#27471)
* vulkan: add f16 B-type matmul pipelines and warp tile size tuning for Intel coopmat1

* simplify mmp selection in mul_mat_id per review comment

* vulkan: enable f16 B-type pipeline for dense matmul on all vendors (MoE stays Intel-only)

* vulkan: add missing ocp_fp4 branches, restrict required_subgroup_size to Intel, fix mmp selection

* refine mmp selection in mul_mat_id

* add f16B pipeline selection just like q8_1

* update f16B pipeline selection in dense function
b10870
2026-09-09 09:03:32 +03:00
Georgi GerganovandGitHub 30b6a755e2 tests : use less threads for data initialization (#28325)
* tests : use 1 thread for data initialization

* cont : scale threads with number of elements

* cont : adjust
b10869
2026-09-09 07:50:38 +03:00
304665fe7a Add IQ type handling for MoE (#28476)
Co-authored-by: cwriter <cwriter@localhost>
b10868
2026-09-09 09:25:41 +08:00
Ruben OrtlamandGitHub f3f1a8f276 llama: disable lazy tensor loading by default on iGPUs (#28326)
* llama: add lazy mode auto, fix iGPU regression

* revert changes except disabling lazy load on iGPUs in AUTO
b10867
2026-09-08 18:05:09 +02:00
Sarah WuandGitHub 9113cc1880 ggml : fix msvc+clang ggml_vld1q_u32 (#28284) 2026-09-08 17:40:26 +03:00
uvosandGitHub d4389a4dd9 Revert "ggml-cuda : restore prop.integrated on HIP builds (#24233)" (#28604)
This reverts commit c7d8722922.
b10865
2026-09-08 16:19:53 +02:00
Foad Abo DahoodandGitHub 5d806aa257 server : apply checkpoint min-step eviction only when the checkpoint list is full (#28302)
The spacing eviction in create_checkpoint() keeps the oldest checkpoint and
erases every later one within checkpoint_min_step of it. For prompts shorter
than checkpoint_min_step this drops the checkpoint at n_tokens - 4 that the
next request resumes from, so hybrid/recurrent models re-prefill from the
previous checkpoint instead. Apply the spacing rule only once the list is at
n_ctx_checkpoints, and replace an existing checkpoint at the same n_tokens
instead of appending a duplicate.
b10864
2026-09-08 16:01:03 +03:00
Foad Abo DahoodandGitHub 88ada91c18 metal : fix idle threads in mul_mv_iq3_xxs for ne00 < 1024 (#28086)
* metal : fix half-idle simdgroup in kernel_mul_mv_iq3_xxs_f32 for ne00 < 1024

* metal : keep N_R0_IQ3_XXS = 4, dispatch a separate 8-row split kernel for ne00/32 < 32

The plain kernel is unchanged from master (4 rows per simdgroup, one thread per
chunk). The row-split mapping now lives in a separate kernel_mul_mv_iq3_xxs_f32_split
instantiation with N_R0_IQ3_XXS_SPLIT = 8, and the host selects it only when
ne00/32 < 32 and divides 32, so wide matrices keep the master kernel bit for bit.

* metal : select the iq3_xxs row split with a function constant instead of a separate kernel
b10863
2026-09-08 15:54:42 +03:00
Aman GuptaandGitHub 415e909d84 spec: single device drafter should create meta backend wrapper (#28390) 2026-09-08 20:44:33 +08:00
Sigbjørn SkjæretandGitHub 03fa73cb27 ci : disable npm gha cache (#28600)
* disable npm gha cache

* lies
2026-09-08 14:06:15 +02:00
Daniel BeveniusandGitHub 1744c6bde8 ci : add PYTEST_WORKERS=1 to fix server-self-hosted job (#28603)
* ci : add PYTEST_WORKERS=1 to fix server-self-hosted job

This commit adds the `PYTEST_WORKERS=1` environment variable to the
hf-jobs-t4-small:cuda13 runner steps.

This is an attempt to address CI failure of this job that I might have
introduced in Commit 42f0225fea
("server : use pytest-xdist for server tests (#28298)").

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/34126971262/job/101757819134

* apply same changes to server-metal steps
2026-09-08 13:36:03 +02:00
Pepper GrayandGitHub ca86fb222e llama : add missing headers (#28566)
* fix compile-error: add missing header

Bug: #28557
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28559
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28560
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28561
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28562
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28564
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

---------

Signed-off-by: Pepper Gray <hello@peppergray.xyz>
b10859
2026-09-08 12:59:53 +02:00