Commit Graph
10880 Commits
Author SHA1 Message Date
Lumpiasty 06a29462c1 ggml: add ggml_backend_free_scratch to drop Vulkan compute prealloc
Add an optional backend interface method free_scratch (with a public
ggml_backend_free_scratch wrapper) that frees transient/scratch device
memory a backend holds outside of any allocated buffer, keeping the
backend usable - the scratch is reallocated lazily on the next compute.

Implement it for the Vulkan backend (ggml_backend_vk_free_scratch): free
the prealloc_x/y/split_k/add_rms_partials and sync_staging device buffers
and reset their sizes, so an idle/cold model does not hold the vision or
matmul compute preallocations in VRAM. All other backends leave the hook
null (no-op).

Assisted-by: Claude
2026-09-09 00:38:35 +02:00
Lumpiasty ec691cd027 llama: evict recurrent/SSM state on device release
The recurrent (SSM/conv) state of hybrid models (e.g. Qwen3.5) was left
resident when a model's device buffers were released for on-demand VRAM
sharing - llama_memory_recurrent::release_device_buffers() was a no-op
default. Implement it (and restore_device_buffers) with the same
capture-host-shadow / free / reallocate pattern as llama_kv_cache, so
llama_memory_hybrid now evicts both its attention KV and its recurrent
state. The state is read-write, so its shadow is recaptured on every
release.

Assisted-by: Claude
2026-09-09 00:36:41 +02:00
Lumpiasty 0d68afc4c1 server: also free the compute-graph scheduler on KV eviction
Under LLAMA_SLEEP_EVICT_KV a cold model still held the scheduler's worst-case
compute buffer (hundreds of MiB, e.g. ~700 MiB for gemma-26B) plus the resident
experts. release_device(evict_kv) now also frees the sched (sched.reset() +
sched_need_reserve), and restore_device() rebuilds it via sched_reserve(), so a
fully-evicted model holds essentially no VRAM (gemma cold: 7147 -> 51 MiB).
Trades a heavier re-warm (weights+KV H2D + sched reserve) for the freed VRAM;
weights-only mode is unchanged (fast switch, keeps KV+compute).

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty 4597050be9 server: restore weights/KV on wake, not at decode, so KV eviction is safe
With LLAMA_SLEEP_EVICT_KV the KV cache device buffers are freed when a model
goes cold. update_slots() touches the KV before the decode-time vram_ensure_warm
(e.g. SWA models create a checkpoint that reads the KV via ggml_backend_tensor_get),
so the KV must already be resident by then. Move the restore to the sleep-wake
handler (handle_sleeping_state(false)), which runs before update_slots, fixing a
GGML_ASSERT(buffer) crash on iSWA models (gemma) when KV eviction is enabled.

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty 857fbbfb39 llama: guard buffer iteration against released device buffers
release_device_weights()/release_device_buffers() leave a null buffer in
ctxs_bufs while the device memory is released for on-demand VRAM sharing.
The memory-breakdown / total_size / clear paths iterated these and called
ggml_backend_buffer_get_size()/get_type()/clear() on the null buffer, tripping
GGML_ASSERT(buffer) and aborting on shutdown of a cold model. Skip null buffers
in llama_model::memory_breakdown and llama_kv_cache::{memory_breakdown,
total_size,clear}.

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty 35d35d5c31 server: coordinate model load with the VRAM arbiter to avoid load-time OOM
Before uploading a model's weights, acquire the shared VRAM token (ring the
doorbell so any resident model releases first). Previously load uploaded weights
to VRAM before the arbiter was active, so loading a large model while another
(e.g. the warm 4B task model) held VRAM could exceed the budget and OOM.

- add vram_arena_open() (idempotent flock/doorbell setup) and
  vram_acquire_for_load(), called from load_model() before
  common_init_from_params().
- a coordinated load now stays warm holding the token and serves its first
  request without a re-warm (drop the init-time go_cold cold-start).

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty ad657136a2 server: keep the sleep-wake path active when the VRAM arbiter is enabled
The request handler only calls wait_until_no_sleep() (which wakes a server out
of its sleeping state) when sleep_idle_seconds >= 0. But the VRAM arbiter's
cross-process doorbell can put a server to sleep even when idle-sleep is
disabled, so without this a doorbell-slept server would never wake and requests
to it would hang until timeout.

Do not bypass the wake path when LLAMA_SLEEP_VRAM_ONLY is set, so the arbiter no
longer depends on --sleep-idle-seconds being configured.

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty 811b4fab77 server: optionally evict the KV cache too (Phase 2 of VRAM sharing)
Extend on-demand device residency to the KV cache so that when a model's KV
plus another model would not fit in VRAM, the KV can also be evicted to a host
shadow (D2H on release, H2D on restore) instead of only the weights.

- llama_memory_i: add release_device_buffers()/restore_device_buffers()
  (default no-op). Implemented in llama_kv_cache (D2H shadow of the live
  ctxs_bufs, freed and reallocated like the weights); llama_memory_hybrid and
  llama_kv_cache_iswa delegate to their child caches.
- llama_context::release_device(evict_kv): also evict the memory's device
  buffers when requested; restore_device() rebuilds them. Public API
  llama_context_release_device gains an evict_kv flag.
- server: LLAMA_SLEEP_EVICT_KV=1 enables it. Off by default (weights-only),
  since the KV shadow adds a D2H/H2D copy of the live cache each cycle.

Validated on RX 580 (Vulkan), 4B @ 32k ctx: weights-only cold VRAM 1750 MB
(KV stays); weights+KV cold VRAM 726 MB (KV freed, ~1 GB reclaimed). KV
survives the round-trip: prompt cache reused after the cycle (prompt_n 4 vs
42), correct output.

Assisted-by: Claude
2026-09-09 00:36:40 +02:00
Lumpiasty 811eb55e12 server: on-demand VRAM sharing to time-share one GPU between models
Add release/restore of a model's GPU weight buffers (keeping a host shadow
and the KV cache) so several always-loaded llama-server processes can
time-share a single GPU without reloading or losing the prompt cache.

- llama-model: release_device_weights()/restore_device_weights() capture a
  compact host shadow (stable iteration order, view-skipping) and free then
  realloc the device weight buffers; weights_resident() query.
- llama-context: release_device()/restore_device() wrappers; decode() auto-
  restores; public C API llama_context_release_device/restore_device.
- server: LLAMA_SLEEP_VRAM_ONLY makes idle-sleep release only the VRAM weights
  (not a full unload/reload). A cross-process flock token in LLAMA_VRAM_ARENA
  enforces "resident iff holds token"; an inotify doorbell forces the holder
  to release on contention. The warden thread only touches the task queue, so
  releases run on the loop thread and never race a decode.

Validated on RX 580 (Vulkan): two models share 8GB, never both resident,
correct output under contention, KV cache preserved (no re-prefill).

Assisted-by: Claude
2026-09-09 00:36:34 +02:00
Lumpiasty 786b4077a4 ggml-vulkan: tiled transpose fast-path for concat with transposed source
The generic concat shader reads a transposed source (nb[1]==type_size) with an
uncoalesced stride, which is catastrophically slow on discrete GPUs (~5.6ms for
a 16MB concat on the RX 580 vs ~130us of memory bandwidth), a ~9% prefill
hotspot on Qwen3.5-4B (delta-net state concat).

Add a fast path for concat along dim 0 where one source is stored transposed and
the other source + dst are contiguous along dim 0: copy the contiguous source
with copy.comp and transpose the other source into the matching dst sub-region
with the existing tiled copy_transpose shader (shared-memory 32x32 transpose,
coalesced read+write). Reuses pipeline_cpy_* / pipeline_cpy_transpose_* with
custom push constants + doffset, no new shader. Falls back to the generic path
otherwise (gated on type/shape/contiguity + 16-bit doffset bound).

Assisted-by: opencode
2026-09-09 00:32:14 +02:00
LumpiastyandClaude Opus 5 6d483cc3c1 vulkan: record why flash-attn shmem staging stays off for GCN
Staging K/V through shared memory looks like an obvious win on GCN: without it
each rowgroup re-reads the whole K/V block from global memory, and with row_split
4 that pulls one 16KB block through a 16KB vector L1 four times.

Measured, it loses: -6.7% pp2048 at depth 16k and -7.4% at 32k on Polaris at head
size 128. The kvsh stride of D/4+1 dwords is 4 mod 32, which costs an 8-way LDS
bank conflict on wave64 - the +1 padding is tuned for warp32 - and the extra
shared memory eats occupancy this shader is already short of.

Comment only, no behaviour change. Leaving a note so the next person does not
spend a GPU on it.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
2026-09-09 00:21:39 +02:00
Lumpiasty c2d0492791 ggml-vulkan: enable flash-attn mask_opt for GCN large head sizes
mask_opt was disabled on AMD GCN for head sizes <= 256, but it is
beneficial there in high-context prefill: on fully-visible causal
blocks it skips the per-block mask load+add, and it skips fully-masked
blocks entirely. The attention op on GCN is compute-bound on the
softmax path (no matrix cores), so this cuts real work.

Verified lossless (perplexity bit-identical with it on vs off) and a
measured prefill win on Qwen3.5-35B-A3B (head_dim 256) on an RX 580:
pp2048 unchanged at short context, +8% @ 16k, +12% @ 32k, growing with
depth. Enable for GCN when HSK/HSV >= 256; the existing large-mask
conditions keep it off for decode.

Assisted-by: Claude
2026-09-09 00:21:39 +02:00
LumpiastyandClaude Opus 5 dbc449afd0 sched: skip the MoE routing-ids readback when the batch uses every expert
Offloaded expert weights are uploaded per split, and to decide which experts to
upload the scheduler reads the routing ids back to the host. The ids are produced
on the same device we are about to upload to, so the readback forces a full
pipeline flush - 47 of them per eval on a 48-layer MoE.

Once the batch draws enough experts the readback stops telling us anything. At
2048 tokens x 10 experts over 256 experts every expert comes back used, so the
bitset is all ones and the copy is a single whole-tensor range anyway.

Skip the readback when the batch guarantees that. Uploading an expert that no id
selects cannot change the result, since mul_mat_id only reads the rows the ids
point at, so this stays exact. Decode is unaffected: n_ids there is the number of
experts per token, far below the threshold, so it keeps the bitset path.

Measured on an RX 580 (Vulkan, Polaris) running Laguna-S-2.1 118B IQ2_M with all
experts host-resident: pp2048 +4.7% at depth 32k, +0.9% at 16k, tg32 +3%.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
2026-09-09 00:21:39 +02:00
Ruben Ortlam 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
2026-09-08 18:05:09 +02:00
Sarah Wu 9113cc1880 ggml : fix msvc+clang ggml_vld1q_u32 (#28284) 2026-09-08 17:40:26 +03:00
uvos d4389a4dd9 Revert "ggml-cuda : restore prop.integrated on HIP builds (#24233)" (#28604)
This reverts commit c7d8722922.
2026-09-08 16:19:53 +02:00
Foad Abo Dahood 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.
2026-09-08 16:01:03 +03:00
Foad Abo Dahood 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
2026-09-08 15:54:42 +03:00
Aman Gupta 415e909d84 spec: single device drafter should create meta backend wrapper (#28390) 2026-09-08 20:44:33 +08:00
Sigbjørn Skjæret 03fa73cb27 ci : disable npm gha cache (#28600)
* disable npm gha cache

* lies
2026-09-08 14:06:15 +02:00
Daniel Bevenius 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 Gray 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>
2026-09-08 12:59:53 +02:00
Ankit Khandelwal 64e9bceb2c vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL (#27220)
* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL

* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL

- implement fusion in unary.comp behind UNARY_MUL_FUSION ifdef,
  specialized pipelines per op instead of runtime branching
- fuse adjacent nodes only, ordering handled by graph_optimize
- drop runtime consumer scan and pending_unary_mul deferral

* vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL

1. GELU: gelu_mul_f32/f16 pipelines registered, CREATE_UNARY_MUL(gelu), GELU in dispatch + fuse gate + perf fusion name
2. Renamed/moved: gate is now ggml_vk_can_fuse_unary_mul(cgraph, unary_idx, mul_idx), placed with the other can-fuse helpers
3. norepeat both variants: each op gets plain (spec {0}) + _norepeat (spec {1}) pipelines from the same SPIR-V, selected via ggml_are_same_shape(src0, src1); the shape gate now allows broadcast (other dims equal-or-1)
4. graph_optimize: lambda deleted; standard "// UNARY + MUL: pull the consuming MUL forward" block added alongside the SSM_CONV/ROPE/MUL_MAT reorderings, with the same "other src must be weights or already processed" readiness check

* vulkan : align unary_mul fusion with binary kernel layout, relax gelu test tolerance

- schedule the fused kernel like mul.comp (256 threads x 2 unrolled
  iterations), recovering a 10-18% prompt-processing regression
- allow 5e-7 f32 error for gelu_mul: the shader evaluates gelu with an
  exp-based tanh identity while the CPU reference uses tanhf (~1 ulp)

* vulkan : use ggml_can_repeat in UNARY+MUL fusion shape check

The fused kernel indexes src1 via per-dim fastmod (generic_binary_head.glsl),
which is exact whenever the other operand tiles into the unary result -- not
just when its dims are equal or 1. Replace the hand-rolled loop with
ggml_can_repeat(other, unary) so the check matches the kernel's actual
capability and reuses the standard helper. Argument order matters: reversed,
it would wrongly admit graphs where the unary result is mul->src[1] and the
other operand is larger, producing truncated output.

Also add a rep_ne0 layout to the fused unary+mul backend tests covering a
non-1 repeat factor along dim 0.

* vulkan : fuse UNARY+MUL pairs separated by zero-compute nodes

gemma4's per-layer embedding gating builds gelu -> view_2d_slice -> mul,
where the intervening view is a zero-compute node aliasing an input that
was computed much earlier. Strict adjacency requirements meant neither
CUDA nor the vulkan unary+mul fusion handled this pattern.

Extend ggml_vk_graph_optimize to detect a UNARY whose consuming MUL is
separated only by unscheduled zero-compute nodes (GGML_OP_NONE, VIEW,
RESHAPE, TRANSPOSE, PERMUTE) and schedule those nodes ahead of the pair,
making it adjacent so the existing fusion applies. The reorder is guarded
by ggml_vk_can_fuse_unary_mul, a source-availability check for every
interleaved node, and the protected fusion patterns (topk_moe*, snake);
if fusion is later rejected the reordered graph still executes correctly,
just unfused.

Add a view_mid layout to the fused unary+mul backend tests replicating
the gemma4 pattern.

* vulkan : support OP-on-B in UNARY+MUL fusion

Some models apply the unary activation to the smaller MUL operand, e.g.
qwen3next/qwen35moe shared-expert gating builds ffn_shexp * sigmoid(gate)
with a [1,n_tokens] gate tensor. This shape was correctly rejected before:
the fused kernel derives its iteration extent from the unary tensor and
would leave most of the destination unwritten, and the generic same-shape
requirement in ggml_can_fuse blocked the pair outright.

Add UNARY_MUL_B_FUSION shader variants computing dst = src0 * OP(src1):
the OP operand rides the existing per-dim fastmod indexing, while the
iteration extent now comes from mul. Route {UNARY, MUL} pairs through a
local can-fuse variant that drops the generic same-shape rule and instead
requires the unary result to tile into mul->src[0] (ggml_can_repeat);
pairs with the unary as src0 keep the previous direction check, and
equal-shape pairs keep using the original pipelines.

Add a "gate" layout to the fused unary+mul backend tests covering the
shared-expert gate shape for gelu/sigmoid/silu/softplus in f32 and f16.

* vulkan : fold unary+mul view-hoisting into graph_optimize dep checks

Replace the dedicated UNARY + EMPTY* + MUL scanning block with two small
extensions to the existing scheduling logic:

- a consuming MUL may now join its in-set UNARY across a gap of unused
  zero-compute nodes (NONE/VIEW/RESHAPE/TRANSPOSE/PERMUTE), instead of
  requiring strict adjacency
- while doing so, such zero-compute blockers are ignored for this pair

Fusion validity is still decided later by ggml_vk_can_fuse at dispatch
time, so a rejected pair simply executes adjacent-but-unfused. Note the
relaxation must stay scoped to this pattern: exempting zero-compute
blockers globally reproduces silent output corruption on gemma3n.

* vulkan : select unary_mul OP-on-B via specialization constant

Replace the UNARY_MUL_B_FUSION compile-time shader variants with an
op_on_b specialization constant on the existing unary_mul SPIR-V,
mirroring how the norepeat flag is handled. The four {op}_mul_b_{f32,f16}
shader artifacts are gone - the OP-on-B pipelines reuse the base SPIR-V
with two-entry {norepeat, op_on_b} spec lists - and the duplicated store
expression is collapsed into a single runtime branch that the driver
prunes per specialization.

The constant is declared only under UNARY_MUL_FUSION so every other
binary pipeline keeps its single-entry specialization list.

* vulkan : replace unary_mul pipeline switches with a lookup table

Collapse the four nested selection switches in ggml_vk_unary_mul into a
single indexed lookup against a pipeline_unary_mul[4][2][2][2] table
([unary op][f16][norepeat][op_on_b]), whose trailing dims mirror the
{norepeat, op_on_b} spec constant list. The op axis uses a small shared
index helper that also replaces the switch in ggml_vk_can_fuse_unary_mul,
making it the only place that maps ops to the table.

Pipeline names are unchanged. Adding another supported op now requires
one macro invocation line and one helper case instead of edits in four
separate switches.

* vulkan : use ggml_can_fuse_subgraph for unary_mul pairs

Replace the hand-rolled pair validation in ggml_vk_can_fuse_unary_mul_pair
(bounds, op match, compute flags, single-use elision) with the shared
ggml_can_fuse_subgraph helper; backend-specific shape/type rules remain in
ggml_vk_can_fuse_unary_mul. Unlike ggml_can_fuse, the subgraph helper has
no same-shape requirement, so it covers both operand slots including
OP-on-B gates, and additionally rejects intermediates flagged as graph
outputs and validates view-source confinement.

The outputs parameter takes absolute node indices into the cgraph.

* Fix Whitespace

* vulkan : drop redundant unary_mul gap check in graph_optimize

The zero-compute nodes separating a UNARY from its consuming MUL are
already scheduled ahead of the pair by pass 2 of an earlier
optimization window, so the scoped gap tolerance added for this pattern
is unreachable in practice - disabling it leaves gemma-3n dispatch
counts unchanged (841 GELU_MUL per pass). Remove the flag, the empty
blocker exemption, and the now-unused gap helper, restoring the strict
adjacency requirement of the UNARY -> MUL pull-forward.

Keep the relaxation scoped out entirely: generalizing "zero-compute
nodes never block" beyond this pattern previously reproduced silent
output corruption on gemma3n.

* vulkan: fix whitespace (tab in indent)

* vulkan: fix whitespace (extra blank line)

* vulkan : move op_on_b spec constant to unary.comp

op_on_b is only used by the fused unary*mul path. Keep
generic_binary_head.glsl generic by defining it in unary.comp
instead. Same constant_id=1 and guard, no functional change.

* vulkan : make RMS_NORM/UNARY fusion gap-tolerant for views

Strict j==c+1 blocked RMS_NORM->MUL and UNARY->MUL when a
VIEW sits between (e.g. rms_norm -> view -> mul). Allow
c==back() with an empty-or-scheduled gap, matching the
review suggestion to check src linkage instead of adjacency.
Scoped to the two blessed pairs; safe because gaps can only
contain zero-compute nodes.

* vulkan : trim comments in UNARY+MUL fusion

Assisted-by: Muse Spark
2026-09-08 09:35:02 +02:00
miyan f014bfef8b Fix Vulkan-Hpp handle usage on 32-bit targets. (#22892)
On 32-bit platforms, Vulkan non-dispatchable handles such as VkBuffer are
represented as uint64_t, and Vulkan-Hpp disables implicit conversions for
type safety. This exposes two issues in ggml-vulkan:

1. vk::Buffer is streamed directly into std::ostream in debug/memory logs.
2. vk::Buffer is cast to VkBuffer before being passed to Vulkan-Hpp
   CommandBuffer::copyBuffer APIs.

Fix these by add the operator<< for vk::Buffer, and
by passing vk::Buffer directly to Vulkan-Hpp copyBuffer calls.
2026-09-08 09:34:12 +02:00
Piotr Wilkin (ilintar) 895c045fd1 chat : split specialized parsers into common/parsers (#27764)
* chat : split specialized parsers into common/parsers

Move the 14 dedicated template parsers out of chat.cpp into one file each under
common/parsers, mirroring the src/models split. chat.cpp keeps the template
detection in common_chat_try_specialized_template() and drops from 3915 to 1513
lines.

common/parsers/parsers.h holds the shared helpers and one declaration per
parser. foreach_function/foreach_parameter become inline there since nothing in
chat.cpp uses them any more; common_chat_template_direct_apply_impl and
common_chat_template_generation_prompt_impl lose static and carry their default
arguments in the header. Parser-specific helpers move with their parser:
is_lfm2_template, deepseek_v4_sort_tool_results and the gemma4 turn builder.

No functional change.

Assisted-by: Claude Opus 5

* chat : enumerate parser sources instead of globbing

file(GLOB) does not re-run CMake when a source file is added or removed, so an
incremental build silently keeps building the old set. List the parsers in
common/parsers/sources.cmake and include it from common/CMakeLists.txt.

Assisted-by: Claude Opus 5

* split helpers, add newlines
2026-09-08 09:29:37 +03:00
lhez 7d701b5929 opencl: properly handle non-contiguous inputs to conv2d (#28503)
* opencl: fix conv2d non-contiguous strides

* opencl: format
2026-09-08 09:26:34 +03:00
Georgi Gerganov 5a6caa05fc ggml : update ggml_prec specification (#26675)
* ggml : update ggml_prec specification

[no ci]

* cont : add GGML_PREC_BF16

* cont : rework API

* cont : use new API

* cont : swap arg order

* cont : support for MUL_MAT_ID

* cont : fix accidental remove of "break;"

* cont : return bools, add doc TAG_GGML_PREC, clean-up

* cont : add search tag

* cont : ws
2026-09-08 09:06:24 +03:00
Frank Dai 9dcf84e5ae model : support Kimi-K3 recurrent-state rollback (#28466) 2026-09-08 11:31:48 +08:00
Todor Boinovski 050dde50c9 hexagon: add RELU and LEAKY_RELU ops (#28585)
* hexagon: add RELU op

* hexagon: add LEAKY_RELU op too
2026-09-07 17:04:25 -07:00
Sigbjørn Skjæret 67672dc5b7 ci : bump ty to 0.0.78 (#28548)
* bump ty to 0.0.78

* type fixes

* more type fixes

* add --exit-zero-on-warning

* remove Callable again
2026-09-07 21:10:06 +02:00
Pascal f114f91f9e tests : initialize the L2_NORM batch array (#28553)
* tests: bind the L2_NORM batch count to a local

GCC cannot prove the loop fills norms up to the index read after it
while the bound is a class member, so it reports a maybe uninitialized
use. Reading the count once into a local restores the tracking.

* tests: initialize the L2_NORM batch array

The read after the fill loop is only provably defined once the array
carries an initializer, which GCC 12 requires on the aarch64 Release
build where warnings are fatal.
2026-09-07 19:54:13 +02:00
Piotr Wilkin (ilintar) e71b80510c Revert "CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)" (#28551)
This reverts commit 0c963452ea.

Assisted-by: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Q7rfnkjzgnfhvJsdeXhdoH
2026-09-07 16:28:19 +02:00
Zhaolun Yin ccc3646c63 nix : update deprecated expressions (#28145)
* fixed warnings

* fixed nixfmt warning
2026-09-07 15:59:45 +02:00
Pascal c0b1871bc7 webgpu: format the GET_ROWS case block (#28542)
Brace on its own line and body indented one level, matching the
surrounding cases, so the webgpu clang-format check passes.
2026-09-07 15:55:14 +02:00
Xuan-Son NguyenandPascal 160bd031b2 server: fix LRU hang on multiple requests same model (#28539)
* server: fix LRU hang on multiple requests same model

* server: keep a queued model out of the victim pool until its waiters leave

A waiter that gave up while its model was still loading left the
model idle with no request behind it, and nothing recounted the free
slots, so a second request queued behind it stayed queued forever.
tick() was only driven by requests: join, claim and the end of a
proxied request.

Keep the queue entry alive after a successful claim so the model
coming up is never picked as a victim before its waiters use it, and
recount the slots on every status change and whenever a waiter
abandons the queue. The model is then evicted as soon as it comes up
with nobody left to serve.

---------

Co-authored-by: Pascal <admin@serveurperso.com>
2026-09-07 15:50:46 +02:00
Titaniumtown dbeb37548e sycl: add a batched L2_NORM kernel (#28222)
* sycl: add a batched L2_NORM kernel

* sycl: batch consecutive L2_NORM siblings in the graph dispatch

Measured on Intel Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, f16 KV,
npp=128 ntg=128 npl=2, GGML_SYCL profiler:

    L2_NORM dispatches       12480 -> 6240
    L2_NORM device time      68.77 -> 39.14 ms   (-43%)
    total device time        6782 -> 6748 ms     (-0.5%)
    wall decode t/s          flat

* tests: add L2_NORM_BATCH coverage
2026-09-07 15:24:14 +02:00
Kevin HopperandKevin Hopper 7a333e7240 vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST) (#26578)
* vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST)

CUDA has these ops from the DeepSeek-V4 merge and Metal gained them in
PR 26459. Vulkan was the last major backend running the unfused primitive
chain. On DeepSeek-V4-Flash the unfused Sinkhorn comb chain alone takes
about 32% of decode op time on gfx1151 (Strix Halo), spread over roughly
16k dispatches per token.

dsv4_hc_comb runs the full 20-iteration Sinkhorn in registers. A token's
4x4 comb matrix lives in 16 consecutive subgroup lanes, with idst in bits
0-1 and isrc in bits 2-3 to match the CPU reference layout, so
subgroupShuffleXor by 1|2 reduces rows and by 4|8 reduces columns. One
dispatch replaces about 137 strictly ordered node executions per site.
The shuffle masks never cross a 16-lane boundary, so a subgroup of size
64 packs 4 independent tokens.

dsv4_hc_pre and dsv4_hc_post handle the elementwise stream collapse and
fan-out, with per-token coefficients staged in shared memory.

GGML_VK_DISABLE_DSV4_HC disables all three ops. The _COMB, _PRE and
_POST variants gate each op independently so a single kernel can be
bisected against the unfused graph.

Adds eval cases at the production n_iter=20 across batch sizes that
cross subgroup and workgroup boundaries.

* vulkan: dsv4 hc review fixes

Drop the per-op env-var disables and device flags, the stride divisibility
check (ggml guarantees it) and the workgroup-count fallback in supports_op.
Trim the comb shader comments to the lane layout.

---------

Co-authored-by: Kevin Hopper <no-reply@maestro.press>
2026-09-07 15:24:03 +02:00
ravel7524andCarl Philipp Klemm 0c963452ea CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)
* adjust ncols_picker for routed MoE in mul_mat_q_case function

* Adding CDNA, RDNA2 and RDNA4

* fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support

* feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support

* refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files

* HIP: mmq: enable typical moe ncols on RDNA4

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
2026-09-07 15:22:42 +02:00
AuroraRAS 4735997382 ggml: add gfx90c HIP support (#26454)
* ggml: add gfx90c HIP support

* ggml: make gfx90c HIP support compliant with specifications
2026-09-07 15:21:42 +02:00
fairydreamingandStanisław Szymczyk d23c47f2a9 convert : refactor Hy4-preview conversion - move HC tensor mapping to the global map (#28451)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-09-07 15:20:58 +02:00
Pranesh Gonegandlaandpraneshgo 73ab7599b5 CUDA: branchless Q4_K/Q5_K unpack to speed up mmvq, L2 prefetch on DGX Spark (#26705)
* Update Q4_K and Q5_K to use branchless computation, which stops the scale unpack being re-executed for every column in mmvq, improving perf at batch sizes > 1

* Gating the change off from DGX Spark due to no gain

* Adding prefetch gated to Spark, making branchless change in Q4_K and Q5_K general and modifying switch points based on latest perf data

* Guard the mmvq L2 prefetch against MUSA as well as HIP

* Define the mmvq L2 prefetch only under the Spark guard

* Update switch point for Q4_K to accommodate more models

* Remove stale comments

* Add block_size to ggml_cuda_type_traits and create a separate mmvq_should_prefetch function

* Rename block_size to bs for cleaner indentation

* Fix build error on non-Spark CUDA arch with appropriate conditional around new function added

---------

Co-authored-by: praneshgo <227579474+praneshgo@users.noreply.github.com>
2026-09-07 19:36:58 +08:00
Jeff BolzandAiChiTuDouPian 0cae43063c vulkan: support type-aligned GET_ROWS (#28253)
* vulkan: fall back to CPU for GET_ROWS with misaligned offsets

The Vulkan GET_ROWS shader asserts when a tensor's backing-buffer offset
plus view_offs is misaligned w.r.t. minStorageBufferOffsetAlignment
(see init_pushconst_tensor_offsets). Previously this caused a hard crash
on models using ggml_view + ggml_get_rows (e.g. Qwen3-TTS, Qwen3-VL).

Return false from supports_op() in the misaligned case so the scheduler
falls back to CPU, matching the existing pattern for PAD_REFLECT_1D and
other unsupported op/shape combinations.

Repro: llama-tts -m Qwen3-TTS-*.gguf -mm mmproj-*.gguf -ngl 99
Crash: GGML_ASSERT(dst->op != GGML_OP_GET_ROWS || (a_offset == 0 && ...)) failed

* vulkan: trim comment for GET_ROWS misalign fallback

* vulkan: fix file corruption in gated_linear_attn struct

* vulkan: properly handle misaligned offsets in GET_ROWS quantized path

- get_rows_quant.comp was missing get_aoffset()/get_boffset()/get_doffset()
  calls that are already present in get_rows.comp, causing GGML_ASSERT crashes
  when GET_ROWS operates on views with non-zero view_offs, as produced by
  KV cache slices in Qwen3-TTS and Qwen3-VL.
- Remove the defensive misalignment GGML_ASSERT in init_pushconst_tensor_offsets
  for the binary push-constants specialization, since both get_rows.comp and
  get_rows_quant.comp now correctly apply per-tensor base offsets.
- Remove the workaround CPU fallback in supports_op() for GET_ROWS, since the
  Vulkan backend now handles misaligned offsets natively (no more bailout).
- Add backend test coverage with view_src0=true (ggml_view_4d into a padded
  tensor) for F32, F16, Q4_0, Q4_K, Q8_0, and I32 types, exercising both the
  non-quantized (get_rows.comp) and quantized (get_rows_quant.comp) paths
  with non-zero view_offs that reproduce the original Qwen3-TTS crash.

* tests: trim redundant comments in test_get_rows vs0 region

* tests: trim redundant comments in test_get_rows vs0 region (follow-up)

* vulkan: bind tensor base for binary ops, pass full view_offs via push constants

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, MUL, etc.),
bind the view_src base and pass the full view_offs divided by type_size via
push constant misalign_offsets. This avoids truncation when misalign_bytes is
not a multiple of quantized block size.

ggml_vk_tensor_subbuffer gains a use_view_offs parameter. When false, the
binding points to vk_tensor_offset (base) and size includes view_offs.
init_pushconst_tensor_offsets<binary> computes a/b/d_offset directly from
tensor->view_offs, which is always row-aligned and therefore exact.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).

* vulkan: bind aligned offset for binary ops, pass adjusted misalign via push constants

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.

ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types).

ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (target - adjusted_misalign) instead of the view_src base, keeping the
offset small enough for 16-bit/8-bit push constant fields.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).

* vulkan: bind aligned offset for binary ops, fix UMA offset mismatch

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.

Added ggml_vk_tensor_physical_offset to unify physical offset lookup across
UMA and non-UMA devices. On UMA, resolves via ggml_vk_host_get(tensor->data);
otherwise uses vk_tensor_offset(t) + t->view_offs. Both get_misalign_bytes and
the new ggml_vk_get_adjusted_misalign helper build on top of this function,
so buffer bindings and push constant offsets are always consistent regardless
of device memory model.

ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types) while
remaining small enough for 16-bit/8-bit push constant fields
(adjusted_misalign < lcm(align, type_size)).

ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (physical_offset - adjusted_misalign) on both UMA and discrete GPUs,
fixing a bug where the UMA host_get path previously skipped the adjusted
misalign binding and returned the target offset directly.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA GeForce RTX 5060 Ti).

* finish misalignment fix

* supports_op changes for openvino/webgpu

---------

Co-authored-by: AiChiTuDouPian <15327701848@qq.com>
2026-09-07 12:22:10 +02:00
Daniel Bevenius 1173700b9c examples : print ggml_version and ggml_commit in test-cmake [no ci] (#28538)
This commit adds the printing of the ggml version and commit to the
test-cmake example.

The motivation is just to be able to quickly verify that the correct
version of ggml is being used.

Example output:
```console
test-cmake] llama.cpp version: 0.4.0-dev, build: 10837 (5202104b5)
[test-cmake] ggml version: 0.23.0, commit: 5202104b5
[test-cmake] Initializing backend...
...
```
2026-09-07 12:11:40 +02:00
Sigbjørn Skjæret 5202104b59 caps : recheck typed content if template checks for string (#28511) 2026-09-07 09:14:32 +02:00
DevVexusandClaude Opus 5 9a7570587c convert : write explicit recurrent_layers for Qwen3-Next / Qwen3.5 (#28208)
Problem
- Loader prefers `<arch>.attention.recurrent_layers`, falls back to `full_attention_interval` if missing
- Converter only ever writes the interval. gguf-py has no constant/writer for the array
- Interval can only describe evenly spaced full-attention layers. Any non-uniform `layer_types` gets reconstructed wrong
- No error, no warning. Model loads, runs, wrong layers get wrong ops. Full-attn layers marked recurrent lose their KV cache
- Every published Qwen3.5 checkpoint is uniform so nobody's hit it yet

Repro
12 layers, periods 4/3/5:

    layer:  0 1 2 3 4 5 6 7 8 9 10 11
    actual: L L L F L L F L L L  L  F
    loader: L L L F L L L F L L  L  F
                        ^ ^

Layer 6 is full attn, loaded as recurrent. Layer 7 the reverse.
52-layer non-uniform stack: 15/52 mis-typed.

Fix
- `constants.py`: add `Keys.Attention.RECURRENT_LAYERS` (name already registered in llama-arch.cpp)
- `gguf_writer.py`: add `add_recurrent_layers()`, same shape as `add_rope_pattern()`
- `conversion/qwen.py`: emit array from `layer_types` in `Qwen3NextModel.set_gguf_parameters` (covers 3-Next, 3.5, 3.5-MoE)

Notes
- Array is padded with `false` for MTP blocks. `get_key_or_arr` checks length against `n_layer_all`, which includes MTP. Matches the fallback's `i < n_layer()` guard
- Interval is still written. Old builds only understand the interval
- `layer_types` length != `num_hidden_layers` now raises in converter instead of producing a GGUF that fails at load

Tested
- End-to-end on a 62-layer non-uniform Qwen3.8-27B (2 linear layers removed). Loader reads the array, 62 blocks, 0 mismatches. Without fix: interval fallback, mis-typed
- MTP padding NOT tested on a real MTP model. Reasoned from qwen35.cpp + get_key_or_arr. Would appreciate a check

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
2026-09-07 10:12:50 +03:00
Siavash Norouzi b74f590eaf ggml-cuda: fix divergent barrier in f16 flash attention (#27870)
* ggml-cuda: fix divergent barrier in f16 flash attention

* ggml-cuda: avoid duplicate metadata pointer setup
2026-09-07 09:23:21 +03:00
Aman Gupta 992cb503cd ggml: allow backend inputs to not create another split (#28387) 2026-09-07 09:10:40 +03:00
Jeff Bolz 9ac8c408a3 vulkan: rms_norm fusion opportunities (#28024)
Support RMS_NORM + MUL + ADD (+ MUL) and RMS_NORM + VIEW + SET_ROWS.
Extend ROPE + VIEW + SET_ROWS to support IMROPE.

Worth around 4% in gemma4 on my system.
2026-09-07 09:08:28 +03:00
Daniel Bevenius 2092353c8b ci : add container image checking and tagging (wip) (#28394)
This commit contains a suggestion for handling container images which
are currently not semver tagged, they only have build numbers in there
tags.

The proposed solution here is to first add a check to make sure that
there are container images built for the build number of the release and
if not fail the build. The container images are build nightly but they
can be triggered manually as well.
If the the container images check passes then the make-release workflow
will re-tag the images with the semver.
2026-09-07 07:23:39 +02:00
Anjielon 8fe90e1fbf vulkan: add TQ1_0 support (mm, mat-vec, mat-vec-id, dequant, get_rows) (#27765)
* vulkan: add TQ1_0 support (mm, mat-vec, dequant, get_rows)

* vulkan: pack TQ1_0 powers of 3 into a 32-bit constant

Replaces the constant array with a packed 32-bit value (7 bits per entry,
max 81 < 128) extracted with shift/mask, as suggested in review — avoids a
constant array that may not be kept in registers.

test-backend-ops on gfx1151: tq1_0 MUL_MAT 11/11, MUL_MAT_ID 6/6,
GET_ROWS 4/4, unchanged.

* vulkan: address review - shared TQ1_0 decode helpers, fix standalone dequant shader

Review feedback from jeffbolznv, all points:

- Move the packed-pow3 decode into shared helpers in types.glsl
  (tq1_0_byte_of / tq1_0_digit_of / tq1_0_trit) and use them from
  dequant_funcs.glsl, mul_mm_funcs.glsl, dequant_funcs_cm2.glsl and
  dequant_tq1_0.comp instead of repeating the logic. The cm2 path also
  drops its constant array for the packed-constant extraction.
- Translate all remaining comments to English.
- dequant_tq1_0.comp: use dequant_head.glsl. The shader previously declared
  its own single-field push constant while the pipeline is created with the
  5-field layout, so p.ne read the wrong field - confirmed broken, as
  suspected in review.
- Fix wg_denoms for the standalone dequant pipeline: one invocation decodes
  4 elements with local_size 256, so a workgroup covers 256*4 elements, not
  256*16. With the old value the dispatcher launched a quarter of the
  required workgroups.

Verified by temporarily forcing the dequant + f16 matmul path for TQ1_0
(hack not committed): test-backend-ops MUL_MAT passes through the rewritten
standalone shader, and the standard MUL_MAT / MUL_MAT_ID / GET_ROWS
tq1_0 cases still pass on Vulkan (AMD gfx1151).

* vulkan: address review — English comments, shared tq1_0_trit, trim TQ1_0 test cases

- mul_mat_vec_tq1_0.comp: drop leftover non-English comment and the local
  POW3_PACKED constant; all decode sites now call tq1_0_trit() from types.glsl
- types.glsl / dequant_funcs_cm2.glsl: ASCII-only, drop stale reviewer note
- test-backend-ops: remove the oversized MUL_MAT_ID case (432 MiB A tensor,
  ~172 GFLOP reference); move the two remaining ones next to the other
  backend-specific mul_mat_id one-offs and document why they are needed

* metal: decline TQ1_0 for GET_ROWS and mat-mul in supports_op

The new TQ1_0 cases in test-backend-ops exposed that the Metal backend
claimed support for GET_ROWS/MUL_MAT/MUL_MAT_ID with TQ1_0 sources while
having no such kernels (ggml_metal_library_compile_pipeline aborted on the
missing kernel_get_rows_tq1_0). Decline the type so the ops fall back to
the CPU, matching the existing NVFP4 handling on the same lines.

Assisted-by: Claude Fable 5

* vulkan: trim the TQ1_0 comments

Addresses @0cc4m's review: keep only what the code does not already say.

Removed the block-format recaps (the layout is right there in the struct) and
the step-by-step decode walkthrough. Kept the two facts a reader cannot infer:
the 8-bit truncation is part of the format, not an optimisation, and the powers
of 3 are packed into one uint so they do not end up in a constant array that
may miss the registers.

No functional change.

* vulkan: address review — trim comments, fold Metal check, drop unused _v

Per @0cc4m's review:

- dequant_funcs.glsl, dequant_funcs_cm2.glsl: drop the "see types.glsl"
  pointers — they apply to every quant and say nothing specific.
- dequant_tq1_0.comp: drop the wg_denoms note. It is a precondition, not
  information.
- mul_mm_funcs.glsl: same pointer removed.
- types.glsl: the comment on tq1_0_trit is down to the one fact the code
  cannot show — the 8-bit truncation is part of the format, matching the C
  reference, not an optimisation.
- dequant_funcs_cm2.glsl: removed dequantFuncTQ1_0_v and its define. You were
  right that it is optional: it wrapped four scalar decodes and vectorised
  nothing, and mul_mm_cm2.comp already guards the path with
  `#if defined(dequantFuncA_v)` (DATA_A_F32 omits it the same way).
- ggml-metal-device.m: folded TQ1_0 into the existing NVFP4 check instead of a
  separate block, and dropped both comments.
- test-backend-ops.cpp: the two mul_mat_id cases stay — they cover the
  block-stride loop and the per-expert base offset that k == 256 alone never
  reaches — but the comment is now one line instead of five.

Kept: the one-line labels on the three block regions in mul_mat_vec_tq1_0.comp
and on tq1_0_byte_of(). Those state the 5-trits-per-byte packing, which the
loop bounds do not show. Happy to remove them too if you prefer.

Re-verified on AMD gfx1151 (Vulkan), test-backend-ops, 2/2 backends passed:
MUL_MAT 9 TQ1_0 cases, MUL_MAT_ID 5, GET_ROWS 4 — all OK, no failures.
The coopmat2 path is unchanged apart from the removed _v define.
2026-09-07 06:35:30 +02:00