]> git.djapps.eu Git - pkg/ggml/sources/whisper.cpp/log
pkg/ggml/sources/whisper.cpp
6 weeks agoCANN: Define `cann_graph_update_required` before macro (llama/17434)
Raul Torres [Mon, 24 Nov 2025 02:02:52 +0000 (02:02 +0000)]
CANN: Define `cann_graph_update_required` before macro (llama/17434)

**Description of the problem**

`cann_graph_update_required` is redundantly defined and
initialized as `false` inside two mutually exclusive macro branches.

**Proposed solution**

Define it right before the macro so that it could serve both
branches.

6 weeks agoggml-hexagon: Initial Hexagon v68/v69 support (llama/17394)
M. Mediouni [Mon, 24 Nov 2025 00:54:49 +0000 (01:54 +0100)]
ggml-hexagon: Initial Hexagon v68/v69 support (llama/17394)

* ggml-hexagon: fix build error with GCC

Add stdexcept include to fix GCC build errors

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: check VTCM acquire failures

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: disable destination bypass on older than v73

v68 errors out if having bypass enabled when the VTCM is the destination.

At least on v68 this made things actually work... not a proper fix though, so to look at later...

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: add initial v68/v69 support

v68 is the Hexagon revision notably used on the Snapdragon 8cx
Gen 3 and the QCM6490.

Also add support for v69.

8MB isn't a supported page size, so relax asked for page size constraint
for HAP_compute_res_attr_set_vtcm_param_v2 to optimal.

Signed-off-by: Mohamed Mediouni <redacted>
---------

Signed-off-by: Mohamed Mediouni <redacted>
6 weeks agoggml-hexagon: add `hex_supported_buffer` for better buffer supported check (llama...
nullname [Sun, 23 Nov 2025 22:26:36 +0000 (06:26 +0800)]
ggml-hexagon: add `hex_supported_buffer` for better buffer supported check (llama/17212)

* hexagon: add buffer support checks for hexagon sessions

* refactor: simplify buffer support checks in hexagon operations

* hexagon: update buffer support checks to use tensor structure

* refactor: streamline buffer initialization for DSP queue in hexagon operations

* refactor: simplify buffer initialization in DSP queue for hexagon operations

* refactor: optimize hex_supported_buffer function by fold expression

* wip

* refactor: simplify dspqueue_buffers_init function and its usage in hexagon operations

* fix: improve nan handling at hvx_vec_fast_sigmoid_fp32_guard

* refactor: optimize hvx_vec_inverse_fp32_guard for better nan handling

* refactor: update hvx_vec_fast_sigmoid_fp32_guard to use adjusted exponent limits

* refactor: modify hvx_vec_fast_sigmoid_fp32_guard to accept parameters for improved flexibility

* refactor: update hvx_vec_exp_fp32_guard to accept max_exp and inf parameters to save some instructions

* refactor: move hvx_vec_inverse_fp32_guard implementation to hvx-inverse.c for better perf

6 weeks agocuda : support non-contiguous i32 to i32 copy (llama/17326)
Sigbjørn Skjæret [Sun, 23 Nov 2025 10:13:34 +0000 (11:13 +0100)]
cuda : support non-contiguous i32 to i32 copy (llama/17326)

* support non-contiguous i32 to i32 copy

* add tests

* rename cpy_flt to cpy_scalar and reindent params

6 weeks agovulkan: remove a couple unnecessary switches (llama/17419)
Jeff Bolz [Sun, 23 Nov 2025 05:29:40 +0000 (23:29 -0600)]
vulkan: remove a couple unnecessary switches (llama/17419)

6 weeks agoHIP: RDNA4 tensor core support for MMF (llama/17077)
yulo [Fri, 21 Nov 2025 23:03:24 +0000 (07:03 +0800)]
HIP: RDNA4 tensor core support for MMF (llama/17077)

* mmf for rdna4

* align the padding for rdna4

* forbit mul_mat_f for rdna4

* fix as comment

* remove device kernels

* add constexpr for early return

* update based on review comment

* change based on the review comment

* pass compile error

* keep code consistency

---------

Co-authored-by: zhang hui <redacted>
6 weeks agoopencl: refine condition for kqv mm (llama/17392)
lhez [Fri, 21 Nov 2025 22:34:48 +0000 (14:34 -0800)]
opencl: refine condition for kqv mm (llama/17392)

6 weeks agovulkan: disable async for older Intel devices (llama/17369)
Jeff Bolz [Fri, 21 Nov 2025 08:58:17 +0000 (02:58 -0600)]
vulkan: disable async for older Intel devices (llama/17369)

* vulkan: disable async for older Intel devices

* update detection logic

* use name string for detection

6 weeks agoCANN: Refactor `evaluate_and_capture_cann_graph` (llama/17333)
Raul Torres [Fri, 21 Nov 2025 08:23:29 +0000 (08:23 +0000)]
CANN: Refactor `evaluate_and_capture_cann_graph` (llama/17333)

* CANN: Refactor `evaluate_and_capture_cann_graph`

**Description of the problem**

* `matched_graph` is obtained even if graph mode is disabled.
* End of graph capture and graph replay are unnecessarily placed in different `if` blocks.

**Proposed solution**

* Obtain `matched_graph` only if graph mode is enabled.
* Place end of graph capture and graph reply inside the same `if` block.
* Unify graph related comments.

* Remove trailing whitespace

6 weeks agoggml-hexagon: fix swiglu failure at `test-backend-ops` (llama/17344)
nullname [Thu, 20 Nov 2025 23:45:05 +0000 (07:45 +0800)]
ggml-hexagon: fix swiglu failure at `test-backend-ops` (llama/17344)

* refactor: use hvx_vec_exp_fp32_guard_inf for overflow handling in hvx_exp_f32

* feat: add fast sigmoid function with overflow guard for fp32

* refactor: replace hvx_vec_inverse_fp32 with hvx_vec_inverse_fp32_guard_inf for improved overflow handling

* feat: enhance hvx_add_scalar_f32 with overflow handling using infinity guard

* wip

* add HVX_Vector_Alias

wip

* wip

* fix: improve handling of src1 tensor in glu_swiglu_fp32_per_thread function

* fix nc

* wip

* wip

* handle nan at inverse

* wip

* fix neg

* wip

* rename

* fix hvx_vec_inverse_fp32_guard_inf to handle infinity and NaN cases correctly

* wip

* fix hvx_vec_inverse_fp32_guard_inf to handle NaN cases correctly

* wip

* wip

* wip

* fix output sign

6 weeks agoggml : Fix transposed SOLVE_TRI result (llama/17323)
Piotr Wilkin (ilintar) [Thu, 20 Nov 2025 10:58:21 +0000 (11:58 +0100)]
ggml : Fix transposed SOLVE_TRI result (llama/17323)

* Did someone transpose the SOLVE_TRI result matrix? Perhaps...

* Update ggml/src/ggml-cpu/ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update ggml/src/ggml-cpu/ops.cpp

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

Co-authored-by: Sigbjørn Skjæret <redacted>
6 weeks agoDGX Spark: UMA support (llama/17368)
Scott Fudally [Thu, 20 Nov 2025 10:32:02 +0000 (02:32 -0800)]
DGX Spark: UMA support (llama/17368)

* DGX Spark: UMA support

* Updates from PR feedback

* More PR feedback cleanup

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

Co-authored-by: Georgi Gerganov <redacted>
* Remove trailing whitespace

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

---------

Co-authored-by: Georgi Gerganov <redacted>
6 weeks agoggml : remove useless and error-prone variadic macros (llama/17399)
Adrien Gallouët [Thu, 20 Nov 2025 10:18:27 +0000 (11:18 +0100)]
ggml : remove useless and error-prone variadic macros (llama/17399)

Signed-off-by: Adrien Gallouët <redacted>
6 weeks agokleidiai: fix zero-size array declaration (llama/17240)
sudhiarm [Thu, 20 Nov 2025 09:45:49 +0000 (09:45 +0000)]
kleidiai: fix zero-size array declaration (llama/17240)

6 weeks agoggml-cpu:add RISC-V RVV (Zvfh) optimization for FP16 vector scaling (llama/17314)
ixgbe [Thu, 20 Nov 2025 06:09:18 +0000 (14:09 +0800)]
ggml-cpu:add RISC-V RVV (Zvfh) optimization for FP16 vector scaling (llama/17314)

* ggml-cpu:add RISC-V RVV (Zvfh) optimization for FP16 vector scaling

Signed-off-by: Wang Yang <redacted>
* fix comment

* fix comment 2

---------

Signed-off-by: Wang Yang <redacted>
6 weeks agovulkan: implement ADD1, ARANGE, FILL, SOFTPLUS, STEP, ROUND, CEIL, FLOOR, TRUNC ...
Giuseppe Scrivano [Wed, 19 Nov 2025 16:29:45 +0000 (17:29 +0100)]
vulkan: implement ADD1, ARANGE, FILL, SOFTPLUS, STEP, ROUND, CEIL, FLOOR, TRUNC (llama/17319)

* vulkan: initialize array

* vulkan: implement ADD1

* vulkan: implement ARANGE

* vulkan: implement FILL

* vulkan: implement SOFTPLUS

* vulkan: implement STEP

* vulkan: implement ROUND

* vulkan: implement CEIL

* vulkan: implement FLOOR

* vulkan: implement TRUNC

* docs: update Vulkan ops

Signed-off-by: Giuseppe Scrivano <redacted>
6 weeks agovulkan: support larger argsort (llama/17313)
Jeff Bolz [Wed, 19 Nov 2025 16:25:50 +0000 (10:25 -0600)]
vulkan: support larger argsort (llama/17313)

* vulkan: support larger argsort

This is an extension of the original bitonic sorting shader that puts the
temporary values in global memory and when more than 1024 threads are needed
it runs multiple workgroups and synchronizes through a pipelinebarrier.

To improve the memory access pattern, a copy of the float value is kept with
the index value. I've applied this same change to the original shared memory
version of the shader, which is still used when ncols <= 1024.

* Reduce the number of shader variants. Use smaller workgroups when doing a single pass, for a modest perf boost

* reduce loop overhead

* run multiple cols per invocation, to reduce barrier overhead

6 weeks agovulkan: Add copy_transpose shader (llama/17371)
Jeff Bolz [Wed, 19 Nov 2025 15:50:43 +0000 (09:50 -0600)]
vulkan: Add copy_transpose shader (llama/17371)

6 weeks agocuda: fix rope fusion for gemma3 (llama/17378)
Aman Gupta [Wed, 19 Nov 2025 10:25:05 +0000 (18:25 +0800)]
cuda: fix rope fusion for gemma3 (llama/17378)

6 weeks agoFix too relaxed check on CUDA "fast copy" (can_be_transposed) condition (llama/17332)
Piotr Wilkin (ilintar) [Wed, 19 Nov 2025 09:36:33 +0000 (10:36 +0100)]
Fix too relaxed check on CUDA "fast copy" (can_be_transposed) condition (llama/17332)

* Fix too relaxed check on CUDA "fast copy" (can_be_transposed) condition

* Argh.

* Making CISC happy ;)

* Integrate CONT tests

* Use loopy loop

* Skip new tests for (B)F16 for now.

6 weeks agovulkan: force full subgroups for flash attention to fix intel subgroup crash (llama...
Ruben Ortlam [Wed, 19 Nov 2025 07:46:26 +0000 (08:46 +0100)]
vulkan: force full subgroups for flash attention to fix intel subgroup crash (llama/17356)

6 weeks agoggml-cpu: Don't pass -mpowerpc64 when -mcpu already implies it (llama/17308)
Jeremy Rand [Wed, 19 Nov 2025 06:19:00 +0000 (06:19 +0000)]
ggml-cpu: Don't pass -mpowerpc64 when -mcpu already implies it (llama/17308)

6 weeks agoCANN: fix acl_tensor_ptr usage in ASCEND_310P ROPE (llama/17347)
Chenguang Li [Tue, 18 Nov 2025 08:41:52 +0000 (16:41 +0800)]
CANN: fix acl_tensor_ptr usage in ASCEND_310P ROPE (llama/17347)

* cann: fix acl_tensor_ptr usage in ASCEND_310P ROPE implementation

Fix compilation errors in the ASCEND_310P-specific ROPE operation code
by adding .get() calls when passing acl_tensor_ptr smart pointers to
functions expecting raw aclTensor* pointers.

This fixes the code that was missed in the previous refactoring commit
(8981848) which changed ggml_cann_create_tensor() return type from
aclTensor* to acl_tensor_ptr.

* cann: format code

6 weeks agovulkan: support noncontig i32 copy (llama/17328)
Jeff Bolz [Tue, 18 Nov 2025 06:41:24 +0000 (00:41 -0600)]
vulkan: support noncontig i32 copy (llama/17328)

6 weeks agovulkan: add log RTE support to fix Nvidia CI (llama/17320)
Ruben Ortlam [Mon, 17 Nov 2025 20:37:49 +0000 (21:37 +0100)]
vulkan: add log RTE support to fix Nvidia CI (llama/17320)

* vulkan: add log RTE support to fix Nvidia CI

* actually use the rte shader

6 weeks agocmake : fix ARM feature verification (llama/17170)
Adrien Gallouët [Mon, 17 Nov 2025 20:37:29 +0000 (21:37 +0100)]
cmake : fix ARM feature verification (llama/17170)

* cmake : fix ARM feature verification

Use check_cxx_source_compiles to prevent conflicts with
the existing GGML_NATIVE detection code.

Signed-off-by: Adrien Gallouët <redacted>
* cmake : unset __ARM_FEATURE when feature is disabled

Signed-off-by: Adrien Gallouët <redacted>
* cmake : fix scope, this is really a macro

Signed-off-by: Adrien Gallouët <redacted>
* arm_neon.h is useless

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

Signed-off-by: Adrien Gallouët <redacted>
6 weeks agoggml : add missing AVX512 feature checks (llama/17270)
Adrien Gallouët [Mon, 17 Nov 2025 11:12:00 +0000 (12:12 +0100)]
ggml : add missing AVX512 feature checks (llama/17270)

_mm512_cvtepu8_epi16        requires  __AVX512BW__
_mm512_srli_epi16           requires  __AVX512BW__
__builtin_ia32_inserti32x8  requires  __AVX512DQ__

Signed-off-by: Adrien Gallouët <redacted>
6 weeks agoggml : remove dirty flag from version string (ggml/1391)
Daniel Bevenius [Mon, 24 Nov 2025 11:51:50 +0000 (12:51 +0100)]
ggml : remove dirty flag from version string (ggml/1391)

This commit removes the "-dirty" suffix from the GGML version string.

The motivation for this change is to ensure that the version string
works with different ways of checking out ggml and using it in projects.
By removing the dirty flag from the version string, we avoid potential
artifacts like shared libraries getting a -dirty suffix in their names.

Instead, if the project is built from a dirty git state, the dirty flag
will be appended to the commit hash in the GGML_BUILD_COMMIT variable.
This will enable users to still identify that the build was made from
from a modified/dirty state even though the version might match a "real"
version.

For example, the commit can be produces as follows:
```c++
    printf("commit: %s\n", ggml_commit());
```
Which would print the following for a dirty build:
```console
commit: 781baf2a-dirty
```

Refs: https://github.com/ggml-org/ggml/pull/1363#issuecomment-3569691546

6 weeks agogo : Enable VAD for Go bindings (#3563)
Josh Montoya [Wed, 10 Dec 2025 11:31:36 +0000 (03:31 -0800)]
go : Enable VAD for Go bindings (#3563)

* reset context.n so that NextSegment can be called for multiple Process calls

* enable VAD params

6 weeks agogo : reset context.n in Process() (#3503)
Josh Montoya [Mon, 8 Dec 2025 16:33:07 +0000 (08:33 -0800)]
go : reset context.n in Process() (#3503)

7 weeks agovad : fix buffer overflow in sample reduction loop (#3558)
Joseph Sellers [Sat, 6 Dec 2025 11:28:32 +0000 (11:28 +0000)]
vad : fix buffer overflow in sample reduction loop (#3558)

The buffer size calculation loop (line ~6661) uses `n_samples - 1` as
the upper bound for segment_end_samples, but the copy loop (line 6696)
uses `n_samples`. This inconsistency allows the copy loop to compute
segment_length values up to 1 sample larger per segment than what was
allocated, causing heap corruption.

Symptom: `malloc(): corrupted top size` or `malloc(): invalid size
(unsorted)` crashes after VAD completes sample reduction.

Fix: Use consistent bounds (`n_samples - 1`) in both loops.

Fixes #3403

7 weeks agotests : update VAD tests to use Silero V6.2.0 (#3534)
Daniel Bevenius [Sat, 6 Dec 2025 09:58:58 +0000 (10:58 +0100)]
tests : update VAD tests to use Silero V6.2.0 (#3534)

* tests : update VAD tests to use Silero V6.2.0

This commit updates the VAD tests to use the Silero V6.2.0 instead of
V5.1.2. I'm was not sure if we needed to keep testing for both versions,
but opted to just update to the latest version for simplicity.

* wasm : use C++17 for emscripten builds

This commit updates the CMakeLists.txt file to explicitly set the C++
standard to C++17 when building with Emscripten.

The motivation for this change is that building with Emscripten
will currently fail locally and on CI with the following error:
```console
[ 75%] Building CXX object examples/CMakeFiles/common.dir/common-ggml.cpp.o
In file included from /home/danbev/work/ai/whisper.cpp/examples/stream.wasm/emscripten.cpp:5:
/home/danbev/work/utils/emsdk/upstream/emscripten/cache/sysroot/include/emscripten/bind.h:11:2: error:
      "embind requires -std=c++17 or newer"
   11 | #error "embind requires -std=c++17 or newer"
      |  ^
In file included from /home/danbev/work/ai/whisper.cpp/examples/whisper.wasm/emscripten.cpp:4:
/home/danbev/work/utils/emsdk/upstream/emscripten/cache/sysroot/include/emscripten/bind.h:11:2: error:
      "embind requires -std=c++17 or newer"
   11 | #error "embind requires -std=c++17 or newer"
      |  ^
```

2 months agoexamples : fix typo in vad-speech-segments command [no ci] (#3535)
Daniel Bevenius [Thu, 20 Nov 2025 12:35:11 +0000 (13:35 +0100)]
examples : fix typo in vad-speech-segments command [no ci] (#3535)

This commit corrects a typo the command-line argument for specifying the
VAD model in the vad-speech-segments example.

2 months agoreadme : minor (#3516)
gzq [Thu, 20 Nov 2025 11:57:55 +0000 (19:57 +0800)]
readme : minor (#3516)

2 months agometal : fix compile on macos 11 (#3533)
YangLe [Thu, 20 Nov 2025 11:54:54 +0000 (19:54 +0800)]
metal : fix compile on macos 11 (#3533)

2 months agosync : llama.cpp
Georgi Gerganov [Mon, 17 Nov 2025 14:31:08 +0000 (16:31 +0200)]
sync : llama.cpp

2 months agosync : ggml
Georgi Gerganov [Mon, 17 Nov 2025 14:26:39 +0000 (16:26 +0200)]
sync : ggml

2 months agometal : support I32 -> I32 copy (llama/17317)
Georgi Gerganov [Mon, 17 Nov 2025 09:52:00 +0000 (11:52 +0200)]
metal : support I32 -> I32 copy (llama/17317)

2 months agometal : faster argsort (llama/17315)
Georgi Gerganov [Mon, 17 Nov 2025 09:51:48 +0000 (11:51 +0200)]
metal : faster argsort (llama/17315)

* metal : faster argsort

* cont : keep data in registers

2 months agometal : add cumsum (llama/17305)
Georgi Gerganov [Mon, 17 Nov 2025 09:51:13 +0000 (11:51 +0200)]
metal : add cumsum (llama/17305)

2 months agoCANN: Use smart pointers to manage ACL objects (llama/17238)
hipudding [Mon, 17 Nov 2025 00:43:59 +0000 (08:43 +0800)]
CANN: Use smart pointers to manage ACL objects (llama/17238)

* CANN: Use smart pointers to manage ACL objects

Previously, ACL objects were managed via manual destruction, which
led to multiple memory-leak issues during runtime. This patch replaces
manual memory management with smart pointers so that ACL objects
are properly released and ownership is clearly defined.

Note that the ownership of an ACL object belongs to the function
that creates it. Other internal functions should operate on these ACL
objects using raw pointers to avoid unintended ownership transfers.

Additionally, since aclTensorList automatically frees its contained
aclTensor objects, any aclTensor added to a tensor list must release
ownership to avoid double free operations.

This PR also removes the asynchronous task submission mechanism.
Due to changes in recent CANN versions, tiling time has significantly
decreased. Even with a dual-thread submission model, the dispatch
overhead still falls on the critical path, making async submission
less beneficial. Moreover, aclGraph support provides a much better
path to reducing operator dispatch latency.

* CANN: resolve review comments

2 months agovulkan: add LOG operation support for F32 and F16 (llama/17183)
Pavels Zaicenkovs [Sun, 16 Nov 2025 21:50:09 +0000 (22:50 +0100)]
vulkan: add LOG operation support for F32 and F16 (llama/17183)

* vulkan: add LOG operation support for F32 and F16

Part of #14909.

* vulkan: Fix LOG operation types

* docs: Update operation support documentation for Vulkan LOG operation

* vulkan: fix log_f16 shader

* docs: restore missing LOG test cases and regenerate ops.md

2 months agovulkan: fix MMQ quantize_y condition (llama/17301)
Ruben Ortlam [Sun, 16 Nov 2025 18:38:17 +0000 (19:38 +0100)]
vulkan: fix MMQ quantize_y condition (llama/17301)

2 months agometal : remove obosolete asserts (llama/17295)
Georgi Gerganov [Sun, 16 Nov 2025 07:50:26 +0000 (09:50 +0200)]
metal : remove obosolete asserts (llama/17295)

2 months agoopencl: fix rms_norm_mul (llama/17250)
lhez [Sun, 16 Nov 2025 01:40:14 +0000 (17:40 -0800)]
opencl: fix rms_norm_mul (llama/17250)

* opencl: use subgrroup reduce for reduction in rms_norm_mul

* opencl: add comment about workgroup size

2 months agoopencl: add kernel to handle mat mul in attention to improve encoding speed (llama...
shaofeiqi [Sun, 16 Nov 2025 01:33:10 +0000 (17:33 -0800)]
opencl: add kernel to handle mat mul in attention to improve encoding speed (llama/17181)

* Add mul_mm_f16_f32_kq_kqv kernel

* Add ggml_cl_mul_mat_kq_kqv_adreno func

* fix whitespace

* remove unused variable

* remove redundant

* refactor and clean up

* remove trailing whitespace

2 months agosycl : unify unary kernels with a generic implementation and enable wide operator...
shani-f [Sat, 15 Nov 2025 23:52:42 +0000 (01:52 +0200)]
sycl : unify unary kernels with a generic implementation and enable wide operator support (llama/17213)

* SYCL: add generic unary op implementation for multiple ops (ABS/SGN/…); unify non-contiguous access

* SYCL: update documentation and sycl.csv to reflect new unary op support

* update ops.md after syncing SYCL.csv changes

* Fix SYCL.csv merge conflict

* Update ops.md after fixing SYCL.csv conflicts

* Fix SYCL.csv tail after merge conflict and regenerate ops.md

* Fix line endings and final newline in SYCL.csv

* Remove TOPK_MOE entries from SYCL.csv as requested

* Update ops.md after removing TOPK_MOE from SYCL.csv

* Regenerated SYCL.csv and synced ops.md with upstream

* Update ops.md using create_ops_docs.py

2 months agovulkan: Fuse mul_mat_id+add_id+mul and mul_mat+add+add. (llama/17287)
Jeff Bolz [Sat, 15 Nov 2025 18:54:23 +0000 (12:54 -0600)]
vulkan: Fuse mul_mat_id+add_id+mul and mul_mat+add+add. (llama/17287)

These both show up in gpt-oss. Also, cleanup the mul_mat_vec fusion code a bit.

2 months agovulkan: Replace 16-bit unpack8 calls to work around legacy Windows AMD driver bug...
Ruben Ortlam [Sat, 15 Nov 2025 14:18:58 +0000 (15:18 +0100)]
vulkan: Replace 16-bit unpack8 calls to work around legacy Windows AMD driver bug (llama/17285)

2 months agovulkan: implement ABS and NEG (llama/17245)
Giuseppe Scrivano [Sat, 15 Nov 2025 11:00:29 +0000 (12:00 +0100)]
vulkan: implement ABS and NEG (llama/17245)

* docs: update Vulkan ops

* vulkan: add NEG op

* vulkan: add ABS op

---------

Signed-off-by: Giuseppe Scrivano <redacted>
2 months agovulkan: Use ggml_vk_tensor_subbuffer in mul_mat_vec(id) paths (llama/17244)
Jeff Bolz [Sat, 15 Nov 2025 10:56:15 +0000 (04:56 -0600)]
vulkan: Use ggml_vk_tensor_subbuffer in mul_mat_vec(id) paths (llama/17244)

* vulkan: Use ggml_vk_tensor_subbuffer in mul_mat_vec(id) paths

* set allow_misalign

2 months agovulkan: skip all-negative-inf blocks in FA (llama/17186)
Jeff Bolz [Sat, 15 Nov 2025 09:37:25 +0000 (03:37 -0600)]
vulkan: skip all-negative-inf blocks in FA (llama/17186)

2 months agovulkan: change graph_compute to be async and enable get_tensor_async (llama/17158)
Jeff Bolz [Sat, 15 Nov 2025 08:06:41 +0000 (02:06 -0600)]
vulkan: change graph_compute to be async and enable get_tensor_async (llama/17158)

* vulkan: change graph_compute to be async and enable get_tensor_async

This allows some additional CPU/GPU overlap for large pp workloads. Also seems
to help a bit for token gen, maybe getting rid of a small bubble between
graph_compute and get_tensor.

Async set and copy functions seem to be very rarely used, so I didn't enable
them because I didn't have a good way to test them.

The async commands need to be ordered against each other, so put them all on
the compute queue. The non-async commands still use the transfer queue.

The fence for graph_compute/get_tensor_async is submitted and waited on in
ggml_vk_synchronize.

* fix thread safety errors

* teardown context cleanly

* Handle async read to non-pinned dst

2 months agometal : support argsort for ne00 > 1024 (llama/17247)
Georgi Gerganov [Fri, 14 Nov 2025 07:36:06 +0000 (09:36 +0200)]
metal : support argsort for ne00 > 1024 (llama/17247)

* metal : refactor argsort

* cont : sort chunks

* cont : merge sorted buckets

* cont : cleanup

2 months agometal : make the FA extra sizes consistent (llama/17143)
Georgi Gerganov [Fri, 14 Nov 2025 07:13:34 +0000 (09:13 +0200)]
metal : make the FA extra sizes consistent (llama/17143)

2 months agoggml-cpu: handle 3d tensors in repack mat_mul (llama/17241)
Alberto Cabrera Pérez [Thu, 13 Nov 2025 20:53:00 +0000 (20:53 +0000)]
ggml-cpu: handle 3d tensors in repack mat_mul (llama/17241)

* ggml-cpu: handle 3d tensors in repack mul_mat

* Removed unnecessary branch, removed need for <algorithm>

* Fixed dst_ptr pointer in chunk + clang_format

* GGML_ASSERT to check wdata within bounds

* Accidental ggml.h inclusion

* Improved GGML_ASSERT on wdata boundaries

* Address performance regression in Qwen and llama.cpp due to chunking

2 months agoggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063)
Piotr Wilkin (ilintar) [Thu, 13 Nov 2025 18:54:47 +0000 (19:54 +0100)]
ggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063)

* Add ops needed for new hybrid models: SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM

* Update ggml/include/ggml.h

Co-authored-by: Georgi Gerganov <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Code review

* Whitespace

* Update tests/test-backend-ops.cpp

Co-authored-by: Diego Devesa <redacted>
* This is actually sigmoid, duh.

* Add CONST, remove TRI_KEEP, other changes from review

* Update tests/test-backend-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml.c

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml.c

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml-cuda/unary.cu

Co-authored-by: Aman Gupta <redacted>
* Remove extra script

* Update ggml/src/ggml.c

Co-authored-by: Diego Devesa <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Diego Devesa <redacted>
* moving changes from laptop [no ci]

* pre-rebase

* 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>
* Refactor tests

* ggml : cleanup

* cont : fix ggml_fill srcs

* tests : add note

* ggml : add ggml_fill_inplace

* ggml : add asserts

* ggml : fix ggml_fill constant cast

* cont : ggml_tri minor

* Use TENSOR_LOCALS

* Fix regression from #14596, regenerate

* Don't make commits at night...

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Diego Devesa <redacted>
Co-authored-by: Aman Gupta <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agovulkan: remove shell call from vulkan-shaders-gen tool, revert file check (llama...
Ruben Ortlam [Thu, 13 Nov 2025 13:51:21 +0000 (14:51 +0100)]
vulkan: remove shell call from vulkan-shaders-gen tool, revert file check (llama/17219)

* vulkan: remove shell call from vulkan-shaders-gen tool

* use string vector for command execution

* Fix condition

* use string, remove const_cast

* Fix dependency file quotation on Windows

---------

Co-authored-by: Jeff Bolz <redacted>
2 months agosched : fix reserve ignoring user tensor assignments (llama/17232)
Diego Devesa [Thu, 13 Nov 2025 12:14:02 +0000 (04:14 -0800)]
sched : fix reserve ignoring user tensor assignments (llama/17232)

2 months agoggml-cpu : add RISC-V vector intrinsic support for silu and cvar operations (llama...
ixgbe [Thu, 13 Nov 2025 12:13:32 +0000 (20:13 +0800)]
ggml-cpu : add RISC-V vector intrinsic support for silu and cvar operations (llama/17227)

Signed-off-by: Wang Yang <redacted>
2 months agometal: accelerated conv2d (llama/17175)
bagheera [Thu, 13 Nov 2025 11:32:44 +0000 (05:32 -0600)]
metal: accelerated conv2d (llama/17175)

* metal: accelerated conv2d

* cont : cleanup

---------

Co-authored-by: bghira <redacted>
Co-authored-by: Georgi Gerganov <redacted>
2 months agoRevert "ggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)" (llama/17233)
Georgi Gerganov [Thu, 13 Nov 2025 10:59:37 +0000 (12:59 +0200)]
Revert "ggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)" (llama/17233)

This reverts commit 1c398dc9eca9c366ce98deb0e6f3538e444ebc8a.

2 months agoggml-cpu : use template for argsort (llama/17222)
Diego Devesa [Thu, 13 Nov 2025 08:59:05 +0000 (00:59 -0800)]
ggml-cpu : use template for argsort (llama/17222)

2 months agoCANN: Add cross_entropy_loss op support (llama/16886)
TecJesh [Thu, 13 Nov 2025 01:39:51 +0000 (09:39 +0800)]
CANN: Add cross_entropy_loss op support (llama/16886)

* update L2_NORM op support

* update L2_NORM op support

* remove extra whitespace

* cann: update cross_entropy_loss op support

* remove trailing whitespaces

* rebase the latest code in the main repository and remove the l2_norm operator that already exists in another pull request.

* undo the l2_norm operator deletion

2 months agoCUDA: fuse rope + set_rows (llama/16884)
Aman Gupta [Thu, 13 Nov 2025 00:50:01 +0000 (08:50 +0800)]
CUDA: fuse rope + set_rows (llama/16884)

* CUDA: add fused rope

* move k forward_expand up

* create helper function instead of re-using params

* make assert statement more in line with comment

* rope_norm: coalesced writes to global mem

2 months agoCUDA: static assert to prevent misuse of memcpy_1 (llama/17198)
Johannes Gäßler [Wed, 12 Nov 2025 22:13:55 +0000 (23:13 +0100)]
CUDA: static assert to prevent misuse of memcpy_1 (llama/17198)

2 months agoggml : use std::sort in ggml_argsort CPU implementation (llama/17211)
Georgi Gerganov [Wed, 12 Nov 2025 18:43:38 +0000 (20:43 +0200)]
ggml : use std::sort in ggml_argsort CPU implementation (llama/17211)

* ggml : use std::sort in ggml_argsort CPU implementation

* cont : add missing header

2 months agoggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)
Alberto Cabrera Pérez [Wed, 12 Nov 2025 12:52:19 +0000 (12:52 +0000)]
ggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)

* ggml-cpu: handle 3d tensors in repack mul_mat

* Removed unnecessary branch, removed need for <algorithm>

* Fixed dst_ptr pointer in chunk + clang_format

* GGML_ASSERT to check wdata within bounds

* Accidental ggml.h inclusion

* Improved GGML_ASSERT on wdata boundaries

2 months agoCANN: Add L2_NORM op support (llama/16856)
TecJesh [Wed, 12 Nov 2025 07:11:42 +0000 (15:11 +0800)]
CANN: Add L2_NORM op support (llama/16856)

* update L2_NORM op support

* update L2_NORM op support

* remove extra whitespace

2 months agofix ci crash about SSM_CONV (llama/17169)
Neo Zhang Jianyu [Wed, 12 Nov 2025 06:44:29 +0000 (14:44 +0800)]
fix ci crash about SSM_CONV (llama/17169)

* fix ci crash

* Update ggml-sycl.cpp

* Update ggml/src/ggml-sycl/ggml-sycl.cpp

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

Co-authored-by: Zhang Jianyu <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agohexagon: various Op fixes (llama/17135)
Max Krasnyansky [Tue, 11 Nov 2025 23:25:04 +0000 (15:25 -0800)]
hexagon: various Op fixes (llama/17135)

* hexagon: explicitly check for ops with zero nrows

llm_graph_context::build_inp_out_ids() can generate tensors with zero nrows.
Somehow other backends seems to handle this without obvious explicit checks.
In the hexagon case we need to check explicitly and skip them.

* hexagon: introduce fastdiv, fix test-backend-ops for ADD/SUB/MUL

Co-authored-by: chraac <redacted>
* hexagon: use fastdiv in ADD_ID

* hexagon: use ggml_op_is_empty and ggml_is_empty to check for NOPs

---------

Co-authored-by: chraac <redacted>
2 months agodisable rms norm mul rope for chips with no fp16 rte (llama/17134)
Eve [Tue, 11 Nov 2025 18:53:30 +0000 (18:53 +0000)]
disable rms norm mul rope for chips with no fp16 rte (llama/17134)

2 months agoggml-cpu : add RISC-V RVV (Zvfh) optimization for FP16 to FP32 conversion (llama...
ixgbe [Tue, 11 Nov 2025 11:41:51 +0000 (19:41 +0800)]
ggml-cpu : add RISC-V RVV (Zvfh) optimization for FP16 to FP32 conversion (llama/17161)

Signed-off-by: Wang Yang <redacted>
2 months agoggml-cpu: templateify ggml_compute_forward_rope_f32 and _f16 (llama/16805)
duduta [Tue, 11 Nov 2025 11:33:24 +0000 (13:33 +0200)]
ggml-cpu: templateify ggml_compute_forward_rope_f32 and _f16 (llama/16805)

* extract rotate_pairs logic from ggml_compute_forward_rope_f32

* templateify ggml_compute_forward_rope_f32 and _f16

* abort when rope type not supported, remove GLM from test-rope

* add imrope branch to switch

* add rope tests for perf

* Update ggml/src/ggml-cpu/ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml-cpu/ops.cpp

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

Co-authored-by: Georgi Gerganov <redacted>
2 months agokleidiai: add optimized per-channel kernels for Q8_0 (llama/16993)
Charles Xu [Tue, 11 Nov 2025 11:20:31 +0000 (12:20 +0100)]
kleidiai: add optimized per-channel kernels for Q8_0 (llama/16993)

2 months agocmake : add version to all shared object files (llama/17091)
Mike Abbott [Tue, 11 Nov 2025 11:19:50 +0000 (04:19 -0700)]
cmake : add version to all shared object files (llama/17091)

When compiling llama.cpp in Yocto, it fails QA checks because the generated so files aren't versioned.  This applies a version to all generated so files, allowing the package to build without errors.

2 months agoopencl: add fastdiv and use it in set_rows, ported from cuda (llama/17090)
lhez [Mon, 10 Nov 2025 23:00:13 +0000 (15:00 -0800)]
opencl: add fastdiv and use it in set_rows, ported from cuda (llama/17090)

* opencl: add fastdiv for mm q8_0

* opencl: use uint4 for fastdiv vals

* opencl: use fastdiv for set_rows

* opencl: do not use fastdiv for q8_0 mm

2 months agocpu: skip NOPs to avoid barriers (llama/17133)
Max Krasnyansky [Mon, 10 Nov 2025 20:44:49 +0000 (12:44 -0800)]
cpu: skip NOPs to avoid barriers (llama/17133)

* cpu: skip NOPs to avoid barriers

* cpu: use ggml_op_is_empty

2 months agometal : cap threadgroups size of set_rows (llama/17146)
Georgi Gerganov [Mon, 10 Nov 2025 19:33:35 +0000 (21:33 +0200)]
metal : cap threadgroups size of set_rows (llama/17146)

2 months agoggml-cpu : inspect -march and -mcpu to found the CPU (llama/16333)
Adrien Gallouët [Mon, 10 Nov 2025 19:03:36 +0000 (20:03 +0100)]
ggml-cpu : inspect -march and -mcpu to found the CPU (llama/16333)

Signed-off-by: Adrien Gallouët <redacted>
2 months agovulkan: check glslc executable string (llama/17144)
Ruben Ortlam [Mon, 10 Nov 2025 15:59:26 +0000 (16:59 +0100)]
vulkan: check glslc executable string (llama/17144)

2 months agovulkan: fix validation issue introduced by #16868 (llama/17145)
Ruben Ortlam [Mon, 10 Nov 2025 15:59:10 +0000 (16:59 +0100)]
vulkan: fix validation issue introduced by #16868 (llama/17145)

2 months agometal : enable tensor API for A19 (llama/17087)
Georgi Gerganov [Mon, 10 Nov 2025 13:38:42 +0000 (15:38 +0200)]
metal : enable tensor API for A19 (llama/17087)

2 months agoarm64: add i8mm route with SVE ggml_vec_dot_q4_K_q8_K and ggml_vec_dot_q6_K_… (#15277)
fj-y-saito [Mon, 10 Nov 2025 13:12:59 +0000 (22:12 +0900)]
arm64: add i8mm route with SVE ggml_vec_dot_q4_K_q8_K and ggml_vec_dot_q6_K_… (#15277)

* add i8mm route with SVE ggml_vec_dot_q4_K_q8_K and ggml_vec_dot_q6_K_q8_K

* Surround SVE function with compiler directive

* fix compile switch

* fix coding style

* ggml : fix indent

---------

Co-authored-by: Georgi Gerganov <redacted>
2 months agocuda/vulkan : bicubic interpolation (llama/17022)
Acly [Mon, 10 Nov 2025 09:19:39 +0000 (10:19 +0100)]
cuda/vulkan : bicubic interpolation (llama/17022)

* vulkan : implement upscale with bicubic interpolation

* cuda : implement upscale with bicubic interpolation

* tests : add ggml_interpolate with GGML_SCALE_MODE_BICUBIC to backend tests

* adapt OpenCL backend to not support the OP in that case so tests don't fail

* print scale mode & flags in test-backend-ops

2 months agovulkan: fix memory allocations (llama/17122)
Ruben Ortlam [Sun, 9 Nov 2025 15:14:41 +0000 (16:14 +0100)]
vulkan: fix memory allocations (llama/17122)

2 months agovad : Silero VAD v6.2.0 (#3524)
KITAITI Makoto [Mon, 17 Nov 2025 13:26:17 +0000 (22:26 +0900)]
vad : Silero VAD v6.2.0 (#3524)

* Add ggml-silero-v6.2.0 to download candidates

* Make default VAD model ggml-silero-v6.2.0

* Make VAD model in documentations ggml-silero-v6.2.0

2 months agoruby : VAD separately from ASR (#3518)
KITAITI Makoto [Thu, 13 Nov 2025 01:15:26 +0000 (10:15 +0900)]
ruby : VAD separately from ASR (#3518)

* Add Whisper::VAD::Context

* Add test for Whisper::VAD::Context

* Add Whisper::VAD::Segment

* Add Whisper::VAD::Segments

* Add Whisper::VAD::Context#detect

* Define Whisper::VAD::Segments#each

* Define Whisper::VAD::Segment#start_time and #end_time

* Define Whisper::VAD::Segment#deconstruct_keys

* Add tests for Whisper::VAD family

* Add signatures for VAD family

* Add document on VAD in README

* Define Whisper::VAD::Segments#length

* Add test for Whisper::VAD::Segments#length

* Add signature of Segments#length

* Make vad_segments responsible to initialize VAD::Segments

* Remove meaningless argument check

* Check NULL of segments member

* Add tests for Whisper::VAD::Segments

* Initialize Whisper::VAD::Segment on .allocate

* Add tests for Whisper::VAD::Segment

* Check NULL of context member

* Add test for Whisper::VAD::Context.allocate

2 months agosync : llama.cpp
Georgi Gerganov [Sun, 9 Nov 2025 20:01:21 +0000 (22:01 +0200)]
sync : llama.cpp

2 months agosync : ggml
Georgi Gerganov [Sun, 9 Nov 2025 16:49:56 +0000 (18:49 +0200)]
sync : ggml

2 months agovulkan: iGPU memory reporting fix (llama/17110)
Ruben Ortlam [Sun, 9 Nov 2025 08:54:47 +0000 (09:54 +0100)]
vulkan: iGPU memory reporting fix (llama/17110)

* vulkan: use all device-local heaps for memory availability reporting

Co-authored-by: Giuseppe Scrivano <redacted>
* use all available heaps for iGPU memory reporting

* Allow multiple memory types per buffer request for devices with split heaps

---------

Co-authored-by: Giuseppe Scrivano <redacted>
2 months agovulkan: fix mmq out of bounds reads (llama/17108)
Ruben Ortlam [Sun, 9 Nov 2025 08:52:57 +0000 (09:52 +0100)]
vulkan: fix mmq out of bounds reads (llama/17108)

* vulkan: fix mmq out of bounds reads, streamline outdated matmul host code

* fix mul_mat_id quantization call

* Fix compiler warnings

2 months agovulkan: fuse mul_mat_id + mul (llama/17095)
Jeff Bolz [Sun, 9 Nov 2025 08:48:42 +0000 (02:48 -0600)]
vulkan: fuse mul_mat_id + mul (llama/17095)

* vulkan: fuse mul_mat_id + mul

This comes up in qwen3 moe.

* split mul_mat_id fusion tests into a separate class

2 months agometal : retain src and dst buffers during async ops (llama/17101)
Georgi Gerganov [Sun, 9 Nov 2025 06:28:51 +0000 (08:28 +0200)]
metal : retain src and dst buffers during async ops (llama/17101)

2 months agovulkan: Use spec constants for conv2d s/d/p and kernel W/H (llama/16978)
Jeff Bolz [Sat, 8 Nov 2025 19:24:29 +0000 (13:24 -0600)]
vulkan: Use spec constants for conv2d s/d/p and kernel W/H (llama/16978)

* vulkan: Use spec constants for conv2d s/d/p and kernel W/H

Also add some additional unroll hints, which seems to help.

* lock around map lookup

2 months agoRevert "CUDA: add expert reduce kernel (ggml/16857)" (llama/17100)
Aman Gupta [Sat, 8 Nov 2025 13:05:19 +0000 (21:05 +0800)]
Revert "CUDA: add expert reduce kernel (ggml/16857)" (llama/17100)

2 months agoCUDA: skip fusion for repeating adds in bias (llama/17080)
Aman Gupta [Sat, 8 Nov 2025 08:58:05 +0000 (16:58 +0800)]
CUDA: skip fusion for repeating adds in bias (llama/17080)

2 months agovulkan: Increase BK to 32; use BK/4 for non-CM mul_mm.comp (llama/16636)
SavicStefan [Sat, 8 Nov 2025 08:28:22 +0000 (09:28 +0100)]
vulkan: Increase BK to 32; use BK/4 for non-CM mul_mm.comp (llama/16636)

Signed-off-by: Stefan Savic <redacted>
Co-authored-by: Stefan Savic <redacted>
2 months agoggml: disable vxe for cross-compilation by default (llama/16966)
Aleksei Nikiforov [Sat, 8 Nov 2025 08:00:20 +0000 (09:00 +0100)]
ggml: disable vxe for cross-compilation by default (llama/16966)

Otherwise compilation will fail due to enabling -mvx -mzvector
and not setting corresponding -march options.

2 months agovulkan: fuse rms_norm + mul + rope (+ view + set_rows) (llama/16977)
Jeff Bolz [Sat, 8 Nov 2025 07:52:15 +0000 (01:52 -0600)]
vulkan: fuse rms_norm + mul + rope (+ view + set_rows) (llama/16977)

This change combines the rms_norm+mul and rope+view+set_rows fusions to
allow fusing the whole sequence together. This comes up in Qwen3, Bailing,
and some other models.