]> git.djapps.eu Git - pkg/ggml/sources/llama.cpp/log
pkg/ggml/sources/llama.cpp
6 weeks agoggml webgpu: faster normal quant and some k-quant matrix operations, better shader...
Reese Levine [Tue, 10 Mar 2026 16:14:27 +0000 (09:14 -0700)]
ggml webgpu: faster normal quant and some k-quant matrix operations, better shader parameter handling (#20173)

* K quant speedup (#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <redacted>
Co-authored-by: Reese Levine <redacted>
* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <redacted>
* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <redacted>
6 weeks agoReduce level of content parser warning message to avoid log spam on non-debug verbosi...
Piotr Wilkin (ilintar) [Tue, 10 Mar 2026 14:21:51 +0000 (15:21 +0100)]
Reduce level of content parser warning message to avoid log spam on non-debug verbosity (#20347)

6 weeks agoexamples : fix empty items in json_schema_to_grammar.py [no ci] (#19968)
Ray Xu [Tue, 10 Mar 2026 13:38:18 +0000 (21:38 +0800)]
examples : fix empty items in json_schema_to_grammar.py [no ci] (#19968)

* Fix logic for retrieving schema items in `json_schema_to_grammar.py`

If `schema['items']` is `{}` and `prefixItems not in schema', as `{}` is Falsy, the original code here will raise an error.

I think if `schema['items']` is `{}`, them items should just be `{}`

* Apply suggestion from @CISC

Co-authored-by: Sigbjørn Skjæret <redacted>
* Add tests for arrays with empty items

Add two unit tests to `tests/test-json-schema-to-grammar.cpp` that validate handling of arrays when 'items' is an empty schema and when 'prefixItems' is present alongside an empty 'items'. Both tests expect the same generated grammar, ensuring the JSON Schema->grammar conversion treats an empty 'items' schema (and the presence of 'prefixItems') correctly and covering this edge case.

---------

Co-authored-by: Sigbjørn Skjæret <redacted>
6 weeks agodocs: update CPU backend ops to mark POOL_1D as supported (#20304)
a3894281 [Tue, 10 Mar 2026 13:31:24 +0000 (15:31 +0200)]
docs: update CPU backend ops to mark POOL_1D as supported (#20304)

6 weeks agomodels : fix assert in mamba2 (cont) (#20335)
Georgi Gerganov [Tue, 10 Mar 2026 13:00:08 +0000 (15:00 +0200)]
models : fix assert in mamba2 (cont) (#20335)

* models : fix assert in mamba2 (cont)

* cont : add n_group mod

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Co-authored-by: Sigbjørn Skjæret <redacted>
6 weeks agoserver : make 2 checkpoints near the end of the prompt (#20288)
Georgi Gerganov [Tue, 10 Mar 2026 12:28:23 +0000 (14:28 +0200)]
server : make 2 checkpoints near the end of the prompt (#20288)

* server : make 2 checkpoints near the end of the prompt

* cont : adjust checkpoints

6 weeks agocommon : fix incorrect uses of stoul (#20313)
Sigbjørn Skjæret [Tue, 10 Mar 2026 10:40:26 +0000 (11:40 +0100)]
common : fix incorrect uses of stoul (#20313)

6 weeks agokleidiai : support for concurrent sme and neon kernel execution (#20070)
Charles Xu [Tue, 10 Mar 2026 07:25:25 +0000 (08:25 +0100)]
kleidiai : support for concurrent sme and neon kernel execution (#20070)

6 weeks agoggml-cpu: add RVV repack GEMM and GEMV for quantization types (#19121)
Taimur Ahmad [Tue, 10 Mar 2026 06:49:52 +0000 (11:49 +0500)]
ggml-cpu: add RVV repack GEMM and GEMV for quantization types (#19121)

* ggml-cpu: add rvv ggml_quantize_mat_4x8 for q8_0

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv repacking for iq4_nl

* ggml-cpu: add generic impl for iq4_nl gemm/gemv

* ggml-cpu: add rvv repacking for q8_0

* ggml-cpu: refactor; add rvv repacking for q4_0, q4_K

* ggml-cpu: refactor; add rvv repacking for q2_K

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: refactor rvv repack

---------

Co-authored-by: Rehan Qasim <redacted>
6 weeks agometal: handle command buffer failures gracefully in synchronize (#20306)
Julian Pscheid [Tue, 10 Mar 2026 06:32:24 +0000 (23:32 -0700)]
metal: handle command buffer failures gracefully in synchronize (#20306)

Replace GGML_ABORT("fatal error") in ggml_metal_synchronize() with
error flag + return. This aligns synchronize error handling with
graph_compute, which already returns GGML_STATUS_FAILED for the same
condition.

When a command buffer fails (e.g., iOS GPU access revocation during
backgrounding, macOS eGPU disconnect, OOM), the backend enters an
error state instead of killing the host process. Subsequent
graph_compute calls return GGML_STATUS_FAILED immediately. Recovery
requires recreating the backend.

Failed extra command buffers are properly released on the error path
to avoid Metal object leaks.

6 weeks agollama-quant : fail early on missing imatrix, refactor type selection, code cleanup...
ddh0 [Tue, 10 Mar 2026 06:16:05 +0000 (01:16 -0500)]
llama-quant : fail early on missing imatrix, refactor type selection, code cleanup (#19770)

* quantize : imatrix-fail early + code cleanup

* fix manual override printing

it's in the preliminary loop now, so needs to be on its own line

* revert header changes per ggerganov

* remove old #includes

* clarify naming

rename `tensor_quantization` to `tensor_typo_option` to descirbe its
functionality

* fix per barto

6 weeks agocommon: consolidate PEG string parsers (#20263)
Aldehir Rojas [Mon, 9 Mar 2026 23:29:21 +0000 (18:29 -0500)]
common: consolidate PEG string parsers (#20263)

* common : consolidate PEG string parsers
* cont : fix json_string_content()

6 weeks agomodel: fix step3.5 n_rot (#20318)
Xuan-Son Nguyen [Mon, 9 Mar 2026 22:42:24 +0000 (23:42 +0100)]
model: fix step3.5 n_rot (#20318)

6 weeks agollama: dynamic head_dim and n_rot for SWA (#20301)
Xuan-Son Nguyen [Mon, 9 Mar 2026 21:22:39 +0000 (22:22 +0100)]
llama: dynamic head_dim and n_rot for SWA (#20301)

* llama: dynamic head_dim and n_rot for SWA

* also add gguf_writer wrappers

* fix build

* build_rope_shift arg reorder

6 weeks agoserver: Parse port numbers from MCP server URLs in CORS proxy (#20208)
Evan Huus [Mon, 9 Mar 2026 16:47:54 +0000 (12:47 -0400)]
server: Parse port numbers from MCP server URLs in CORS proxy (#20208)

* Parse port numbers from MCP server URLs

* Pass scheme to http proxy for determining whether to use SSL

* Fix download on non-standard port and re-add port to logging

* add test

---------

Co-authored-by: Xuan Son Nguyen <redacted>
6 weeks agometal : extend mul_mv_ext to BF16, Q2_K, Q3_K (#20250)
Paul Flynn [Mon, 9 Mar 2026 14:48:12 +0000 (10:48 -0400)]
metal : extend mul_mv_ext to BF16, Q2_K, Q3_K (#20250)

Enable mul_mv_ext small-batch kernels (BS 2-8) for BF16, Q2_K,
and Q3_K quantization types. These types previously fell through
to the slower single-row mul_mv path.

BF16 uses the float4 dequantize path (like F16). Q2_K and Q3_K
use the float4x4 K-quant path (like Q4_K/Q5_K/Q6_K).

Co-authored-by: Claude Opus 4.6 <redacted>
6 weeks agoserver : fix checkpoints n_tokens calculation (#20287)
Georgi Gerganov [Mon, 9 Mar 2026 14:47:06 +0000 (16:47 +0200)]
server : fix checkpoints n_tokens calculation (#20287)

6 weeks agometal : add upscale (#20284)
Georgi Gerganov [Mon, 9 Mar 2026 14:45:11 +0000 (16:45 +0200)]
metal : add upscale (#20284)

6 weeks agoserver : warn swa-full is not supported for non-SWA models (#20291)
Georgi Gerganov [Mon, 9 Mar 2026 14:44:25 +0000 (16:44 +0200)]
server : warn swa-full is not supported for non-SWA models (#20291)

6 weeks agoserver : fix off-by-1 in server_tokens::size_up_to_pos() (#20279)
Georgi Gerganov [Mon, 9 Mar 2026 14:43:38 +0000 (16:43 +0200)]
server : fix off-by-1 in server_tokens::size_up_to_pos() (#20279)

* server : fix off-by-1 in server_tokens::size_up_to_pos()

* cont : fix typo [no ci]

6 weeks agocommon: map developer role to system (#20215)
Piotr Wilkin (ilintar) [Mon, 9 Mar 2026 13:25:11 +0000 (14:25 +0100)]
common: map developer role to system (#20215)

* Map developer role to system
* Simplify

6 weeks agomodels : fix assert in mamba2 graph (#20270)
Georgi Gerganov [Mon, 9 Mar 2026 11:15:15 +0000 (13:15 +0200)]
models : fix assert in mamba2 graph (#20270)

6 weeks agoserver : add kill switch when server is stuck (#20277)
Georgi Gerganov [Mon, 9 Mar 2026 08:33:12 +0000 (10:33 +0200)]
server : add kill switch when server is stuck (#20277)

6 weeks agoggml-cuda: disable gdn for musa (#20278)
Aman Gupta [Mon, 9 Mar 2026 08:15:36 +0000 (16:15 +0800)]
ggml-cuda: disable gdn for musa (#20278)

6 weeks agollama-quant : left-align tensor names in output (#20117)
ddh0 [Mon, 9 Mar 2026 07:28:41 +0000 (02:28 -0500)]
llama-quant : left-align tensor names in output (#20117)

6 weeks agocontributing: limit open PRs for new contributors to 1 (#20036)
Aman Gupta [Mon, 9 Mar 2026 07:05:34 +0000 (15:05 +0800)]
contributing: limit open PRs for new contributors to 1 (#20036)

6 weeks agoggml-vulkan: add SGN operator, auto-generate Vulkan.csv and ops.md (#20219)
Bertay Eren [Mon, 9 Mar 2026 06:24:16 +0000 (09:24 +0300)]
ggml-vulkan: add SGN operator, auto-generate Vulkan.csv and ops.md (#20219)

6 weeks agovulkan: skip zero size tensors in backend copies (#20233)
Ruben Ortlam [Mon, 9 Mar 2026 06:23:45 +0000 (07:23 +0100)]
vulkan: skip zero size tensors in backend copies (#20233)

6 weeks agocuda : display total and free VRAM capacity during device initialization (#20185)
Michael Huang [Mon, 9 Mar 2026 04:45:43 +0000 (21:45 -0700)]
cuda : display total and free VRAM capacity during device initialization (#20185)

6 weeks agollama-bench: introduce `-hf` and `-hff` flags & use `--mmap 1` by default (#20211)
Aaron Teo [Mon, 9 Mar 2026 01:05:44 +0000 (09:05 +0800)]
llama-bench: introduce `-hf` and `-hff` flags & use `--mmap 1` by default (#20211)

6 weeks agoPEG parser for LFM2 (#20251)
Piotr Wilkin (ilintar) [Mon, 9 Mar 2026 00:11:22 +0000 (01:11 +0100)]
PEG parser for LFM2 (#20251)

* PEG parser for LFM2

* Simplify using python_value()

7 weeks agoserver : do not create checkpoints right after mtmd chunks (#20232)
Georgi Gerganov [Sun, 8 Mar 2026 20:16:46 +0000 (22:16 +0200)]
server : do not create checkpoints right after mtmd chunks (#20232)

7 weeks agograph : remove redundant scale_w parameter (#20235)
Sigbjørn Skjæret [Sun, 8 Mar 2026 17:58:28 +0000 (18:58 +0100)]
graph : remove redundant scale_w parameter (#20235)

7 weeks agocommon : gracefully handle incomplete output (#20191)
Aldehir Rojas [Sun, 8 Mar 2026 16:17:02 +0000 (11:17 -0500)]
common : gracefully handle incomplete output (#20191)

* common : handle incomplete UTF-8 at end of input in PEG parser

* cont : if reached end prematurely, emit needs_more_input to propagate partial output

* cont: refactor peg parse context to add lenient flag

* cont : remove partial flag, keep lenient flag

7 weeks agoFix compile bug (#20203)
Piotr Wilkin (ilintar) [Sun, 8 Mar 2026 16:15:49 +0000 (17:15 +0100)]
Fix compile bug (#20203)

* Fix compile bug

* Update common/chat-auto-parser-helpers.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Co-authored-by: Sigbjørn Skjæret <redacted>
7 weeks agoFix structured outputs (#20223)
Piotr Wilkin (ilintar) [Sun, 8 Mar 2026 16:14:43 +0000 (17:14 +0100)]
Fix structured outputs (#20223)

* Fix structured outputs

* Update common/chat-auto-parser-generator.cpp

Co-authored-by: Aldehir Rojas <redacted>
---------

Co-authored-by: Aldehir Rojas <redacted>
7 weeks agoggml-vulkan: Add ELU op support (#20183)
GiantPrince [Sun, 8 Mar 2026 11:38:17 +0000 (07:38 -0400)]
ggml-vulkan: Add ELU op support (#20183)

* ggml-Vulkan: add ELU support

* ggml-Vulkan: remove extra spaces and variables

* ggml-Vulkan: fix format issue

* ggml-Vulkan: fix format issue

* fix whitespace issue

* Update Vulkan.csv and ops.md

7 weeks agovulkan: Fix data races in coopmat1 mul_mat(_id) (#20084)
Jeff Bolz [Sun, 8 Mar 2026 11:33:48 +0000 (06:33 -0500)]
vulkan: Fix data races in coopmat1 mul_mat(_id) (#20084)

* vulkan: Fix data races in coopmat1 mul_mat(_id)

Add barriers between coopmat store and regular loads. We sort of got away with
this because it was the same subgroup accessing the values, but it's still a
race and may not work.

* switch to subgroup control barriers

7 weeks agollama: end-to-end tests (#19802)
Johannes Gäßler [Sun, 8 Mar 2026 11:30:21 +0000 (12:30 +0100)]
llama: end-to-end tests (#19802)

* tests: add end-to-end tests per model architecture

* fixup for rebase

* fix use-after-free in llama-model-loader.cpp

* fix CI

* fix WebGPU

* fix CI

* disable CI for macOS-latest-cmake-arm64

* use expert_weights_scale only if != 0.0f

* comments

7 weeks agoreadme : update infra list (#20212)
Christopher Maher [Sun, 8 Mar 2026 10:42:28 +0000 (03:42 -0700)]
readme : update infra list (#20212)

7 weeks agoRevert to OAI-compatible args (#20213)
Piotr Wilkin (ilintar) [Sun, 8 Mar 2026 10:33:03 +0000 (11:33 +0100)]
Revert to OAI-compatible args (#20213)

* Revert to OAI-compatible args

* Apply workaround::func_args_not_string

7 weeks agoserver : correct index on finish in OAI completion streams (#20226)
decahedron1 [Sun, 8 Mar 2026 09:08:57 +0000 (04:08 -0500)]
server : correct index on finish in OAI completion streams (#20226)

7 weeks ago[SYCL] supprt Flash Attention for fp32/fp16/Q4/Q5/Q8 (#20190)
Neo Zhang [Sun, 8 Mar 2026 04:00:07 +0000 (12:00 +0800)]
[SYCL] supprt Flash Attention for fp32/fp16/Q4/Q5/Q8 (#20190)

* support flash-attention for fp32/fp16/Q4/Q5/Q8

* rm warining

* update for JIT

7 weeks agoggml: add GATED_DELTA_NET op (#19504)
Aman Gupta [Sat, 7 Mar 2026 07:41:10 +0000 (15:41 +0800)]
ggml: add GATED_DELTA_NET op (#19504)

* ggml: add GATED_DELTA_NET op

* remove the transpose

* add KDA

* add qwen35 dense

* llama : check for fused gated delta net backend support

---------

Co-authored-by: Georgi Gerganov <redacted>
7 weeks agoopencl: add l2_norm (#20160)
lhez [Sat, 7 Mar 2026 02:03:05 +0000 (18:03 -0800)]
opencl: add l2_norm (#20160)

7 weeks agoAutoparser: True streaming (#20177)
Piotr Wilkin (ilintar) [Sat, 7 Mar 2026 00:55:33 +0000 (01:55 +0100)]
Autoparser: True streaming (#20177)

* Relax atomicity constraint for nicer, more pleasent, True Streaming parsing

* Whitespace

* Remove redundant atomics

7 weeks agoAutoparser: add optional argument reshuffle capability (#20171)
Piotr Wilkin (ilintar) [Fri, 6 Mar 2026 21:34:15 +0000 (22:34 +0100)]
Autoparser: add optional argument reshuffle capability (#20171)

* Allow reshuffled arguments in tagged argument parser format tool calls.

* Remove shuffle just keep the optional parsers in any order

* Remove unnecessary import

7 weeks agoquants : Add memsets and other fixes for IQ quants (#19861)
Bartowski [Fri, 6 Mar 2026 21:06:56 +0000 (16:06 -0500)]
quants : Add memsets and other fixes for IQ quants (#19861)

* Add memsets and other fixes for IQ quants

* Make memset unconditional, change Laux back to L

* Move another memset

7 weeks agoAdd @pwilkin to CODEOWNERS for autoparser code (#20174)
Piotr Wilkin (ilintar) [Fri, 6 Mar 2026 20:25:41 +0000 (21:25 +0100)]
Add @pwilkin to CODEOWNERS for autoparser code (#20174)

7 weeks agoAutoparser - complete refactoring of parser architecture (#18675)
Piotr Wilkin (ilintar) [Fri, 6 Mar 2026 20:01:00 +0000 (21:01 +0100)]
Autoparser - complete refactoring of parser architecture (#18675)

* Autoparser - full single commit squish

* Final pre-merge changes: minor fixes, Kimi 2.5 model parser

7 weeks agohexagon: add f32 ssm_conv op (#20122)
Todor Boinovski [Fri, 6 Mar 2026 17:59:26 +0000 (09:59 -0800)]
hexagon: add f32 ssm_conv op (#20122)

* hexagon: add ssm_conv op

* hexagon: hvx kernel is functional

* hexagon: improvements to ssm-conv hvx kernel

* hexagon: added dma to ssm-conv hvx kernel

* hexagon: ssm-conv dynamically compute gather scratchpad

* hex-ssm-conv: add local context and fix various issues (spad indexing, etc)

---------

Co-authored-by: Max Krasnyansky <redacted>
7 weeks agoserver : preserve anthropic thinking blocks in conversion (#20120)
Tom Vaucourt [Fri, 6 Mar 2026 16:41:12 +0000 (17:41 +0100)]
server : preserve anthropic thinking blocks in conversion (#20120)

* server : preserve anthropic thinking blocks in conversion (#20090)

* server : add tests for anthropic thinking block conversion

---------

Co-authored-by: root <redacted>
7 weeks agocpu: skip redudant ROPE cache updates (#20149)
Max Krasnyansky [Fri, 6 Mar 2026 16:32:40 +0000 (08:32 -0800)]
cpu: skip redudant ROPE cache updates (#20149)

7 weeks agoggml-cuda: add mem check for fusion (#19916)
Aman Gupta [Fri, 6 Mar 2026 16:05:43 +0000 (00:05 +0800)]
ggml-cuda: add mem check for fusion (#19916)

* ggml-cuda: add mem check for fusion

* Replace NaNs with -FLT_MAX

* fix typo

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
7 weeks agoggml: update comments for backends which have no memory to report (#20157)
Aaron Teo [Fri, 6 Mar 2026 15:24:38 +0000 (23:24 +0800)]
ggml: update comments for backends which have no memory to report (#20157)

Signed-off-by: Aaron Teo <redacted>
7 weeks agoggml-cpu: Fix gcc 15 ICE on ppc64le (#20083) (#20130)
shalinib-ibm [Fri, 6 Mar 2026 15:22:39 +0000 (20:52 +0530)]
ggml-cpu: Fix gcc 15 ICE on ppc64le (#20083) (#20130)

This patch addresses an Internal Compiler Error (Segmentation fault)
observed with gcc 15 by replacing the intrinsic + cast by doing
a cat on the data first and then calling the intrinsic. This bypasses the
buggy compiler path while maintaining identical instruction selection.

Performance Verification:
Assembly analysis on RHEL 9 (GCC 15.1.1) confirms that both the original
code and this fix generate the identical Power10 prefixed load instruction:
    `plxv 40, 2(14)`

This ensures zero performance regression while unblocking builds on
newer toolchains.

Reproduced on:
- Alpine Linux + GCC 15.2.0-r2
- RHEL 9  + GCC 15.1.1 (gcc-toolset-15)

Signed-off-by: Shalini Salomi Bodapati <redacted>
7 weeks agoCUDA: use shared mem for ssm_conv (#20128)
Aman Gupta [Fri, 6 Mar 2026 15:09:59 +0000 (23:09 +0800)]
CUDA: use shared mem for ssm_conv (#20128)

* CUDA: use shared mem for ssm_conv

* fuse silu + ssm_conv

* fuse unary + mul

* enable for fp16

* formatting

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
7 weeks agocontext: ignore zero scale LoRAs when checking sameness (#20166)
Tim Neumann [Fri, 6 Mar 2026 13:05:52 +0000 (14:05 +0100)]
context: ignore zero scale LoRAs when checking sameness (#20166)

7 weeks agoCheckpoint every n tokens: squash (#20087)
Piotr Wilkin (ilintar) [Fri, 6 Mar 2026 10:39:26 +0000 (11:39 +0100)]
Checkpoint every n tokens: squash (#20087)

7 weeks agowebui: Agentic Loop + MCP Client with support for Tools, Resources and Prompts (...
Aleksander Grygier [Fri, 6 Mar 2026 09:00:39 +0000 (10:00 +0100)]
webui: Agentic Loop + MCP Client with support for Tools, Resources and Prompts (#18655)

7 weeks agoggml-cpu: fix data race for debug asserts (#20148)
Johannes Gäßler [Fri, 6 Mar 2026 08:12:49 +0000 (09:12 +0100)]
ggml-cpu: fix data race for debug asserts (#20148)

7 weeks agokv-cache : fix M-RoPE checkpoints (#20132)
Georgi Gerganov [Fri, 6 Mar 2026 06:46:51 +0000 (08:46 +0200)]
kv-cache : fix M-RoPE checkpoints (#20132)

7 weeks agocli : Don't clear system prompt when using '/clear' (#20067)
Roj234 [Fri, 6 Mar 2026 05:41:11 +0000 (13:41 +0800)]
cli : Don't clear system prompt when using '/clear' (#20067)

* Enhance /clear command to include system prompt

Add system prompt to messages when clearing chat history.

* Use lambda

7 weeks agoopencl: add neg, exp and diag (#20127)
lhez [Fri, 6 Mar 2026 05:16:39 +0000 (21:16 -0800)]
opencl: add neg, exp and diag (#20127)

* opencl: add `neg`

* opencl: add `exp`

* opencl: add `diag`

7 weeks agohexagon: add fp16 support for binary ops: add,sub,mul,div (#20139)
YardenTal44 [Fri, 6 Mar 2026 02:29:13 +0000 (04:29 +0200)]
hexagon: add fp16 support for binary ops: add,sub,mul,div (#20139)

* hexagon: add fp16 support for binary ops: add,sub,mul,div

* hexagon: fix test-backend-ops failures for fp16 binary ops on older arches (<v79)

* hexagon: decide on n_threads (aka n_jobs) early to avoid overallocating scratchpad

* snapdragon: fix readme link

---------

Co-authored-by: Max Krasnyansky <redacted>
7 weeks agomodels : kda chunk size = 16 (#19827)
ymcki [Thu, 5 Mar 2026 15:01:23 +0000 (23:01 +0800)]
models : kda chunk size = 16 (#19827)

* models : add llm_build_delta_net_base

* cont : keep qwen35 and qwen35moe graphs intact

* cont : add comments [no ci]

* add kimi linear to delta-net-base

* removed unnecessary ggml_cont from g_exp_t

* removed ggml_cont from g_diff_exp_t. moved ggml_cont for o to kimi-linear.cpp

* removed unnecessary diag mask

* cont : simplify

* cont : avoid graph splits

* scale q after mul instead of beginning

* scale q after mul instead of beginning

* identical ppl

* cont : fix scale and decay mask

* minor : remove TODO

* block implementation for kda

* remove space at the end of line 101

* concat+pad

* pad+binary row concat

* chunk size 16 for kda

* removed minor differences to master

---------

Co-authored-by: Georgi Gerganov <redacted>
7 weeks agoCUDA: Improve performance via less synchronizations between token (#17795)
Andreas Kieslinger [Thu, 5 Mar 2026 11:53:21 +0000 (12:53 +0100)]
CUDA:  Improve performance via less synchronizations between token (#17795)

* Adds CPU-to-CUDA copy capability to
ggml_backend_cuda_cpy_tensor_async()

* Adds function to relax sync requirements between input copies on
supported backends (CUDA for now)

* Exchanges synchronous copy with async copy function.

* Adds macro guards to allow compilation in non-CUDA builds

* Reworked backend detection in ggml-backend.cpp to avoid linking
conflicts

* Relax requirement of checks in async CUDA copies from backend and buffer type to just buffer type, to avoid linking issues

* Minor cleanup

* Makes opt-in to relax use of explicit syncs more general. Backends like
vulkan which require a synchronization between HtoD copies and graph
execution could also adopt this change now.

* Reintroduces stricter check for CPU->CUDA backend async copy via
GGML_DEVICE_TYPE_CPU.

* Corrects initialization of ggml_backend_sync_mode in
ggml_backend_sched_split initialization

* Simplifies synchronizations to adhere to `saaasg` pattern.

* Apply suggestion from @ggerganov (src->buffer to buf_src)

Co-authored-by: Georgi Gerganov <redacted>
* Apply suggestion from @ggerganov (src->buffer to buf_src) v2

Co-authored-by: Georgi Gerganov <redacted>
---------

Co-authored-by: Georgi Gerganov <redacted>
7 weeks agomodel : update Qwen3.5 model type detection (#20126)
Eric Zhang [Thu, 5 Mar 2026 11:47:14 +0000 (19:47 +0800)]
model : update Qwen3.5 model type detection (#20126)

* model : fix Qwen3.5 model type detection

* Update src/llama-model.cpp

whoops, my bad

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Co-authored-by: Sigbjørn Skjæret <redacted>
7 weeks agocli : add command and file auto-completion (#19985)
Sigbjørn Skjæret [Thu, 5 Mar 2026 09:47:28 +0000 (10:47 +0100)]
cli : add command and file auto-completion (#19985)

7 weeks agoconvert : register Qwen 3.5 ForCausalLM for text only (#20119)
Sigbjørn Skjæret [Thu, 5 Mar 2026 09:30:02 +0000 (10:30 +0100)]
convert : register Qwen 3.5 ForCausalLM for text only (#20119)

7 weeks agowebui: Improvements for Models Selector UI (#20066)
Aleksander Grygier [Thu, 5 Mar 2026 07:52:22 +0000 (08:52 +0100)]
webui: Improvements for Models Selector UI (#20066)

7 weeks agochore : correct typos [no ci] (#20041)
Marcel Petrick [Thu, 5 Mar 2026 07:50:21 +0000 (08:50 +0100)]
chore : correct typos [no ci] (#20041)

* fix(docs): correct typos found during code review

Non-functional changes only:
- Fixed minor spelling mistakes in comments
- Corrected typos in user-facing strings
- No variables, logic, or functional code was modified.

Signed-off-by: Marcel Petrick <redacted>
* Update docs/backend/CANN.md

Co-authored-by: Aaron Teo <redacted>
* Revert "Auxiliary commit to revert individual files from 846d1c301281178efbc6ce6060ad34c1ebe45af8"

This reverts commit 02fcf0c7db661d5ff3eff96b2b2db9fdb7213256.

* Update tests/test-backend-ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Signed-off-by: Marcel Petrick <redacted>
Co-authored-by: Aaron Teo <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
7 weeks agohexagon: Flash Attention optimizations (dma, mpyacc, multi-row) and MatMul updates...
Max Krasnyansky [Thu, 5 Mar 2026 05:55:29 +0000 (21:55 -0800)]
hexagon: Flash Attention optimizations (dma, mpyacc, multi-row) and MatMul updates (#20118)

* ggml-hexagon: enhance hvx_dot_f16_f16_aa_rx4 for improved performance by expanding vector handling and optimizing accumulation

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: optimize hvx_dot_f16_f16_aa_rx4 and enhance hvx_vec_reduce_sum_f32x4 for improved performance and reduced complexity

* ggml-hexagon: add hvx_dot_f16_f16_aa_rx32 for enhanced vector processing in flash attention

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* optimize hvx_dot_f16_f16_aa_rx4 and hvx_dot_f16_f16_aa_rx32 by removing unused scale parameter and improving vector accumulation

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: refactor hvx_dot_f16_f16_aa_rx4 for improved readability and return HVX_Vector for better integration

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: initialize sums variable in hvx_dot_f16_f16_aa_rx32 for clarity

* ggml-hexagon: fix compiling error

* fix hvx_dot_f16_f16_aa_rx4 to handle leftover elements correctly using masking

* refactor hvx_dot_f16_f16_aa_rx4 to accept vector and leftover element counts as parameters for improved clarity and flexibility

* wip

* fa: instrumentation and dma reordering

* hex-fa: use block-size 64 to improve DMA pipelining

* hex-fa: optimize vec-dot for v79 and above

* hex-fa: use block size 64

* hex-fa: avoid scalar fp32->fp16 conversions

* hex-fa: simplify dot_f16 functions using optimized vec_mpyacc

* hex-fa: rewrite mad_f32_f16 using hvx_vec_mpyacc

* hex-mm: use mpyacc in matmul dot functions

---------

Co-authored-by: chraac <redacted>
7 weeks agoopencl: add `SET`, support i32 for `CPY`, minor refactor for cpy (#20101)
lhez [Thu, 5 Mar 2026 05:32:26 +0000 (21:32 -0800)]
opencl: add `SET`, support i32 for `CPY`, minor refactor for cpy (#20101)

7 weeks agohexagon: add llama-completion runner script (#20095)
Todor Boinovski [Wed, 4 Mar 2026 23:04:59 +0000 (15:04 -0800)]
hexagon: add llama-completion runner script (#20095)

7 weeks ago[WebGPU] Fix wait logic for inflight jobs (#20096)
Nikhil Jain [Wed, 4 Mar 2026 19:54:55 +0000 (11:54 -0800)]
[WebGPU] Fix wait logic for inflight jobs (#20096)

* Enable tmate debugging for investigating thread safety issue

* Refactor wait and submit to operate on vector<wgpu::FutureWaitInfo>, and fix wait to delete only the future that is completed.

* Cleanup

* Remove clear change and run clang-format

* Cleanup

7 weeks agoAdd concat op to webgpu. (#20068)
Masashi Yoshimura [Wed, 4 Mar 2026 19:19:00 +0000 (04:19 +0900)]
Add concat op to webgpu. (#20068)

7 weeks agotools : add missing clocale include in mtmd-cli [no ci] (#20107)
Sigbjørn Skjæret [Wed, 4 Mar 2026 13:18:04 +0000 (14:18 +0100)]
tools : add missing clocale include in mtmd-cli [no ci] (#20107)

7 weeks agoggml: fix ggml_is_contiguous_n for ne == 1 (#20092)
Johannes Gäßler [Wed, 4 Mar 2026 11:04:31 +0000 (12:04 +0100)]
ggml: fix ggml_is_contiguous_n for ne == 1 (#20092)

7 weeks agoggml : use a simple std::thread in AMX without OpenMP (#20074)
Adrien Gallouët [Wed, 4 Mar 2026 10:57:09 +0000 (11:57 +0100)]
ggml : use a simple std::thread in AMX without OpenMP (#20074)

Disabling OpenMP generally provides better inference performance (at
least in my testing) but the loading becomes slightly slower.

Benchmark results for `convert_B_packed_format()`:

Before this commit:

         N      K |  No OpenMP     OpenMP |    Diff |  Speedup
    ------------------------------------------------------------
       512   2880 |    640.9us    263.5us |  -58.9% |    0.41x
      2880   4096 |     2.55ms    261.7us |  -89.8% |    0.10x
    201088   2880 |   256.44ms    21.61ms |  -91.6% |    0.08x
    ------------------------------------------------------------

    Total: 325.43ms vs 31.05ms

After:

         N      K |  No OpenMP     OpenMP |    Diff |  Speedup
    ------------------------------------------------------------
       512   2880 |     1.49ms    263.5us |  -82.3% |    0.18x
      2880   4096 |     1.55ms    261.7us |  -83.1% |    0.17x
    201088   2880 |    24.03ms    21.61ms |  -10.1% |    0.90x
    ------------------------------------------------------------

    Total: 78.97ms vs 31.05ms

Tested with unsloth/gpt-oss-20b-GGUF:Q4_K_M.

Signed-off-by: Adrien Gallouët <redacted>
7 weeks agoimpl : use 6 digits for tensor dims (#20094)
ddh0 [Wed, 4 Mar 2026 08:53:38 +0000 (02:53 -0600)]
impl : use 6 digits for tensor dims (#20094)

Many models have vocabulary sizes, and thus tensor shapes, with more
than 5 digits (ex: Gemma 3's vocab size is 262,208).

I already fixed this for `llama_format_tensor_shape` but missed it for
`llama_format_tensor_shape` until now. Oops.

7 weeks agoFix locale-dependent float printing in GGUF metadata (#17331)
SamareshSingh [Wed, 4 Mar 2026 08:30:40 +0000 (02:30 -0600)]
Fix locale-dependent float printing in GGUF metadata (#17331)

* Set C locale for consistent float formatting across all binaries.

* Add C locale setting to all tools binaries

Add std::setlocale(LC_NUMERIC, "C") to all 16 binaries in the tools/
directory to ensure consistent floating-point formatting.

* Apply suggestion from @JohannesGaessler

---------

Co-authored-by: Johannes Gäßler <redacted>
7 weeks agocompletion : Fix a typo in warning message (#20082)
standby24x7 [Wed, 4 Mar 2026 05:44:49 +0000 (14:44 +0900)]
completion : Fix a typo in warning message (#20082)

resuse -> reuse

7 weeks agodocs: Fix intel documentation link (#20040)
Mickael Desgranges [Tue, 3 Mar 2026 13:50:00 +0000 (14:50 +0100)]
docs: Fix intel documentation link (#20040)

7 weeks agokleidiai : add sme fp16 compute path for q4_0 gemm on aarch64 (#20043)
Charles Xu [Tue, 3 Mar 2026 09:40:26 +0000 (10:40 +0100)]
kleidiai : add sme fp16 compute path for q4_0 gemm on aarch64 (#20043)

7 weeks agoopencl: add optimized q4_1 mm kernel for adreno (#19840)
shaofeiqi [Tue, 3 Mar 2026 03:49:41 +0000 (19:49 -0800)]
opencl: add optimized q4_1 mm kernel for adreno (#19840)

* Add Q4_1 OpenCL Kernels

* opencl: refactor transpose

* opencl: format

* opencl: refactor q4_1 unpack

* opencl: move `ggml_cl_mul_mat_q4_1_f32_adreno`

* opencl: refactor `ggml_cl_mul_mat_q4_1_f32_adreno` and kernels

* opencl: rename kernel files and kernes

* opencl: fix build for non adreno

* opencl: move code around and format

---------

Co-authored-by: Li He <redacted>
7 weeks agoggml webgpu: fix workgroup dispatch limit for large batch sizes (#19965)
Abhijit Ramesh [Tue, 3 Mar 2026 03:35:11 +0000 (19:35 -0800)]
ggml webgpu: fix workgroup dispatch limit for large batch sizes (#19965)

* ggml-webgpu: fix workgroup dispatch limit for large batch sizes

WebGPU limits workgroup sizes to 65535 per dimension. Large MUL_MAT
operations with batch sizes exceedeing this limi would fail.

* add compute_2d_workgroups() helper to split total workgroup ID across
X/Y dimensions

* update mul_mat_reg_tile.wgsl to reconstruct linear workgroup ID from 2D
   dispatch

* update mul_mat_subgroup_matrix.wgsl to reconstruct linear workgroup ID
  from 2D dispatch

* update mul_mat.wgsl to compute global index from 2D workgroup
  coordinates

* refactor all three mul_mat dispatch paths to use the shared helper

* ggml-webgpu: add bounds checking for over-dispatched workgroups

2D workgroup dispatch can over-dispatch when total workgroups don't
divide evenly into the 65535 per-dimension limit. Extra workgroups
would compute invalid batch indices, causing memory corruption.

* add batch_idx bound check to mul_mat_reg_tile.wgsl and
mul_mat_subgroup_matrix.wgsl to prevent over-dispatched workgroups
from accessing invalid memory

* fixes test failures with large batch sizes (eg., bs=[128, 1024])

* ggml-webgpu: add back TODO for spliting large sizes into batches

* Optimize 2d workgroup provisioning

* Set some parameters that increase speed

---------

Co-authored-by: Reese Levine <redacted>
7 weeks agoggml webgpu: Clean up per-thread parameter buffer pool and job submission logic ...
Nikhil Jain [Mon, 2 Mar 2026 18:23:34 +0000 (10:23 -0800)]
ggml webgpu: Clean up per-thread parameter buffer pool and job submission logic (#19772)

* Allow webgpu_buf_pool to resize if needed, remove inflight_threads, and replace inflight_threads with num_kernels for submission

* Run clang-format

* Keep track of num batched kernels that have not been submitted yet

* Run clang-format

* Increase buf pool max size

* Increase param buf pool init size

* Remove webgpu buf pool resizing

* Merge with master

* Add buffer pool growth

* Move buffer pool growth outside of lock

* Reduce max pool size to 32

* Run clang-format

* Only resize param buf pool

7 weeks agoggml-webgpu: Support non-contiguous `src0` and overlapping `src0/src1` in binary...
Masashi Yoshimura [Mon, 2 Mar 2026 15:59:53 +0000 (00:59 +0900)]
ggml-webgpu: Support non-contiguous `src0` and overlapping `src0/src1` in binary ops (#19850)

* ggml-webgpu: Add binary op support for overlapping and non-contiguous.

* Add newline to binary.wgsl

* Append the test of binary op for src overlapping  to test_bin_bcast.

* Remove unnecessary newline.

7 weeks agovulkan: tune MMVQ for Intel Windows (#19988)
Ruben Ortlam [Mon, 2 Mar 2026 14:58:25 +0000 (15:58 +0100)]
vulkan: tune MMVQ for Intel Windows (#19988)

7 weeks agoscripts : improve get-wikitext-2.sh (#19952)
Adrien Gallouët [Mon, 2 Mar 2026 14:40:49 +0000 (15:40 +0100)]
scripts : improve get-wikitext-2.sh (#19952)

* scripts : improve get-wikitext-2.sh

Switch to sh, add curl fallback, and avoid redundant downloads

Signed-off-by: Adrien Gallouët <redacted>
* fix indent

Signed-off-by: Adrien Gallouët <redacted>
---------

Signed-off-by: Adrien Gallouët <redacted>
Signed-off-by: Adrien Gallouët <redacted>
7 weeks agoggml-cpu: optimise s390x multiply extend instructions (#20032)
Aaron Teo [Mon, 2 Mar 2026 08:23:56 +0000 (16:23 +0800)]
ggml-cpu: optimise s390x multiply extend instructions (#20032)

8 weeks agovulkan: improve partial offloading performance on AMD (#19976)
Ruben Ortlam [Sun, 1 Mar 2026 16:32:14 +0000 (17:32 +0100)]
vulkan: improve partial offloading performance on AMD (#19976)

* vulkan: fix and enable cpy_tensor_async function

* use transfer_queue for async transfers on AMD, synchronize with timeline semaphore

* update offload_op logic

* fix missing transfer submission

* disable async transfer queue on AMD GCN

* revert op batch size change

* fix cpy_tensor_async checks

8 weeks agocuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels (#19999)
oobabooga [Sun, 1 Mar 2026 05:40:22 +0000 (02:40 -0300)]
cuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels (#19999)

8 weeks agovendors : update miniaudio library to 0.11.24 (#19914)
Dmitry Atamanov [Sat, 28 Feb 2026 15:10:01 +0000 (20:10 +0500)]
vendors : update miniaudio library to 0.11.24 (#19914)

8 weeks agovendor : update cpp-httplib to 0.35.0 (#19969)
Adrien Gallouët [Sat, 28 Feb 2026 12:53:56 +0000 (13:53 +0100)]
vendor : update cpp-httplib to 0.35.0 (#19969)

Signed-off-by: Adrien Gallouët <redacted>
8 weeks agotests : model metadata loading from huggingface (#19796)
Bartowski [Sat, 28 Feb 2026 09:44:38 +0000 (04:44 -0500)]
tests : model metadata loading from huggingface (#19796)

* Add model metadata loading from huggingface for use with other tests

* Add incremental chunking instead of full redownload, fix caching issue and add warning when it fails

* Add support for split models, load metadata from each individual split file, also avoid mmproj

* Code cleanup, revert incremental downloading

* Only compile when cpp-httplib has SSL support

* Fix formatting

8 weeks agoCUDA: add CDNA3 MFMA support for flash attention MMA kernel (#19806)
Jayant Lohia [Fri, 27 Feb 2026 18:37:26 +0000 (00:07 +0530)]
CUDA: add CDNA3 MFMA support for flash attention MMA kernel (#19806)

* CUDA: add CDNA3 MFMA support for flash attention MMA kernel

Add MI300X (gfx942) MFMA tensor core flash attention using
v_mfma_f32_16x16x16_f16 (FP16 in, FP32 accumulate).

- Add FATTN_WARP_SIZE=64 for CDNA wavefront64
- Add CDNA config for head sizes 64, 80, 96, 112, 128
- Add FP16 MFMA intrinsic path in mma.cuh
- Add manual V transpose load for MFMA register layout
- Route CDNA to MMA for prompt processing, VEC for token generation
- Fix Q loading and combine stride granularity for non-power-of-2 heads

Benchmarks (Qwen2.5-1.5B Q4_K_M, MI300X):
  pp512  +7%,  pp1024 +13%,  pp2048 +23%,  pp4096 +39%
  tg128  -10% (FA overhead, VEC used for both)

All 2480 flash attention tests pass.

Ref: https://github.com/ggml-org/llama.cpp/issues/17917

* address review: replace FATTN_WARP_SIZE with constexpr, improve dispatch

- Replace #define FATTN_WARP_SIZE with constexpr int warp_size =
  ggml_cuda_get_physical_warp_size() in each device function
- Use ne[1]*gqa_ratio threshold for MMA vs tile dispatch. Benchmarked
  crossover on MI300X @ d32768 with power-of-2 GQA models:
    hsk=64  (Llama 1B, gqa=4): MMA wins at eff >= 128 (+11%)
    hsk=128 (Llama 3B, gqa=4): MMA wins at eff >= 128 (+4%)
  Unified threshold: eff_nq >= 128 for all head sizes.
- Remove VEC fallback; small batches fall through to tile kernel

* Update ggml/src/ggml-cuda/fattn.cu

* use ggml_cuda_info().devices warp_size instead of hardcoded check

---------

Co-authored-by: Johannes Gäßler <redacted>
8 weeks agoserver: Add pragma once to server-context.h (#19944)
Roj234 [Fri, 27 Feb 2026 17:28:36 +0000 (01:28 +0800)]
server: Add pragma once to server-context.h (#19944)

8 weeks agoserver: Mirroring /v1/responses to /responses to match /v1/chat/completions pattern...
Sami Kama [Fri, 27 Feb 2026 16:44:42 +0000 (08:44 -0800)]
server: Mirroring /v1/responses to /responses to match /v1/chat/completions pattern (#19873)