]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
2 months agovulkan: handle mat_mul with A matrix > 4GB (llama/16176)
Jeff Bolz [Sun, 28 Sep 2025 01:36:34 +0000 (20:36 -0500)]
vulkan: handle mat_mul with A matrix > 4GB (llama/16176)

* vulkan: handle mat_mul with A matrix > 4GB

This change splits mat_mul operations with huge A matrix into chunks in the M
dimension. This works well for stable-diffusion use cases where the im2col
matrix has very large M.

Fix the order of setting the stride in mul_mm_cm2 - setting the dimension
clobbers the stride, so stride should be set after.

* build fixes

2 months agovulkan: support arbitrary KV dimension in flash attention (llama/16160)
Jeff Bolz [Sat, 27 Sep 2025 20:43:39 +0000 (16:43 -0400)]
vulkan: support arbitrary KV dimension in flash attention (llama/16160)

The "Clamp" spec constant is already based on whether KV is a multiple of Bc,
so use that to control whether bounds checking is performed. Add bounds checking
to the scalar and coopmat1 paths. Coopmat2 didn't need any changes (the K/V
tensors are already optionally clamped, nothing else needed to be changed).

2 months agovulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)
Acly [Sat, 27 Sep 2025 20:41:03 +0000 (22:41 +0200)]
vulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)

* don't use VULKAN_HPP_DEFAULT_DISPATCH_LOADER_DYNAMIC_STORAGE which can cause conflicts if application or other libraries do the same

2 months agoCUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)
Aman Gupta [Sat, 27 Sep 2025 16:49:32 +0000 (00:49 +0800)]
CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)

* CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32

This commit adds mul_mat_id support for ncols_dst >= 16. It does this by
packing ncols_dst tiles into the blockDim.y.

My tests on a RTX 3090 show that this is faster than the cuBLAS fallback
for f16 till bs=64, and for f32 till bs=32

* Review: refactor if statement

2 months agoCUDA: refactor and deduplicate vector FA kernels (llama/16208)
Johannes Gäßler [Sat, 27 Sep 2025 16:45:07 +0000 (18:45 +0200)]
CUDA: refactor and deduplicate vector FA kernels (llama/16208)

* CUDA: refactor and deduplicate vector FA kernels

2 months agovulkan: throw system error instead of SIGABRT during init on older devices (llama...
Dmytro Minochkin [Sat, 27 Sep 2025 16:26:46 +0000 (19:26 +0300)]
vulkan: throw system error instead of SIGABRT during init on older devices (llama/16156)

* Throw system error on old Vulkan driver rather than SIGABRT

* Optionally handle any potential error in vulkan init

2 months agovulkan: support GET_ROWS for k-quants (llama/16235)
Jeff Bolz [Sat, 27 Sep 2025 10:36:11 +0000 (06:36 -0400)]
vulkan: support GET_ROWS for k-quants (llama/16235)

The dequantize functions are copy/pasted from mul_mm_funcs.comp with very few
changes - add a_offset and divide iqs by 2. It's probably possible to call
these functions from mul_mm_funcs and avoid the duplication, but I didn't go
that far in this change.

2 months agodevops: add s390x & ppc64le CI (llama/15925)
Aaron Teo [Fri, 26 Sep 2025 18:03:33 +0000 (02:03 +0800)]
devops: add s390x & ppc64le CI (llama/15925)

* devops: move s390x and ppc64le ci build

we have access to ubuntu-24.04-s390x and ppc64le images now

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le for now since they have compiler errors

Signed-off-by: Aaron Teo <redacted>
* devops: stop warnings as errors

Signed-off-by: Aaron Teo <redacted>
* devops: switch to non-macro flag

Signed-off-by: Aaron Teo <redacted>
* devops: going the llama macro route

Signed-off-by: Aaron Teo <redacted>
* devops: add big-endian gguf test models

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le to test s390x, check test build

Signed-off-by: Aaron Teo <redacted>
* devops: dup .gguf.inp files for big-endian tests

Signed-off-by: Aaron Teo <redacted>
* devops: dup .gguf.out files for big-endian too

Signed-off-by: Aaron Teo <redacted>
* devops: add python setup and endian byteswap

Signed-off-by: Aaron Teo <redacted>
* devops: pooring thing does not have s390x python3

Signed-off-by: Aaron Teo <redacted>
* devops: add missing rust compiler for s390x

Signed-off-by: Aaron Teo <redacted>
* devops: try rust actions runner

Signed-off-by: Aaron Teo <redacted>
* Revert "devops: try rust actions runner"

This reverts commit 3f8db04356033d6c1d7eccc75ca396bc5298250c.

Signed-off-by: Aaron Teo <redacted>
* devops: try a different path for rust

Signed-off-by: Aaron Teo <redacted>
* devops: dump home directory and user info

Signed-off-by: Aaron Teo <redacted>
* devops: install gguf-py only

Signed-off-by: Aaron Teo <redacted>
* devops: missed relative path

Signed-off-by: Aaron Teo <redacted>
* devops: remove big-endian files since local swapping is working

Signed-off-by: Aaron Teo <redacted>
* devops: revert test-tokenizer-0 cmakelists

Signed-off-by: Aaron Teo <redacted>
* Fix unicode flags conversion from and to uint16_t

Bitfields are allocated in different order on s390x

Signed-off-by: Aaron Teo <redacted>
* Simplify byteswap command

Signed-off-by: Aaron Teo <redacted>
* Add byteswapping and git-lfs for test-tokenizers-ggml-vocabs

Signed-off-by: Aaron Teo <redacted>
* Fix endianness detection in vocab loader

Signed-off-by: Aaron Teo <redacted>
* Disable test-thread-safety on s390x

In this test a model is downloaded,
then immediately loaded to check if more downloads are needed,
and then used for test.

There is no clean way to separate all those steps
 to add byteswapping between them, so just skip this test.

Signed-off-by: Aaron Teo <redacted>
* Fix q8_0 test in test-quantize-fns

vec_signed uses unexpected rounding mode.
Explicitly use different rounding function.

Signed-off-by: Aaron Teo <redacted>
* devops: add big-endian stories260K

Signed-off-by: Aaron Teo <redacted>
* devops: add s390x test-eval-callback

Signed-off-by: Aaron Teo <redacted>
* devops: fix test does not exist

Signed-off-by: Aaron Teo <redacted>
* devops: fix model not found llama-eval-callback

Signed-off-by: Aaron Teo <redacted>
* Fix q3_K dot product error in test-quantize-fns on s390x

Array q8bytes had only 4 elements allocated, but 8 elements accessed.
This lead to write out of bounds and later read of overwritten values out of bounds
and incorrect result.

Signed-off-by: Aaron Teo <redacted>
* devops: re-enable ppc64le for testing

Signed-off-by: Aaron Teo <redacted>
* devops: activate test-thread-safety for s390x

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le tests

for some reason it keeps failing test-thread-safety tests and I do not
    have a machine that is able to replicate the tests.

Signed-off-by: Aaron Teo <redacted>
* devops: LLAMA_FATAL_WARNINGS=ON

Signed-off-by: Aaron Teo <redacted>
* Correct repository URL for s390x for test-thread-safety model

Signed-off-by: Aaron Teo <redacted>
* Fix fs_get_cache_directory

Ensure it works even if both XDG_CACHE_HOME and HOME are unset.
This might happen in containers.

Signed-off-by: Aaron Teo <redacted>
* Re-enable CI for ppc64le

Signed-off-by: Aaron Teo <redacted>
* Fortify ggml_rope_impl

Only memcpy data from sections argument if it's non-NULL.

Signed-off-by: Aaron Teo <redacted>
* Add TODO in struct unicode_cpt_flags to reimplement it in endian-independent way

* Update URL for big-endian model

* Update .github/workflows/build.yml

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update remaining mentions of BE models to ggml-org/models repo

---------

Signed-off-by: Aaron Teo <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agometal : report OOM errors (llama/16274)
Georgi Gerganov [Fri, 26 Sep 2025 11:14:28 +0000 (14:14 +0300)]
metal : report OOM errors (llama/16274)

2 months agocommon : use cpp-httplib as a cURL alternative for downloads (llama/16185)
Adrien Gallouët [Fri, 26 Sep 2025 11:12:19 +0000 (13:12 +0200)]
common : use cpp-httplib as a cURL alternative for downloads (llama/16185)

* vendor : update httplib

Signed-off-by: Adrien Gallouët <redacted>
* common : use cpp-httplib as a cURL alternative for downloads

The existing cURL implementation is intentionally left untouched to
prevent any regressions and to allow for safe, side-by-side testing by
toggling the `LLAMA_CURL` CMake option.

Signed-off-by: Adrien Gallouët <redacted>
* ggml : Bump to Windows 10

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

Signed-off-by: Adrien Gallouët <redacted>
2 months agoggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)
Aaron Teo [Fri, 26 Sep 2025 10:27:25 +0000 (18:27 +0800)]
ggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)

* ggml-cpu: impl mxfp4 s390x

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: missing s = sumf

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix incorrect kval_mxfp4 type

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: rework mxfp4

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: missing delta calc

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix typo

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix typo for vec_splats

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: expand to 2 blocks per loop

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add unroll to boost perf

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: back to 1 block per loop to test perf

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: back to 1 block per loop to test perf"

This reverts commit 1fe55724e2dc295701101bf838bdd4a512237492.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: rm unroll from single block

Signed-off-by: Aaron Teo <redacted>
---------

Signed-off-by: Aaron Teo <redacted>
2 months agomusa: fix build warnings (llama/15611)
R0CKSTAR [Fri, 26 Sep 2025 00:56:10 +0000 (08:56 +0800)]
musa: fix build warnings (llama/15611)

Signed-off-by: Xiaodong Ye <redacted>
2 months agoCUDA: add a fused top-K MoE kernel (llama/16130)
Aman Gupta [Thu, 25 Sep 2025 14:35:05 +0000 (22:35 +0800)]
CUDA: add a fused top-K MoE kernel (llama/16130)

* CUDA: add a fused top-K MoE kernel

This kernel does the following:
1. softmax over the logits per token [n_experts, n_tokens]
2. argmax reduce over the top-k (n_experts_used) logits
3. write weights + ids to global memory

It is intended as fusion of softmax->top-k->get_rows pipeline for MoE models

* Refactor into ggml_cuda_should_use_topk_moe

* Review: Use better coalescing pattern, use WARP_SIZE, store logits into registers before

* Review: format + micro-optimizations

* Fix bug: fix tie breakers

* Add optional norm + clean-up code

* Use smem for final write

* Add bounds check

* Use better memory pattern for writeback

2 months agoggml : fix loongarch lsx compilation error (llama/15864)
junchao-zhao [Thu, 25 Sep 2025 09:22:55 +0000 (17:22 +0800)]
ggml : fix loongarch lsx compilation error (llama/15864)

2 months agoggml : remove -dev suffix from release version (#1355)
Daniel Bevenius [Fri, 26 Sep 2025 15:34:42 +0000 (17:34 +0200)]
ggml : remove -dev suffix from release version (#1355)

This commit removes the `-dev` suffix from the version string in
CMakeLists.txt and the release script. The version will now be
just be formatted as `MAJOR.MINOR.PATCH`.

2 months agopkg-config: include the new GGML_VERSION as a version (#1348)
Christoph Reiter [Thu, 25 Sep 2025 16:59:38 +0000 (18:59 +0200)]
pkg-config: include the new GGML_VERSION as a version (#1348)

Instead of hardcoding 0.0.0

If the 0.9.2-dev version is supposed to mean "less than" 0.9.2 instead of equal,
then the .pc version format would need to be changed to 0.9.2dev, since pkgconf
uses RPM version comparison. But let's keep it simple for now.

The version in CMake was added in #1336

2 months agoexamples : fix typo mismatch in gpt (#1349)
hebangwen [Thu, 25 Sep 2025 15:39:30 +0000 (23:39 +0800)]
examples : fix typo mismatch in gpt (#1349)

2 months agoggml : bump version to 0.9.3 (#1353) v0.9.3
Daniel Bevenius [Thu, 25 Sep 2025 12:39:05 +0000 (14:39 +0200)]
ggml : bump version to 0.9.3 (#1353)

2 months agoscripts : refactor release script into prepare and finalize stages (#1352)
Daniel Bevenius [Thu, 25 Sep 2025 10:38:50 +0000 (12:38 +0200)]
scripts : refactor release script into prepare and finalize stages (#1352)

This commit splits the release process into two distinct stages:

* prepare_release: This stage handles all updating the version and
  creating a new branch with the version change. This should then be
  use to open a PR for review. Once the PR has been merged the
  finalize_release stage can be run.

* finalize_release: This stage must be run on master and master must
  have the version bump commit (this is checked for). This stage
  handles tagging, and also creates a new branch for the update of the
  new development version. The tag should then be pushed to the remote
  which will trigger the release process on GitHub. The branch should be
  used to open a new PR for the development version update.

2 months agoscripts : fix next dev version calculation [no ci] (#1351)
Daniel Bevenius [Thu, 25 Sep 2025 09:43:51 +0000 (11:43 +0200)]
scripts : fix next dev version calculation [no ci] (#1351)

This commit updates the release script to correctly calculate the next
development version after a release.

The motivation for this is that the script currently increments the new
version for the next development cycle but the version is already set to
the updated version so this should not increment again.

2 months agosync : llama.cpp
Georgi Gerganov [Thu, 25 Sep 2025 08:40:17 +0000 (11:40 +0300)]
sync : llama.cpp

2 months agometal : fuse NORM + MUL + ADD, support non-multiples of 4 (llama/16220)
Georgi Gerganov [Thu, 25 Sep 2025 08:30:16 +0000 (11:30 +0300)]
metal : fuse NORM + MUL + ADD, support non-multiples of 4 (llama/16220)

* metal : fuse NORM + MUL + ADD

* metal : support norms of non-multiple of 4

* cont : fix comment [no ci]

2 months agometal : relax reorder conditions (llama/16216)
Georgi Gerganov [Thu, 25 Sep 2025 08:29:42 +0000 (11:29 +0300)]
metal : relax reorder conditions (llama/16216)

2 months agometal : restore im2col perf (llama/16219)
Georgi Gerganov [Thu, 25 Sep 2025 08:29:08 +0000 (11:29 +0300)]
metal : restore im2col perf (llama/16219)

2 months agosync : llama.cpp
Georgi Gerganov [Thu, 25 Sep 2025 08:20:29 +0000 (11:20 +0300)]
sync : llama.cpp

2 months agorpc : use ggml logging facilities
Radoslav Gerganov [Thu, 25 Sep 2025 07:20:02 +0000 (10:20 +0300)]
rpc : use ggml logging facilities

Use RPC_DEBUG environment variable to enable debug messages.
Add helper macro LOG_DBG() which does an early
check of the env var before calling GGML_LOG_DEBUG().
Make sure we log a debug message for every server function.

2 months agoci: run the x64 and arm ci on the github machines instead (llama/16183)
Eve [Thu, 25 Sep 2025 05:06:06 +0000 (05:06 +0000)]
ci: run the x64 and arm ci on the github machines instead (llama/16183)

* run the x64 ci on regular machines

* set up the same thing for arm

fix test-quantize-perf just like #12306

* try to disable sve

* add another sve run

2 months agollama: print memory breakdown on exit (llama/15860)
Johannes Gäßler [Wed, 24 Sep 2025 14:53:48 +0000 (16:53 +0200)]
llama: print memory breakdown on exit (llama/15860)

* llama: print memory breakdown on exit

2 months agoggml : split graph allocations according to backend max buffer size (llama/15815)
Acly [Wed, 24 Sep 2025 14:17:49 +0000 (16:17 +0200)]
ggml : split graph allocations according to backend max buffer size (llama/15815)

* ggml : make gallocr respect the backend's max buffer size

* if the graph requires more memory than can fit into a single allocation, split it into multiple backend buffers
* vulkan: report the actual max  allocation size in buffer type  interface

* fix missing newline, apple-clang warning

* track size of individual chunks in ggml_dyn_tallocr and raise max chunks.
revert to use suballocation_block_size as max chunk size for vulkan.

* track (chunk, offset) pairs instead of "global" offsets through gallocr.

* simpler, don't need loops to map between local/global offsets
* touches more code

* fix dyn_tallocr_max_size and initialization

* fix memory leak when buffers are reused due to same buffer type appearing multiple times

* make vbuffer allocation follow the same logic as backend_buffer did before

* continue to use leftover unallocated space of previous chunks after a new one has been created

* treat free blocks of each chunk as separate list
* they're still allocated together, but start/end of each chunk is tracked, and allocate/free iterate over sub-ranges
* exhaust freed blocks of all chunks before considering their last blocks with unallocated space
* start with 0 chunks/blocks and create chunks as needed
* allow the last chunk to grow beyond max size

* refactor: move adding new free block and new chunk into separate functions

* allocate chunks individually with a separate free-blocks list for each one

* needs a bit more memory/allocations/indirections, but code is simpler

* fix warnings (missing static) & debug checks

2 months agoggml-cpu: Respect cpumask settings (llama/16164)
Xiangyan Sun [Tue, 23 Sep 2025 08:58:12 +0000 (01:58 -0700)]
ggml-cpu: Respect cpumask settings (llama/16164)

2 months agoggml : fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl (llama/15928)
Sigbjørn Skjæret [Tue, 23 Sep 2025 08:25:20 +0000 (10:25 +0200)]
ggml : fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl (llama/15928)

* fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl

* change initialization to true

2 months agozdnn: refactor codebase + add docs (llama/16178)
Aaron Teo [Tue, 23 Sep 2025 06:53:05 +0000 (14:53 +0800)]
zdnn: refactor codebase + add docs (llama/16178)

* zdnn: initial matmul refactor

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: rm static from funcs

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: update ggml-zdnn.h

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: change header files to hpp

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: switch to common.hpp

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: move mulmat forward around

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: rm inline from utils

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: code cleanup

Signed-off-by: Aaron Teo <redacted>
* docs: add zDNN docs

Signed-off-by: Aaron Teo <redacted>
---------

Signed-off-by: Aaron Teo <redacted>
2 months agoggml-cpu : fix typo in gemm comments [no ci] (llama/16189)
Daniel Bevenius [Tue, 23 Sep 2025 03:59:03 +0000 (05:59 +0200)]
ggml-cpu : fix typo in gemm comments [no ci] (llama/16189)

2 months agoggml : implement set_rows with i32 index (llama/16159)
Sigbjørn Skjæret [Mon, 22 Sep 2025 17:13:00 +0000 (19:13 +0200)]
ggml : implement set_rows with i32 index (llama/16159)

* implement set_rows with i32 index

* template fix

* test quantized path

warnings--

* Apply suggestions from code review

Co-authored-by: Georgi Gerganov <redacted>
* forgotten name change

* deduplicate cuda/sycl and test-fix

* indent++

* vulkan: support set_rows with i32 index type (llama/16162)

* disable i32 index for webgpu for now

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Jeff Bolz <redacted>
2 months agoggml : extend ggml_can_fuse to work with non-sequential nodes (llama/16123)
Georgi Gerganov [Mon, 22 Sep 2025 08:12:37 +0000 (11:12 +0300)]
ggml : extend ggml_can_fuse to work with non-sequential nodes (llama/16123)

* ggml : extend ggml_can_fuse to work with non-sequential nodes in the graph

* cont : fix wrong bounds check condition

* cont : remove unnecessary overload

2 months agoggml : add ggml_op_is_empty (llama/16122)
Georgi Gerganov [Mon, 22 Sep 2025 08:12:09 +0000 (11:12 +0300)]
ggml : add ggml_op_is_empty (llama/16122)

* ggml : add ggml_op_is_empty

* ggml : move to ggml-impl.h

2 months agoVulkan: add conv_transpose_2d operation (llama/16022)
Shin-myoung-serp [Mon, 22 Sep 2025 08:04:01 +0000 (17:04 +0900)]
Vulkan: add conv_transpose_2d operation (llama/16022)

* Vulkan: add conv_transpose_2d operation

* Vulkan: fix typo in conv_transpose_2d shader(s0mp, s0L, s1mp, s1L)

* Vulkan: fix incorrect indentation in conv_transpose_2d shader

* Vulkan: add checking the push constants size limit and reuse conv2d_mm.comp for conv_transpose_2d operation

* Vulkan: revert the order of the index calculation and bound check in conv_2d shader

* Vulkan: explicity check push constants limit in supports_op() for conv_transpose_2d operation.

* Vulkan: remove unnecessary lower bound checks for H/W_idx in the conv_2d shader.

2 months agovulkan: add RTE variants of exp shader (llama/16165)
Jeff Bolz [Mon, 22 Sep 2025 05:37:17 +0000 (00:37 -0500)]
vulkan: add RTE variants of exp shader (llama/16165)

This fixes some failures on Turing where "round to zero" rounds to the max f16
value but the CPU reference value is infinite.

2 months agovulkan: vec dot matrix multiplication fix (llama/16151)
Ruben Ortlam [Mon, 22 Sep 2025 05:22:43 +0000 (07:22 +0200)]
vulkan: vec dot matrix multiplication fix (llama/16151)

* vulkan: fix matrix multiplication index calculation for odd m/n and odd k in combination with batching

* add odd m/n + odd k test with batching

2 months agoopencl: fix concat crash on win arm64 with Adreno (llama/15944)
lhez [Sun, 21 Sep 2025 23:42:10 +0000 (16:42 -0700)]
opencl: fix concat crash on win arm64 with Adreno (llama/15944)

2 months agoopencl: initial `q8_0` mv support (llama/15732)
lhez [Sun, 21 Sep 2025 21:48:44 +0000 (14:48 -0700)]
opencl: initial `q8_0` mv support (llama/15732)

2 months agovulkan: optimize UMA buffer operations and fix driver hangs (llama/16059)
Giuseppe Scrivano [Sun, 21 Sep 2025 06:31:55 +0000 (08:31 +0200)]
vulkan: optimize UMA buffer operations and fix driver hangs (llama/16059)

* vulkan: optimize UMA buffer operations and fix driver hangs

The previous implementation was blocking the GPU for extended periods,
causing the i915 driver to reset the context due to the hangcheck
protection.

[32628.443070] i915 0000:00:02.0: [drm] GPU HANG: ecode 12:1:85dffffb, in llama-server [194114]
[32628.443091] i915 0000:00:02.0: [drm] llama-server[194114] context reset due to GPU hang

* vulkan: implement deferred_memset on UMA

---------

Signed-off-by: Giuseppe Scrivano <redacted>
2 months agovulkan: fix validation error about VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR...
Jeff Bolz [Sun, 21 Sep 2025 06:23:37 +0000 (01:23 -0500)]
vulkan: fix validation error about VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR (llama/16086)

2 months agoggml : prepare for development of 0.9.2-dev
Georgi Gerganov [Sat, 20 Sep 2025 13:44:23 +0000 (16:44 +0300)]
ggml : prepare for development of 0.9.2-dev

2 months agoggml : bump version to 0.9.1
Georgi Gerganov [Sat, 20 Sep 2025 13:44:23 +0000 (16:44 +0300)]
ggml : bump version to 0.9.1

2 months agoscripts : fix sed usage to work on Mac (#1345)
Georgi Gerganov [Sat, 20 Sep 2025 13:44:08 +0000 (16:44 +0300)]
scripts : fix sed usage to work on Mac (#1345)

2 months agotests : adjust to new timestep_embedding operator
Georgi Gerganov [Sat, 20 Sep 2025 10:16:10 +0000 (13:16 +0300)]
tests : adjust to new timestep_embedding operator

2 months agosync : llama.cpp
Georgi Gerganov [Sat, 20 Sep 2025 10:09:41 +0000 (13:09 +0300)]
sync : llama.cpp

2 months agovulkan: use vec dot for matrix matrix multiplications (llama/16056)
Ruben Ortlam [Sat, 20 Sep 2025 08:42:56 +0000 (10:42 +0200)]
vulkan: use vec dot for matrix matrix multiplications (llama/16056)

* vulkan: Change the mul_mm shared memory and register caching system to use vec2 instead of scalars, to enable using dot2 instructions

* use fma instead of dot to fix Nvidia and Apple performance issues

2 months agoggml : refactor forward_dup for cpu backend (llama/16062)
Xuan-Son Nguyen [Fri, 19 Sep 2025 04:31:56 +0000 (11:31 +0700)]
ggml : refactor forward_dup for cpu backend (llama/16062)

* ggml : refactor forward_dup for cpu backend

* clean up a bit

* add quant/dequant perf test

2 months agoggml-amx : fix ggml_amx_init() on generic Linux (llama/16049)
Adrien Gallouët [Thu, 18 Sep 2025 21:07:26 +0000 (23:07 +0200)]
ggml-amx : fix ggml_amx_init() on generic Linux (llama/16049)

Generalize Linux check to `__linux__` to support non-glibc systems (like musl).
Also, return `false` on unknown/untested OS.

Without this commit, the code compiles (with warnings) but fails:

    register_backend: registered backend CPU (1 devices)
    register_device: registered device CPU (Intel(R) Xeon(R) Platinum 8488C)
    build: 6487 (51c4cac6) with x86_64-linux-musl-gcc (GCC) 15.1.0 for x86_64-linux-musl (debug)
    system info: n_threads = 8, n_threads_batch = 8, total_threads = 16
    ....
    print_info: n_ctx_orig_yarn  = 262144
    print_info: rope_finetuned   = unknown
    print_info: model type       = 4B
    Illegal instruction (core dumped)

Signed-off-by: Adrien Gallouët <redacted>
2 months agocmake : fix static linking for OpenMP on Unix-like systems (llama/16031)
Adrien Gallouët [Thu, 18 Sep 2025 21:07:18 +0000 (23:07 +0200)]
cmake : fix static linking for OpenMP on Unix-like systems (llama/16031)

When compiling with GGML_STATIC=ON, the build process would produce a
binary that was still dynamically linked to OpenMP. This defeats the
purpose of a static build:

    $ cmake -B build \
            -DBUILD_SHARED_LIBS=OFF \
            -DLLAMA_CURL=OFF \
            -DGGML_CCACHE=OFF \
            -DGGML_NATIVE=OFF \
            -DGGML_STATIC=ON

    $ ldd llama-server
            linux-vdso.so.1 (0x0000e1a434e3b000)
            libgomp.so.1 => /lib/aarch64-linux-gnu/libgomp.so.1 (0x0000e1a4345a0000)
            libstdc++.so.6 => /lib/aarch64-linux-gnu/libstdc++.so.6 (0x0000e1a434300000)
            libm.so.6 => /lib/aarch64-linux-gnu/libm.so.6 (0x0000e1a434240000)
            libgcc_s.so.1 => /lib/aarch64-linux-gnu/libgcc_s.so.1 (0x0000e1a434200000)
            libc.so.6 => /lib/aarch64-linux-gnu/libc.so.6 (0x0000e1a434030000)
            /lib/ld-linux-aarch64.so.1 (0x0000e1a434df0000)

This commit resolves the issue by modifying `CMAKE_FIND_LIBRARY_SUFFIXES`
to prioritize `.a` files, forcing CMake to link the static version of
the library.

Signed-off-by: Adrien Gallouët <redacted>
2 months agoopencl: optimize mxfp4 kernels (llama/16037)
Shawn Gu [Thu, 18 Sep 2025 19:03:34 +0000 (12:03 -0700)]
opencl: optimize mxfp4 kernels (llama/16037)

- flatten mxfp4 and packed fp4->fp16 bit-wise convert function (replace lut)
- MoE kernel optimizations

---------

Co-authored-by: Li He <redacted>
2 months agorename optimize_graph to graph_optimize (llama/16082)
Jeff Bolz [Thu, 18 Sep 2025 18:46:17 +0000 (13:46 -0500)]
rename optimize_graph to graph_optimize (llama/16082)

2 months agoCUDA: Optimize PAD_REFLECT_1D (llama/15957)
Bowen Han [Thu, 18 Sep 2025 18:26:03 +0000 (11:26 -0700)]
CUDA: Optimize PAD_REFLECT_1D (llama/15957)

* CUDA: Optimize PAD_REFLECT_1D
feat: add more test cases for PAD_REFLECT_1D

* use fast_div to improve performance

* Apply suggestion from JohannesGaessler

Co-authored-by: Johannes Gäßler <redacted>
* Apply suggestion from JohannesGaessler

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

* use a concise expression to further speedup the cuda kernel

---------

Co-authored-by: Johannes Gäßler <redacted>
2 months agoCUDA: fix compilation on CC 6.0 (llama/16091)
Johannes Gäßler [Thu, 18 Sep 2025 17:28:32 +0000 (19:28 +0200)]
CUDA: fix compilation on CC 6.0 (llama/16091)

2 months agometal : use function constants for mul_mv_ext kernels (llama/16074)
Georgi Gerganov [Thu, 18 Sep 2025 13:28:41 +0000 (16:28 +0300)]
metal : use function constants for mul_mv_ext kernels (llama/16074)

* metal : use function constants for mul_mv_ext kernels

ggml-ci

* metal : remove NW template argument

ggml-ci

* metal : adjust constants

ggml-ci

2 months agocuda : add missing F32<->I32 entries in ggml_cuda_cpy_fn (llama/16060)
Sigbjørn Skjæret [Thu, 18 Sep 2025 11:28:22 +0000 (13:28 +0200)]
cuda : add missing F32<->I32 entries in ggml_cuda_cpy_fn (llama/16060)

2 months agometal : improve F32, F16 and BF16 mat-vec multiplication (llama/16057)
Georgi Gerganov [Thu, 18 Sep 2025 09:33:45 +0000 (12:33 +0300)]
metal : improve F32, F16 and BF16 mat-vec multiplication (llama/16057)

* metal : improve F32, F16 and BF16 mat-vec multiplication

ggml-ci

* metal : make the NSG a function constant in mul_mv kernels

ggml-ci

2 months agometal : avoid call free for non-owned buffer (llama/16067)
Jhen-Jie Hong [Thu, 18 Sep 2025 07:06:48 +0000 (15:06 +0800)]
metal : avoid call free for non-owned buffer (llama/16067)

2 months agometal : handle nil cv during pipeline creation (llama/16065)
Georgi Gerganov [Thu, 18 Sep 2025 07:03:24 +0000 (10:03 +0300)]
metal : handle nil cv during pipeline creation (llama/16065)

ggml-ci

2 months agoCANN: Remove print (llama/16044)
Chenguang Li [Thu, 18 Sep 2025 01:26:33 +0000 (09:26 +0800)]
CANN: Remove print (llama/16044)

Signed-off-by: noemotiovon <redacted>
2 months agoGGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (llama/16018)
Reese Levine [Wed, 17 Sep 2025 20:09:40 +0000 (13:09 -0700)]
GGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (llama/16018)

* Add paramater buffer pool, batching of submissions, refactor command building/submission

* Add header for linux builds

* Free staged parameter buffers at once

* Format with clang-format

* Fix thread-safe implementation

* Use device implicit synchronization

* Update workflow to use custom release

* Remove testing branch workflow

* some f32 tests passing

* Disable set_rows until it's implemented

* f32 add all tests passing

* Begin work on set_rows

* Work on set rows

* Add error buffers for reporting unsupported SET_ROWS indices

* Remove extra comments

* Add templated addition, clean up code

* Get addition and multiplication working

* Implement rms_norm

* Add get_rows implementation

* Add new get_rows files

* Refactor use of wg size entry

* Fix compilation

* Try manually unrolled q4_0 quant

* Revert "Try manually unrolled q4_0 quant"

This reverts commit 77f8b96515f7e640ae4b0e44f066321fbc4a6166.

* Move to constant max wg size

* Check for tensor size in supports_op

* Vectorize f32 and change default workgroup size

* Move f32 get_rows from < 4 to % 4 != 0

* fix linter errors

* Add in-place tests

---------

Co-authored-by: Neha Abbas <redacted>
2 months agometal : refactor + optimize v2 (llama/15995)
Georgi Gerganov [Sat, 20 Sep 2025 10:09:06 +0000 (13:09 +0300)]
metal : refactor + optimize v2 (llama/15995)

2 months agosync : llama.cpp
Georgi Gerganov [Sat, 20 Sep 2025 10:07:59 +0000 (13:07 +0300)]
sync : llama.cpp

2 months agoCUDA: fix FA occupancy, optimize tile kernel (llama/15982)
Johannes Gäßler [Wed, 17 Sep 2025 13:32:42 +0000 (15:32 +0200)]
CUDA: fix FA occupancy, optimize tile kernel (llama/15982)

2 months agovulkan: automatically remove unsupported devices (llama/15976)
Eve [Wed, 17 Sep 2025 07:35:37 +0000 (07:35 +0000)]
vulkan: automatically remove unsupported devices (llama/15976)

* remove unsupported vulkan devices

* make this happen during selection instead

* pass by reference

2 months agoCANN: Optimize ggml_cann_set_device (llama/15935)
Chenguang Li [Wed, 17 Sep 2025 06:33:08 +0000 (14:33 +0800)]
CANN: Optimize ggml_cann_set_device (llama/15935)

* CANN: Fix ggml_cann_set_device to avoid redundant device switches

- Added a check to skip aclrtSetDevice if the current device is already set.
- Prevents unnecessary context switches while keeping thread/device consistency.

* CANN: add device default id

2 months agoggml : fix padding in timestep embedding kernels (llama/15932)
Daniel Bevenius [Tue, 16 Sep 2025 13:25:57 +0000 (15:25 +0200)]
ggml : fix padding in timestep embedding kernels (llama/15932)

* ggml : remove adding extra dim timestep embedding

This commit updates the ggml_timestep_embedding function to no longer
add an extra dimension when the specified dimension is odd.

The motivation for this change is that this introduces an unnecessary
dimension when the dimension is odd, which caused an issue in the
kernels which were not expecting this extra dimension and it resulted in
uninitialized memory for the second to last dimension.

* ggml-cuda : fix padding in timestep embedding kernel

This commit removes the zeroing out of the last dimension now that we
are not adding the extra padding dimension.

* ggml-metal : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel

* ggml-opencl : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-sycl : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-vulkan : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-cpu : fix padding in timestep embedding function

This commit removes the zeroing out of the last dimension now that we
are not adding the extra padding dimension.

2 months agoCUDA: fix im2col_3d to respect non-contiguous inputs (views) (llama/15956)
Jake Karnes [Mon, 15 Sep 2025 22:28:31 +0000 (16:28 -0600)]
CUDA: fix im2col_3d to respect non-contiguous inputs (views) (llama/15956)

* fix im2col_3d to respect non-contiguous inputs (views)

The CUDA 3D im2col kernel computed source addresses assuming compact layout (products of dims), ignoring nb[] strides.

This patch switches im2col_3d source indexing to use true strides derived from src1->nb[] (in elements), mirroring the approach used in the 2D CUDA im2col path. Destination indexing is unchanged.

* use ggml_element_size() for src strides

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

Co-authored-by: Johannes Gäßler <redacted>
2 months agoSYCL: Add COUNT_EQUAL operator support (llama/15991)
yael-works [Mon, 15 Sep 2025 16:51:35 +0000 (19:51 +0300)]
SYCL: Add COUNT_EQUAL operator support (llama/15991)

* SYCL: Add COUNT_EQUAL operator support (rebased on master)

* SYCL: remove duplicate op_count_equal definition

* tests: remove test_count_equal_typed and use test_count_equal for all cases

* tests: keep only I32 case for COUNT_EQUAL as suggested

* tests: keep only I32 case for COUNT_EQUAL as requested

2 months agoCUDA: some micro-optimizations in mmf.cuh for mul_mat_id (llama/15926)
Aman Gupta [Mon, 15 Sep 2025 09:35:11 +0000 (17:35 +0800)]
CUDA: some micro-optimizations in mmf.cuh for mul_mat_id (llama/15926)

2 months agometal : remove memory pools (llama/15966)
Georgi Gerganov [Sun, 14 Sep 2025 19:02:32 +0000 (22:02 +0300)]
metal : remove memory pools (llama/15966)

* metal : remove mem pool usage

ggml-ci

* metal : remove mem pool implementation

ggml-ci

* metal : take into account the actual allocated memory of the tensor

ggml-ci

* cont : use ggml_backend_buft_get_alloc_size

ggml-ci

* cont : improve, comments

ggml-ci

* cont : add functions for the extra tensor sizes

* metal : add comments

ggml-ci

* metal : implement .get_alloc_size for the rest of the buffer types

ggml-ci

* metal : remove ggml_metal_heap

ggml-ci

2 months agoVulkan: Clean up mul_mm shader (llama/15987)
Ruben Ortlam [Sun, 14 Sep 2025 14:56:28 +0000 (16:56 +0200)]
Vulkan: Clean up mul_mm shader (llama/15987)

* vulkan: move mul_mm dequantization steps into a separate file and functions

* improve mul_mm vector load code

* fix debug mode issues and warnings

2 months agometal : fix kernel requirements (llama/15983)
Georgi Gerganov [Sun, 14 Sep 2025 12:33:22 +0000 (15:33 +0300)]
metal : fix kernel requirements (llama/15983)

* metal : fix kernel requirements

ggml-ci

* cont : fix supports_op

* cont : fix supports_op for ARGMAX

2 months agoggml-zdnn: rm user mapped buffers (llama/15965)
Aaron Teo [Sun, 14 Sep 2025 05:37:03 +0000 (13:37 +0800)]
ggml-zdnn: rm user mapped buffers (llama/15965)

* ggml-zdnn: rm user mapped buffers

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: rm dead code

Signed-off-by: Aaron Teo <redacted>
* ggml-zdnn: attempt to fix missing extra data buffer free

Signed-off-by: Aaron Teo <redacted>
---------

Signed-off-by: Aaron Teo <redacted>
2 months agovulkan: fix failing dequant shaders (llama/15862)
Jeff Bolz [Sat, 13 Sep 2025 15:29:43 +0000 (16:29 +0100)]
vulkan: fix failing dequant shaders (llama/15862)

* vulkan: fix failing dequant shaders

* add missing const

2 months agovulkan: initialize vulkan-hpp to allow using extension function pointers (llama/15705)
Jeff Bolz [Sat, 13 Sep 2025 15:23:30 +0000 (16:23 +0100)]
vulkan: initialize vulkan-hpp to allow using extension function pointers (llama/15705)

Use this to query register count for shader compiles on NVIDIA. Currently
this is only for performance debug, but it could eventually be used in some
heuristics like split_k.

2 months agometal : refactor kernel loading (llama/15964)
Georgi Gerganov [Sat, 13 Sep 2025 13:24:22 +0000 (16:24 +0300)]
metal : refactor kernel loading (llama/15964)

* metal : refactor bin kernels loading

ggml-ci

* metal : refactor rms kernel loading

ggml-ci

* ci : try to add memory leaks check

ggml-ci

* ci : try to enable memory leak detection for Mac

* cont : seems to be working

2 months agometal : allow ops to run concurrently (llama/15929)
Georgi Gerganov [Sat, 13 Sep 2025 10:54:28 +0000 (13:54 +0300)]
metal : allow ops to run concurrently (llama/15929)

* metal : run graphs ops concurrently

ggml-ci

* cont : add flags for debugging and disabling concurrency

ggml-ci

* cont : refactor and handle fusing

ggml-ci

* cont : simplify - no need to use GPU address

ggml-ci

* cont : prepare mem ranges for reuse + add ggml-metal-common.cpp

ggml-ci

* cont : avoid redundant keywords in cpp [no ci]

* metal : reorder graph for better concurrency

ggml-ci

* metal : fix race on mem pool buffers

ggml-ci

* cont : add env GGML_METAL_GRAPH_OPTIMIZE_DISABLE

ggml-ci

* cont : refactor, optimize, add comments

ggml-ci

* cont : refactor ggml-metal.m

ggml-ci

* minor : update logs [no ci]

2 months agometal : fix memory leaks (llama/15962)
Georgi Gerganov [Sat, 13 Sep 2025 09:45:04 +0000 (12:45 +0300)]
metal : fix memory leaks (llama/15962)

ggml-ci

2 months agoggml-zdnn: fix #15414, activate FP16 and BF16 acceleration and incorrect zTensor...
Aaron Teo [Fri, 12 Sep 2025 18:39:52 +0000 (02:39 +0800)]
ggml-zdnn: fix #15414, activate FP16 and BF16 acceleration and incorrect zTensor free (llama/15839)

2 months agoVulkan iGPU device selection overhaul and PCI ID API support (llama/15947)
Ruben Ortlam [Fri, 12 Sep 2025 11:24:21 +0000 (13:24 +0200)]
Vulkan iGPU device selection overhaul and PCI ID API support (llama/15947)

* vulkan: implement ggml igpu device type, implement pci id support

* fix compiler warning

* prevent printf overflow warning

2 months agovulkan: Make device memory check more portable (llama/15939)
Mathieu Baudier [Fri, 12 Sep 2025 07:06:20 +0000 (09:06 +0200)]
vulkan: Make device memory check more portable (llama/15939)

2 months agoRevert "sycl: add usage of enqueue_functions extension (llama/14244)" (llama/15910)
Neo Zhang Jianyu [Fri, 12 Sep 2025 01:15:12 +0000 (09:15 +0800)]
Revert "sycl: add usage of enqueue_functions extension (llama/14244)" (llama/15910)

* Revert "sycl: add usage of enqueue_functions extension (#14244)"

This reverts commit 8308f98c7fb778e54bf75538f5234d8bd20915e9.

* fix missed revert code, format the code

2 months agoggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type (llama/15797)
Diego Devesa [Thu, 11 Sep 2025 20:47:38 +0000 (13:47 -0700)]
ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type (llama/15797)

* ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type

ggml-backend : add device id to device props

llama : only use iGPU devices if there are no GPU devices

llama : do not use multiple devices from different backends with the same device id

2 months agoCUDA: larger SRAM reads for tile FA, AMD FP16 dot (llama/15927)
Johannes Gäßler [Thu, 11 Sep 2025 19:19:58 +0000 (21:19 +0200)]
CUDA: larger SRAM reads for tile FA, AMD FP16 dot (llama/15927)

* CUDA: larger SRAM reads for tile FA, AMD FP16 dot

* fix logic for availability of v_dot2_f32_f16

2 months agoggml-cpu : add check for ARM MATMUL_INT8/i8mm support (llama/15922)
Daniel Bevenius [Thu, 11 Sep 2025 13:39:12 +0000 (15:39 +0200)]
ggml-cpu : add check for ARM MATMUL_INT8/i8mm support (llama/15922)

This commit adds a check for GGML_MACHINE_SUPPORTS_i8mm when enabling
MATMUL_INT8 features, ensuring that i8mm intrinsics are only used when
the target hardware actually supports them.

The motivation for this is to fix ggml CI build failures where the
feature detection correctly identifies that i8mm is not supported,
adding the +noi8mm flag, but MATMUL_INT8 preprocessor definitions are
still enabled, causing the compiler to attempt to use vmmlaq_s32
intrinsics without i8mm support.

Refs: https://github.com/ggml-org/ggml/actions/runs/17525174120/job/49909199499

2 months agokleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed (llama/15614)
Charles Xu [Thu, 11 Sep 2025 10:45:40 +0000 (12:45 +0200)]
kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed (llama/15614)

* kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed

* removes the Whisper-specific check for GET_ROWS support

2 months agoCANN: Disable acl_graph for prefill stage (llama/15933)
hipudding [Thu, 11 Sep 2025 07:59:37 +0000 (15:59 +0800)]
CANN: Disable acl_graph for prefill stage (llama/15933)

Since the prefill length is not fixed, graphs constructed for the
prefill stage cannot be reused. For this reason, ACL graph
execution is disabled by default during prefill.

2 months agoCUDA: Add `fastdiv` to `k_bin_bcast*`, giving 1-3% E2E performance (llama/15872)
Oliver Simons [Wed, 10 Sep 2025 20:04:03 +0000 (22:04 +0200)]
CUDA: Add `fastdiv` to `k_bin_bcast*`, giving 1-3% E2E performance (llama/15872)

* Add fastdiv and fastmodulo to k_bin_bcast kernel

* Address review comments

* `prod_` instead of `prod` suffix

* Add test case for `k_bin_bcast_unravel` in CUDA backend

2 months agoggml-cpu : fix padding in ggml_timestep_embedding (llama/15917)
Daniel Bevenius [Wed, 10 Sep 2025 15:31:40 +0000 (17:31 +0200)]
ggml-cpu : fix padding in ggml_timestep_embedding (llama/15917)

This commit fixes the zero padding for odd dimensions in
ggml_compute_forward_timestep_embedding_f32.
The motivation for this is that currently if an odd dimension is used,
the padding check incorrectly uses the dimension value for indexing.
For example, with dim=15:

Elements 0-6 are set to cosine values
Elements 7-13 are set to sine values
Element 14 is left uninitialized (contains garbage)
Element 15 is correctly set to zero

This fix changes embed_data[dim] to embed_data[2 * half] so that
element 14 (the first unused element) is properly set to zero as well
as the last element.

Resolves: https://github.com/ggml-org/ggml/issues/1324

2 months agosync : llama.cpp
Georgi Gerganov [Sat, 20 Sep 2025 10:07:06 +0000 (13:07 +0300)]
sync : llama.cpp

2 months agometal : make the backend async (llama/15906)
Georgi Gerganov [Sat, 20 Sep 2025 10:06:32 +0000 (13:06 +0300)]
metal : make the backend async (llama/15906)

2 months agosync : llama.cpp
Georgi Gerganov [Sat, 20 Sep 2025 10:06:06 +0000 (13:06 +0300)]
sync : llama.cpp

2 months agotests : filter out no-ops from coverage report (llama/15900)
Daniel Bevenius [Wed, 10 Sep 2025 12:17:09 +0000 (14:17 +0200)]
tests : filter out no-ops from coverage report (llama/15900)

* tests : filter out no-ops from coverage report

This commit is a follow-up commit for #15745 to address the feedback on
how no-op operations should be filtered out from the coverage report.

The feedback regarding the UNARY and GLU sub-operations not being
handled I not exactly sure what should be done. They are included in the
coverage, for example ABS, ELU, EXP, GELU, GEGLU, GEGLU_ERF etc are in
the list of covered operations:
```console
$ ./build/bin/test-backend-ops --show-coverage
Operations covered by tests (89):
  ✓ ABS
  ✓ ACC
  ✓ ADD
  ✓ ADD1
  ✓ ADD_ID
  ✓ ARANGE
  ✓ ARGMAX
  ✓ ARGSORT
  ✓ CLAMP
  ✓ CONCAT
  ✓ CONV_2D
  ✓ CONV_2D_DW
  ✓ CONV_3D
  ✓ CONV_TRANSPOSE_1D
  ✓ CONV_TRANSPOSE_2D
  ✓ COS
  ✓ COUNT_EQUAL
  ✓ CPY
  ✓ CROSS_ENTROPY_LOSS
  ✓ CROSS_ENTROPY_LOSS_BACK
  ✓ DIAG_MASK_INF
  ✓ DIV
  ✓ DUP
  ✓ ELU
  ✓ EXP
  ✓ FLASH_ATTN_EXT
  ✓ GATED_LINEAR_ATTN
  ✓ GEGLU
  ✓ GEGLU_ERF
  ✓ GEGLU_QUICK
  ✓ GELU
  ✓ GELU_ERF
  ✓ GELU_QUICK
  ✓ GET_ROWS
  ✓ GET_ROWS_BACK
  ✓ GROUP_NORM
  ✓ HARDSIGMOID
  ✓ HARDSWISH
  ✓ IM2COL
  ✓ IM2COL_3D
  ✓ L2_NORM
  ✓ LEAKY_RELU
  ✓ LOG
  ✓ MEAN
  ✓ MUL
  ✓ MUL_MAT
  ✓ MUL_MAT_ID
  ✓ NEG
  ✓ NORM
  ✓ OPT_STEP_ADAMW
  ✓ OPT_STEP_SGD
  ✓ OUT_PROD
  ✓ PAD
  ✓ PAD_REFLECT_1D
  ✓ POOL_2D
  ✓ REGLU
  ✓ RELU
  ✓ REPEAT
  ✓ REPEAT_BACK
  ✓ RMS_NORM
  ✓ RMS_NORM_BACK
  ✓ ROLL
  ✓ ROPE
  ✓ ROPE_BACK
  ✓ RWKV_WKV6
  ✓ RWKV_WKV7
  ✓ SCALE
  ✓ SET
  ✓ SET_ROWS
  ✓ SGN
  ✓ SIGMOID
  ✓ SILU
  ✓ SILU_BACK
  ✓ SIN
  ✓ SOFT_MAX
  ✓ SOFT_MAX_BACK
  ✓ SQR
  ✓ SQRT
  ✓ SSM_CONV
  ✓ SSM_SCAN
  ✓ STEP
  ✓ SUB
  ✓ SUM
  ✓ SUM_ROWS
  ✓ SWIGLU
  ✓ SWIGLU_OAI
  ✓ TANH
  ✓ TIMESTEP_EMBEDDING
  ✓ UPSCALE

Operations without tests (14):
  ✗ ADD_REL_POS
  ✗ CUSTOM
  ✗ DIAG
  ✗ DIAG_MASK_ZERO
  ✗ FLASH_ATTN_BACK
  ✗ GET_REL_POS
  ✗ IM2COL_BACK
  ✗ MAP_CUSTOM1
  ✗ MAP_CUSTOM2
  ✗ MAP_CUSTOM3
  ✗ POOL_1D
  ✗ POOL_2D_BACK
  ✗ WIN_PART
  ✗ WIN_UNPART

Coverage Summary:
  Total operations: 103
  Tested operations: 89
  Untested operations: 14
  Coverage: 86.4%
```

Refs: https://github.com/ggml-org/llama.cpp/pull/15745

* use of ggml_op enum values instead of strcmp

2 months agoCANN: Add ROPE sin/cos cache for reuse (llama/15912)
Chenguang Li [Wed, 10 Sep 2025 10:42:00 +0000 (18:42 +0800)]
CANN: Add ROPE sin/cos cache for reuse (llama/15912)

* CANN: Add ROPE sin/cos cache for reuse

Introduce sin/cos caching mechanism in ROPE to avoid redundant
computation across layers. The cache is built on the first layer
per device and reused by subsequent layers if parameters match.

- Added sin_cache / cos_cache pointers and position_length tracking
- Introduced cache validity flags and properties:
  (ext_factor, theta_scale, freq_scale, attn_factor, is_neox)
- Accelerates ROPE by eliminating repeated sin/cos generation

This change reduces overhead in multi-layer scenarios while
preserving correctness by verifying parameter consistency.

Co-authored-by: hipudding <redacted>
* fix typo

Signed-off-by: noemotiovon <redacted>
---------

Signed-off-by: noemotiovon <redacted>
Co-authored-by: hipudding <redacted>
2 months agoCANN: implement LRU cache for ACL graphs (llama/15814)
Chenguang Li [Wed, 10 Sep 2025 07:29:12 +0000 (15:29 +0800)]
CANN: implement LRU cache for ACL graphs (llama/15814)

* CANN: implement LRU cache for ACL graphs in CANN backend

- Introduce ggml_cann_graph_lru_cache to store multiple ggml_cann_graph objects.
- Graphs are loaded on demand and evicted using LRU policy when capacity is exceeded.
- Updated push, move_to_front, and clear methods to manage cached graphs efficiently.
- Ensures reuse of graphs, reducing graph reconstruction overhead in CANN backend.

* fix typo

* The LRU cache capacity can be configured via an env variable

Signed-off-by: noemotiovon <redacted>
* refactory acl graph

* refactory && fix review comments

Signed-off-by: noemotiovon <redacted>
---------

Signed-off-by: noemotiovon <redacted>
2 months agovulkan: throw the oom error instead of no memory type found (llama/15905)
Ruben Ortlam [Tue, 9 Sep 2025 20:26:03 +0000 (22:26 +0200)]
vulkan: throw the oom error instead of no memory type found (llama/15905)

2 months agovulkan: Fix OOB accesses in soft_max_back (llama/15861)
Jeff Bolz [Tue, 9 Sep 2025 12:41:15 +0000 (07:41 -0500)]
vulkan: Fix OOB accesses in soft_max_back (llama/15861)