Commit Graph
10899 Commits
Author SHA1 Message Date
Sigbjørn Skjæret c7af7092b8 key cache to sanitizer matrix 2026-09-10 18:29:01 +02:00
Sigbjørn Skjæret 52d4268656 ci : add self-hosted-gpu-cuda and server-sanitize to hf-jobs (#28693) 2026-09-10 18:11:32 +02:00
shivamkumard-ctrl 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 Gulfo 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 Bevenius 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 Jackson 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 Garg c32d1dabe8 tests : increase tolerance for Add fusion tests (#28691) b10893 2026-09-10 15:12:40 +03:00
Georgi Gerganov e5a8d439ce tests : drop SYCL special-casing in test-backend-ops.cpp (#28688) b10892 2026-09-10 13:43:43 +03:00
Julian Pscheid 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
Kartik GuliaandSigbjørn Skjæret 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 <[email protected]>

---------

Co-authored-by: Sigbjørn Skjæret <[email protected]>
2026-09-10 13:41:57 +03:00
fairydreamingandStanisław Szymczyk 311d4211bf memory : avoid allocating V cache for indexer (it's not used) (#28330)
Co-authored-by: Stanisław Szymczyk <[email protected]>
b10889
2026-09-10 10:55:46 +02:00
Cordell Blanchardandgabby-zy 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 <[email protected]>
Assisted-by: Claude Code

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

---------

Co-authored-by: gabby-zy <[email protected]>
b10888
2026-09-10 10:06:07 +03:00
Aaron Teo 4ea6d1bb6d ggml-cpu(s390x): add repack support for q4_0 (#28667)
ggml-cpu: clean comments

Signed-off-by: Aaron Teo <[email protected]>
b10887
2026-09-10 09:51:17 +03:00
Aaron Teo 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 <[email protected]>

* ggml-cpu: clean up variable naming for understanding

Signed-off-by: Aaron Teo <[email protected]>

* docs: update support for Q1_0

Signed-off-by: Aaron Teo <[email protected]>

---------

Signed-off-by: Aaron Teo <[email protected]>
b10886
2026-09-10 09:50:28 +03:00
Aaron Teo d7e86430a7 model: fix all granite family parameter counts (#28643)
* model: fix all granite family parameter counts

Signed-off-by: Aaron Teo <[email protected]>

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

Signed-off-by: Aaron Teo <[email protected]>

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

Signed-off-by: Aaron Teo <[email protected]>

---------

Signed-off-by: Aaron Teo <[email protected]>
b10885
2026-09-10 09:49:56 +03:00
Eve 434ddbbc0e ci: fix sanitizer tests (#28583) b10884 2026-09-09 19:46:28 +00:00
Ruben Ortlam 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
Todor BoinovskiandMax Krasnyansky 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 <[email protected]>
2026-09-09 08:40:24 -07:00
Masato Nakasaka 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æret 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æret 9cf3bf256b py : bump numpy to 2.4.6 (#28649) 2026-09-09 14:59:37 +02:00
linsen458-specandlinsen 4850c7727f llama : use int32_t for llama_sampler_chain_n return type (#28631)
Contributes to #4574

Co-authored-by: linsen <[email protected]>
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 <[email protected]>

* 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 <[email protected]>
Co-authored-by: Carl Philipp Klemm <[email protected]>
b10877
2026-09-09 13:25:54 +02:00
Piotr Wilkin (ilintar)andJohannes Gäßler 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 <[email protected]>

* apply code review comments

---------

Co-authored-by: Johannes Gäßler <[email protected]>
b10876
2026-09-09 12:50:08 +02:00
Aaron Teo 14a9d09f75 args: officially deprecate --mmap|mlock|dio (#28334)
Signed-off-by: Aaron Teo <[email protected]>
b10875
2026-09-09 18:36:27 +08:00
Aaron Teo e2d2c0d6aa model: fix granite3 moe unknown parameter count (#28632)
Signed-off-by: Aaron Teo <[email protected]>
b10874
2026-09-09 16:46:22 +08:00
Xuan-Son Nguyen 6de9cdb26b mtmd: propagate video ID to bitmap (#28601) b10873 2026-09-09 10:37:36 +02:00
Pascal 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
WakeUpMorty 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, Fish 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 Gerganov 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
cwriterandcwriter 304665fe7a Add IQ type handling for MoE (#28476)
Co-authored-by: cwriter <cwriter@localhost>
b10868
2026-09-09 09:25:41 +08: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
b10867
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.
b10865
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.
b10864
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
b10863
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 <[email protected]>

* fix compile-error: add missing header

Bug: #28559
Signed-off-by: Pepper Gray <[email protected]>

* fix compile-error: add missing header

Bug: #28560
Signed-off-by: Pepper Gray <[email protected]>

* fix compile-error: add missing header

Bug: #28561
Signed-off-by: Pepper Gray <[email protected]>

* fix compile-error: add missing header

Bug: #28562
Signed-off-by: Pepper Gray <[email protected]>

* fix compile-error: add missing header

Bug: #28564
Signed-off-by: Pepper Gray <[email protected]>

---------

Signed-off-by: Pepper Gray <[email protected]>
b10859
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
b10858
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.
b10857
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
b10856
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
b10855
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) b10853 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
b10852
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.
b10850
2026-09-07 19:54:13 +02:00