]> git.djapps.eu Git - pkg/ggml/sources/llama.cpp/log
pkg/ggml/sources/llama.cpp
2 months agoci : remove vulkaninfo calls (#16169)
Georgi Gerganov [Mon, 22 Sep 2025 07:16:05 +0000 (10:16 +0300)]
ci : remove vulkaninfo calls (#16169)

2 months agoci : use smaller model (#16168)
Georgi Gerganov [Mon, 22 Sep 2025 06:11:39 +0000 (09:11 +0300)]
ci : use smaller model (#16168)

* ci : switch from gemma to qwen3 0.6b

* ci : use smaller model for some tests

2 months agovulkan: add RTE variants of exp shader (#16165)
Jeff Bolz [Mon, 22 Sep 2025 05:37:17 +0000 (00:37 -0500)]
vulkan: add RTE variants of exp shader (#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 agoci : adjust params for less runtime (#16167)
Georgi Gerganov [Mon, 22 Sep 2025 05:31:40 +0000 (08:31 +0300)]
ci : adjust params for less runtime (#16167)

* ci : adjust params for less runtime

* ci : gate BF16 on some hardware

* ci : move extra tests to Arm runner

2 months agovulkan: vec dot matrix multiplication fix (#16151)
Ruben Ortlam [Mon, 22 Sep 2025 05:22:43 +0000 (07:22 +0200)]
vulkan: vec dot matrix multiplication fix (#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 (#15944)
lhez [Sun, 21 Sep 2025 23:42:10 +0000 (16:42 -0700)]
opencl: fix concat crash on win arm64 with Adreno (#15944)

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

2 months agoci : add label for the RISC-V runner (#16150)
Georgi Gerganov [Sun, 21 Sep 2025 16:00:27 +0000 (19:00 +0300)]
ci : add label for the RISC-V runner (#16150)

2 months agoci : migrate ggml ci to self-hosted runners (#16116)
Georgi Gerganov [Sun, 21 Sep 2025 13:50:45 +0000 (16:50 +0300)]
ci : migrate ggml ci to self-hosted runners (#16116)

* ci : migrate ggml ci to a self-hosted runners

* ci : add T4 runner

* ci : add instructions for adding self-hosted runners

* ci : disable test-backend-ops from debug builds due to slowness

* ci : add AMD V710 runner (vulkan)

* cont : add ROCM workflow

* ci : switch to qwen3 0.6b model

* cont : fix the context size

2 months agovulkan: optimize UMA buffer operations and fix driver hangs (#16059)
Giuseppe Scrivano [Sun, 21 Sep 2025 06:31:55 +0000 (08:31 +0200)]
vulkan: optimize UMA buffer operations and fix driver hangs (#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 (#16086)

2 months agosync : ggml upstream/0.0.6527
Georgi Gerganov [Sat, 20 Sep 2025 09:55:47 +0000 (12:55 +0300)]
sync : ggml

2 months agoggml : introduce semantic versioning (ggml/1336)
Daniel Bevenius [Tue, 16 Sep 2025 04:16:52 +0000 (06:16 +0200)]
ggml : introduce semantic versioning (ggml/1336)

* ggml : introduce semantic versioning

This commit introduces semantic versioning for the GGML library.

The motivation for this is that the current versioning, using build
numbers, makes it difficult to track changes and releases for projects
that use ggml.

The release steps are the following:
1. Sync the changes from llama.cpp using sync-llama-am.sh and after the
   PR has been approved and merged move to step 2.
2. Run scripts/release.sh and specify the type of release, major, minor,
   or patch. This script will handle incrementing the version
   (major|minor|patch), create a new commit with the version change,
   create a tag for the version, and prepare for the next development
   iteration.
3. Inspect the commits/tag and push to master. This will trigger the
   github release workflow which is triggered for new tags which will
   then publish a new release on github.

Example usage:
```console
$ ./scripts/release.sh major --dry-run
[dry-run] - No changes will be made

Step 1: Reading current version...
Current version: 0.9.0-dev
New release version: 1.0.0

Step 2: Updating version in ggml/CMakeLists.txt...
  [dry-run] Would update GGML_VERSION_MAJOR to 1
  [dry-run] Would update GGML_VERSION_MINOR to 0
  [dry-run] Would update GGML_VERSION_PATCH to 0
  [dry-run] Would remove -dev suffix

Step 3: Committing version bump...
  [dry-run] Would commit: 'ggml : bump version to 1.0.0'

Step 4: Creating git tag...
  [dry-run] Would create tag: v1.0.0 with message 'Release version 1.0.0'

Step 5: Preparing for next development cycle...
  [dry-run] Would update GGML_VERSION_MINOR to 1
  [dry-run] Would add -dev suffix back

Step 6: Committing development version...
  [dry-run] Would commit: 'ggml : prepare for development of 1.1.0-dev'

[dry-run] Summary (no changes were made):
  • Would have released version: 1.0.0
  • Would have created tag: v1.0.0
  • Would have set next development version: 1.1.0-dev
```

Refs: https://github.com/ggml-org/ggml/issues/1333

* ggml: create branch for release candidate and check master

* ggml : sign the git tag

2 months agoCUDA : conditionally add cuda architectures (ggml/1341)
Gregor Jasny [Wed, 10 Sep 2025 15:21:11 +0000 (17:21 +0200)]
CUDA : conditionally add cuda architectures (ggml/1341)

2 months agovulkan: use vec dot for matrix matrix multiplications (#16056)
Ruben Ortlam [Sat, 20 Sep 2025 08:42:56 +0000 (10:42 +0200)]
vulkan: use vec dot for matrix matrix multiplications (#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 agoserver: fix SSE and OpenAI compatibility for error messages when streaming (#16109)
Benni [Sat, 20 Sep 2025 05:56:30 +0000 (07:56 +0200)]
server: fix SSE and OpenAI compatibility for error messages when streaming (#16109)

* server: fix SSE and OpenAI compatibility for error messages when streaming

* server: remove obsolete event parameter and use required data fieldname instead

2 months agollama-bench: add --devices and --list-devices support (#16039)
ssweens [Fri, 19 Sep 2025 22:15:21 +0000 (15:15 -0700)]
llama-bench: add --devices and --list-devices support (#16039)

* * llama-bench: add --devices support
- Support --devices same as llama-server
- Provide for benchmarking different device combinations
- Include --list-devices like llama-server for convenience

* fix: field display ordering restored

* fix: integrated the rpc devices
- aimed to mimic the server as much as possible

* cleanup: defaults for list-devices
- handle dup device listing with RPC

* cleanup: remove dup device load calls

* docs: update llama-bench
- added the recently added n-cpu-moe option to the docs while in there

* llama-bench: rpc device simplification
* rpc servers unify with other devices earlier, simplifying code
* --list-devices made stateless and simpler
* various cleanup

2 months agochat: Fix streaming parser for granite models (#15682)
shun095 [Fri, 19 Sep 2025 15:57:30 +0000 (00:57 +0900)]
chat: Fix streaming parser for granite models (#15682)

* fix(chat): fix streaming parser for granite models

* tests: add test cases for Granite models chat parser

2 months agofeat: Improve mobile UI for Settings Dialog (#16084)
Aleksander Grygier [Fri, 19 Sep 2025 07:52:27 +0000 (09:52 +0200)]
feat: Improve mobile UI for Settings Dialog (#16084)

* feat: Improve mobile UI for Settings Dialog

* chore: update webui build output

* fix: Linting errors

* chore: update webui build output

2 months agochat : fix build on arm64 (#16101)
Xuan-Son Nguyen [Fri, 19 Sep 2025 06:02:51 +0000 (13:02 +0700)]
chat : fix build on arm64 (#16101)

2 months agoggml : refactor forward_dup for cpu backend (#16062)
Xuan-Son Nguyen [Fri, 19 Sep 2025 04:31:56 +0000 (11:31 +0700)]
ggml : refactor forward_dup for cpu backend (#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 (#16049)
Adrien Gallouët [Thu, 18 Sep 2025 21:07:26 +0000 (23:07 +0200)]
ggml-amx : fix ggml_amx_init() on generic Linux (#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 (#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 (#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 (#16037)
Shawn Gu [Thu, 18 Sep 2025 19:03:34 +0000 (12:03 -0700)]
opencl: optimize mxfp4 kernels (#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 (#16082)
Jeff Bolz [Thu, 18 Sep 2025 18:46:17 +0000 (13:46 -0500)]
rename optimize_graph to graph_optimize (#16082)

2 months agoCUDA: Optimize PAD_REFLECT_1D (#15957)
Bowen Han [Thu, 18 Sep 2025 18:26:03 +0000 (11:26 -0700)]
CUDA: Optimize PAD_REFLECT_1D (#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 (#16091)
Johannes Gäßler [Thu, 18 Sep 2025 17:28:32 +0000 (19:28 +0200)]
CUDA: fix compilation on CC 6.0 (#16091)

2 months agoAdd resumable downloads for llama-server model loading (#15963)
Eric Curtin [Thu, 18 Sep 2025 15:22:50 +0000 (16:22 +0100)]
Add resumable downloads for llama-server model loading (#15963)

- Implement resumable downloads in common_download_file_single function
- Add detection of partial download files (.downloadInProgress)
- Check server support for HTTP Range requests via Accept-Ranges header
- Implement HTTP Range request with "bytes=<start>-" header
- Open files in append mode when resuming vs create mode for new downloads

Signed-off-by: Eric Curtin <redacted>
2 months agometal : use function constants for mul_mv_ext kernels (#16074)
Georgi Gerganov [Thu, 18 Sep 2025 13:28:41 +0000 (16:28 +0300)]
metal : use function constants for mul_mv_ext kernels (#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 (#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 (#16060)

2 months agoserver : include usage statistics only when user request them (#16052)
Radoslav Gerganov [Thu, 18 Sep 2025 10:36:57 +0000 (13:36 +0300)]
server : include usage statistics only when user request them (#16052)

* server : include usage statistics only when user request them

When serving the OpenAI compatible API, we should check if
{"stream_options": {"include_usage": true} is set in the request when
deciding whether we should send usage statistics

closes: #16048

* add unit test

2 months agollama : bump max seq limit from 64 to 256 (#15916)
Georgi Gerganov [Thu, 18 Sep 2025 09:47:56 +0000 (12:47 +0300)]
llama : bump max seq limit from 64 to 256 (#15916)

ggml-ci

2 months agometal : improve F32, F16 and BF16 mat-vec multiplication (#16057)
Georgi Gerganov [Thu, 18 Sep 2025 09:33:45 +0000 (12:33 +0300)]
metal : improve F32, F16 and BF16 mat-vec multiplication (#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 (#16067)
Jhen-Jie Hong [Thu, 18 Sep 2025 07:06:48 +0000 (15:06 +0800)]
metal : avoid call free for non-owned buffer (#16067)

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

ggml-ci

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

Signed-off-by: noemotiovon <redacted>
2 months agoGGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (#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 (#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 (#15995)
Georgi Gerganov [Wed, 17 Sep 2025 17:38:12 +0000 (20:38 +0300)]
metal : refactor + optimize v2 (#15995)

* metal : improve naming

* metal : refactor device

ggml-ci

* cont : props

ggml-ci

* metal : apply ggml_mem_ranges_t

ggml-ci

* metal : remove GGML_METAL_USE_BF16

ggml-ci

* metal : refactor device buffer

ggml-ci

* cont : fix naming

* metal : sync before destroying the backend

ggml-ci

* metal : refactor context

ggml-ci

* metal : migrate ggml-metal.m to ggml-metal.cpp

ggml-ci

* metal : adjust ops API

ggml-ci

* metal : use C++ to store piplienes

ggml-ci

* metal : migrate ops to separate functions

ggml-ci

* metal : add ggml_metal_library_t

ggml-ci

* metal : improve naming

ggml-ci

* metal : cleanp

ggml-ci

* metal : add support for GGML_OP_LOG

ggml-ci

* metal : fix error handling

ggml-ci

2 months agoSvelteKit-based WebUI (#14839)
Aleksander Grygier [Wed, 17 Sep 2025 17:29:13 +0000 (19:29 +0200)]
SvelteKit-based WebUI (#14839)

2 months agoconvert : add Llama4ForCausalLM (#16042)
Xuan-Son Nguyen [Wed, 17 Sep 2025 17:18:21 +0000 (00:18 +0700)]
convert : add Llama4ForCausalLM (#16042)

* convert : add Llama4ForCausalLM

* handle swa

* half working version

* fix use_kq_norm

* fix use_kq_norm

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

2 months agocommon : Fix corrupted memory error on json grammar initialization (#16038)
David Ribeiro Alves [Wed, 17 Sep 2025 08:08:02 +0000 (01:08 -0700)]
common : Fix corrupted memory error on json grammar initialization (#16038)

Initalizing RESERVED_NAME in is_reserved_name() is not thread
safe and leads to corrupted memory when used from multiple threads
as can be seen in the asan trace below. This fixes the initialization
to make it thread-safe.

    #0 0x000100abd018 in std::__1::pair<std::__1::__hash_iterator<std::__1::__hash_node<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, void*>*>, bool> std::__1::__hash_table<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, std::__1::hash<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>, std::__1::equal_to<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>, std::__1::allocator<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>>::__emplace_unique_key_args<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&>(std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&) __hash_table:1565
    #1 0x000100ab0320 in SchemaConverter::visit(nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&) json-schema-to-grammar.cpp:802
    #2 0x000100aafc48 in std::__1::__function::__func<build_grammar(std::__1::function<void (common_grammar_builder const&)> const&, common_grammar_options const&)::$_2, std::__1::allocator<build_grammar(std::__1::function<void (common_grammar_builder const&)> const&, common_grammar_options const&)::$_2>, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> (std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&, nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&)>::operator()(std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&, nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&) function.h:319
    #3 0x000100a2c938 in std::__1::__function::__func<common_chat_params_init_llama_3_x(minja::chat_template const&, templates_params const&, bool)::$_0::operator()(common_grammar_builder const&) const::'lambda'(nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&), std::__1::allocator<common_chat_params_init_llama_3_x(minja::chat_template const&, templates_params const&, bool)::$_0::operator()(common_grammar_builder const&) const::'lambda'(nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&)>, void (nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&)>::operator()(nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&) function.h:319
    #4 0x000100a139f8 in foreach_function(nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&, std::__1::function<void (nlohmann::json_abi_v3_12_0::basic_json<nlohmann::json_abi_v3_12_0::ordered_map, std::__1::vector, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, bool, long long, unsigned long long, double, std::__1::allocator, nlohmann::json_abi_v3_12_0::adl_serializer, std::__1::vector<unsigned char, std::__1::allocator<unsigned char>>, void> const&)> const&) chat.cpp:762
    #5 0x000100a2a7f4 in std::__1::__function::__func<common_chat_params_init_llama_3_x(minja::chat_template const&, templates_params const&, bool)::$_0, std::__1::allocator<common_chat_params_init_llama_3_x(minja::chat_template const&, templates_params const&, bool)::$_0>, void (common_grammar_builder const&)>::operator()(common_grammar_builder const&) function.h:319
    #6 0x000100aa98f4 in build_grammar(std::__1::function<void (common_grammar_builder const&)> const&, common_grammar_options const&) json-schema-to-grammar.cpp:982
    #7 0x0001009c9314 in common_chat_params_init_llama_3_x(minja::chat_template const&, templates_params const&, bool) chat.cpp:1110
    #8 0x0001009b8afc in common_chat_templates_apply_jinja(common_chat_templates const*, common_chat_templates_inputs const&) chat.cpp:1992
    #9 0x0001009b533c in common_chat_templates_apply(common_chat_templates const*, common_chat_templates_inputs const&) chat.cpp:2074
    #10 0x000100810120 in llamacpp_apply_chat_template+0x724 (predict_oai-98384e17fb94e863:arm64+0x100090120)
    ...

==45482==Register values:
 x[0] = 0x00006020004147f8   x[1] = 0x00006080000013c8   x[2] = 0x0000000000000000   x[3] = 0x0000604006289738
 x[4] = 0x0000000000000002   x[5] = 0x0000000000000001   x[6] = 0x04034000004b4000   x[7] = 0x0000000000000001
 x[8] = 0xbebebebebebebebe   x[9] = 0x17d7d7d7d7d7d7d7  x[10] = 0x00000c04000828ff  x[11] = 0x0000000000000001
x[12] = 0x000000002018d383  x[13] = 0x0000000000000000  x[14] = 0xfa0000000000fafa  x[15] = 0x000010700001ffff
x[16] = 0x000000019dc012c0  x[17] = 0x00000001021284f8  x[18] = 0x0000000000000000  x[19] = 0x00000001700acdc0
x[20] = 0x0000000000000002  x[21] = 0x000000002018d384  x[22] = 0x16dd16fd2e731151  x[23] = 0x0000007000020000
x[24] = 0x0000000100c69c08  x[25] = 0x0000000100c69c20  x[26] = 0x00006080000013c7  x[27] = 0x0000000100c69c00
x[28] = 0x00000001700acd60     fp = 0x00000001700aceb0     lr = 0x0000000100abce30     sp = 0x00000001700acd60
AddressSanitizer can not provide additional info.
SUMMARY: AddressSanitizer: SEGV __hash_table:1565 in std::__1::pair<std::__1::__hash_iterator<std::__1::__hash_node<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, void*>*>, bool> std::__1::__hash_table<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, std::__1::hash<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>, std::__1::equal_to<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>, std::__1::allocator<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>>>::__emplace_unique_key_args<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>>, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&>(std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char>> const&)
Thread T5 created by T0 here:
    #0 0x0001020b99d4 in pthread_create+0x5c (libclang_rt.asan_osx_dynamic.dylib:arm64e+0x359d4)
    #1 0x000100873910 in std::sys::pal::unix::thread::Thread::new::h77254fdd87a28e05+0x118 (predict_oai-98384e17fb94e863:arm64+0x1000f3910)
    #2 0x0001007c7a1c in test::run_test::haeb3c2bcd5ed6cf6+0x76c (predict_oai-98384e17fb94e863:arm64+0x100047a1c)
    #3 0x0001007aedb0 in test::console::run_tests_console::he9d142d704f3a986+0x149c (predict_oai-98384e17fb94e863:arm64+0x10002edb0)
    #4 0x0001007c5758 in test::test_main::hf86a5e20735245b9+0x118 (predict_oai-98384e17fb94e863:arm64+0x100045758)
    #5 0x0001007c5da0 in test::test_main_static::h61ee9c8fd30abca0+0x54 (predict_oai-98384e17fb94e863:arm64+0x100045da0)
    ...

==45482==ABORTING

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

* remove unsupported vulkan devices

* make this happen during selection instead

* pass by reference

2 months agoci : revert back to macos-13 for macOS-latest-cmake-x64 (#16040)
Daniel Bevenius [Wed, 17 Sep 2025 07:34:09 +0000 (09:34 +0200)]
ci : revert back to macos-13 for macOS-latest-cmake-x64 (#16040)

This commit reverts the change of the runs-on parameter for the
macOS-latest-cmake-x64 job back to macos-13 that was make in
Commit 51abc96bdc52ba8cd6ad78dcf12ed9a041d7b442 ("ci : update
macos-latest* jobs to use macos-latest (#15938)").

The motivation for this is that using macos-latest will cause an ARM
based runner to be used, and not an x64 based runner.

Refs: https://github.com/ggml-org/llama.cpp/pull/15938#issuecomment-3300805127

2 months agollama-quant : fix the verification of attention layers for encoder-decoder models...
Jie Fu (傅杰) [Wed, 17 Sep 2025 07:30:55 +0000 (15:30 +0800)]
llama-quant : fix the verification of attention layers for encoder-decoder models (#16023)

Signed-off-by: Jie Fu <redacted>
2 months agoexamples : support encoder-decoder models in the simple example (#16002)
Jie Fu (傅杰) [Wed, 17 Sep 2025 07:29:00 +0000 (15:29 +0800)]
examples : support encoder-decoder models in the simple example (#16002)

Signed-off-by: Jie Fu <redacted>
2 months agomodel : add OLMo3 support (#16015)
Shane A [Wed, 17 Sep 2025 07:01:58 +0000 (00:01 -0700)]
model : add OLMo3 support (#16015)

* Add HF to gguf conversion logic for Olmo3

* Add Olmo3 implementation

* Update rope comment

* Fix indentation

Co-authored-by: Sigbjørn Skjæret <redacted>
* Apply suggestion from @CISC

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

Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agoCANN: Optimize ggml_cann_set_device (#15935)
Chenguang Li [Wed, 17 Sep 2025 06:33:08 +0000 (14:33 +0800)]
CANN: Optimize ggml_cann_set_device (#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 agollama-bench: add --n-cpu-moe support (#15952)
jacekpoplawski [Tue, 16 Sep 2025 14:17:08 +0000 (16:17 +0200)]
llama-bench: add --n-cpu-moe support (#15952)

* llama-bench: add --n-cpu-moe support

Support --n-cpu-moe in llama-bench the same way it is supported by
llama-server.

2 months agoci : use macos-latest for arm64 webgpu build (#16029)
Daniel Bevenius [Tue, 16 Sep 2025 13:27:52 +0000 (15:27 +0200)]
ci : use macos-latest for arm64 webgpu build (#16029)

This commit updates the runs-on field for the macOS arm64 webgpu build
job to use macos-latest instead of just latest.

The motivation for this is that this job can wait for a runner to pick
up the job for a very long time, sometimes over 7 hours. This is an
attempt to see if this change can help reduce the wait time.

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/17754163447/job/50454257570?pr=16004

2 months agoggml : fix padding in timestep embedding kernels (#15932)
Daniel Bevenius [Tue, 16 Sep 2025 13:25:57 +0000 (15:25 +0200)]
ggml : fix padding in timestep embedding kernels (#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 agoci : upload xcframework artifact from ios-xcode-build job (#16010)
Daniel Bevenius [Tue, 16 Sep 2025 11:41:38 +0000 (13:41 +0200)]
ci : upload xcframework artifact from ios-xcode-build job (#16010)

This commit updates the github workflows build.yml file to include steps
for uploading and downloading the xcframework artifact. The
macos-latest-swift job now depends on the ios-xcode-build job and
downloads the xcframework artifact produced by it.

The motivation for this changes is that it takes a long time to build
the xcframework and we are currently doing this twice in the workflow.
With this change, we only build it once and reuse the artifact.

2 months agofix: apply clang-format to CUDA macros (#16017)
Bowen Han [Tue, 16 Sep 2025 06:59:19 +0000 (23:59 -0700)]
fix: apply clang-format to CUDA macros (#16017)

clang-format previously broke long CUDA macros (e.g. __launch_bounds__) into
unreadable line breaks inside template declarations, such as:

  template<int D, int ncols, int nwarps, int VKQ_stride,
           typename KQ_acc_t, bool use_logit_softcap>
      __launch_bounds__(nwarps*ggml_cuda_get_physical_warp_size(), 1)

This change adjusts formatting rules so that CUDA macros remain consistent
and aligned with the surrounding template syntax.

2 months agoci : update macos-latest* jobs to use macos-latest (#15938)
Daniel Bevenius [Tue, 16 Sep 2025 03:57:16 +0000 (05:57 +0200)]
ci : update macos-latest* jobs to use macos-latest (#15938)

* ci : update macos-latest* jobs to use macos-latest

This commit updates the jobs that are named macos-latest* to use the
macos-latest label instead explicit versions.

The motivation for this is that there is currently a mixuture of
versions in this workflow and there are jobs that are failing because
they require a newer version.

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/17644792595/job/50140010907#step:5:1759

* ci : add xcodebuild -downloadPlatform iOS command

2 months agocmake : Do not install tools on iOS targets (#15903)
Yuri Khrustalev [Tue, 16 Sep 2025 02:54:44 +0000 (22:54 -0400)]
cmake : Do not install tools on iOS targets (#15903)

2 months agoAdd LLaDA-7b-MoE diffusion model (#16003)
Aman Gupta [Tue, 16 Sep 2025 02:38:28 +0000 (10:38 +0800)]
Add LLaDA-7b-MoE diffusion model (#16003)

2 months agoCUDA: fix im2col_3d to respect non-contiguous inputs (views) (#15956)
Jake Karnes [Mon, 15 Sep 2025 22:28:31 +0000 (16:28 -0600)]
CUDA: fix im2col_3d to respect non-contiguous inputs (views) (#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 agodocker : enable rocWMMA in ROCm images, add gfx1151 (#15997)
Diego Devesa [Mon, 15 Sep 2025 21:38:52 +0000 (14:38 -0700)]
docker : enable rocWMMA in ROCm images, add gfx1151 (#15997)

2 months agoreleases : switch to rocWMMA develop branch, add gfx1151 (#15992)
Diego Devesa [Mon, 15 Sep 2025 21:38:42 +0000 (14:38 -0700)]
releases : switch to rocWMMA develop branch, add gfx1151 (#15992)

* releases : switch to rocWMMA develop branch, add gfx1151

* remove unused variable ROCM_VERSION

2 months agoSYCL: Add COUNT_EQUAL operator support (#15991)
yael-works [Mon, 15 Sep 2025 16:51:35 +0000 (19:51 +0300)]
SYCL: Add COUNT_EQUAL operator support (#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 agollama-run: Fix model download on Windows (#15988)
Nikolay Popov [Mon, 15 Sep 2025 10:08:30 +0000 (13:08 +0300)]
llama-run: Fix model download on Windows (#15988)

* llama-run: Fix model download on Windows
 * fix SSL error (SSL peer certificate or SSH remote key was not OK)
 * fix program crash on std::filesystem::rename

* llama-run: create a separate method to utilize RAII

* llama-run: handle rename exception

2 months agoCUDA: some micro-optimizations in mmf.cuh for mul_mat_id (#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 (#15926)

2 months agofix KLD percentile output (#15999)
ddh0 [Mon, 15 Sep 2025 07:54:57 +0000 (02:54 -0500)]
fix KLD percentile output (#15999)

In `llama-perplexity`, when using `--kl-divergence`, the KL divergence statistics output mistakenly displays the 99th percentile twice. This change fixes that and correctly displays the 90th percentile as originally intended (presumably).

2 months agomodel : add grok-2 support (#15539)
Sigbjørn Skjæret [Sun, 14 Sep 2025 21:00:59 +0000 (23:00 +0200)]
model : add grok-2 support (#15539)

* add grok-2 support

* type fix

* type fix

* type fix

* "fix" vocab for invalid sequences

* fix expert tensor mapping and spaces in vocab

* add chat template

* fix norm tensor mapping

* rename layer_out_norm to ffn_post_norm

* ensure ffn_post_norm is mapped

* fix experts merging

* remove erroneous FFN_GATE entry

* concatenate split tensors and add more metadata

* process all expert layers and try cat instead of hstack

* add support for community BPE vocab

* fix expert feed forward length and ffn_down concat

* commit this too

* add ffn_up/gate/down, unsure if sequence is right

* add ffn_gate/down/up to tensor names

* correct residual moe (still not working)

* mess--

* fix embedding scale being applied twice

* add built in chat template

* change beta fast for grok if default value

* remove spm vocab in favor of community bpe vocab

* change attention temp length metadata type to integer

* update attention temp length metadata

* remove comment

* replace M_SQRT2 with std::sqrt(2)

* add yarn metadata, move defaults to hparams

2 months agoserver : only attempt to enable thinking if using jinja (#15967)
Sigbjørn Skjæret [Sun, 14 Sep 2025 19:17:04 +0000 (21:17 +0200)]
server : only attempt to enable thinking if using jinja (#15967)

2 months agometal : remove memory pools (#15966)
Georgi Gerganov [Sun, 14 Sep 2025 19:02:32 +0000 (22:02 +0300)]
metal : remove memory pools (#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 agorocm.Dockerfile: added gfx1200,gfx1201 architectures to support AMD Radeon RX 9000...
Adam [Sun, 14 Sep 2025 18:43:54 +0000 (04:43 +1000)]
rocm.Dockerfile: added gfx1200,gfx1201 architectures to support  AMD Radeon RX 9000 series (#15994)

* rocm.Dockerfile: added gfx1200,gfx1201 architectures to support  AMD Radeon RX 9000 series

https://rocm.docs.amd.com/projects/install-on-linux/en/docs-6.4.1/reference/system-requirements.html#rdna-os
states the Radeon RX 9000 series is supported support from Ubuntu 24.04.2, and the dockerfile is using 24.04 which is ROCm 6.4.

This fixed the `ROCm error: invalid device function` I was getting when trying to use the rocm container.

2 months agoVulkan: Clean up mul_mm shader (#15987)
Ruben Ortlam [Sun, 14 Sep 2025 14:56:28 +0000 (16:56 +0200)]
Vulkan: Clean up mul_mm shader (#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 agobuild: fix the build failures of Windows HIP release job (#15984)
lcy [Sun, 14 Sep 2025 14:20:35 +0000 (22:20 +0800)]
build: fix the build failures of Windows HIP release job (#15984)

* build: fix the cache keys for Windows HIP release job

Update the cache keys to include the HIP SDK version, preventing the
use of outdated ROCm installation caches.

* build: sync changes from release.yml to build.yml

- Update HIP SDK version to 25.Q3 and ROCm version to 6.4.2
- Update the cache keys to reflect the new versions

* build: remove Windows HIP release for gfx1151
since the current stable rocWMMA does not support gfx1151.

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

* metal : fix kernel requirements

ggml-ci

* cont : fix supports_op

* cont : fix supports_op for ARGMAX

2 months agorpc : fix regression when --device is used (#15981)
Radoslav Gerganov [Sun, 14 Sep 2025 09:28:18 +0000 (12:28 +0300)]
rpc : fix regression when --device is used (#15981)

Fix regression introduced with commit 50f4281a6

2 months agoreleases : update ROCM, add gfx1200, gfx1201, gfx1151 (#15972)
Diego Devesa [Sun, 14 Sep 2025 09:21:59 +0000 (02:21 -0700)]
releases : update ROCM, add gfx1200, gfx1201, gfx1151 (#15972)

* releases : update ROCM, add gfx1200, gfx1201, gfx1151

* releases : set target to 13.3 for macos-x64

* add hipblaslt.dll to release

* add hipblaslt/library to release

2 months agodoc : update documentation for --tensor-split (#15980)
Radoslav Gerganov [Sun, 14 Sep 2025 09:10:07 +0000 (12:10 +0300)]
doc : update documentation for --tensor-split (#15980)

* doc : update documentation for --tensor-split

* Update tools/main/README.md

Co-authored-by: Johannes Gäßler <redacted>
* Update tools/main/README.md

Co-authored-by: Diego Devesa <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
Co-authored-by: Diego Devesa <redacted>
2 months agoggml-zdnn: rm user mapped buffers (#15965)
Aaron Teo [Sun, 14 Sep 2025 05:37:03 +0000 (13:37 +0800)]
ggml-zdnn: rm user mapped buffers (#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 (#15862)
Jeff Bolz [Sat, 13 Sep 2025 15:29:43 +0000 (16:29 +0100)]
vulkan: fix failing dequant shaders (#15862)

* vulkan: fix failing dequant shaders

* add missing const

2 months agovulkan: initialize vulkan-hpp to allow using extension function pointers (#15705)
Jeff Bolz [Sat, 13 Sep 2025 15:23:30 +0000 (16:23 +0100)]
vulkan: initialize vulkan-hpp to allow using extension function pointers (#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 agollama : allow using iGPUs with --device (#15951)
Diego Devesa [Sat, 13 Sep 2025 14:49:49 +0000 (07:49 -0700)]
llama : allow using iGPUs with --device (#15951)

* llama : allow using iGPUs with --device

* mtmd : allow iGPU

* rpc-server : allow iGPU

2 months agometal : refactor kernel loading (#15964)
Georgi Gerganov [Sat, 13 Sep 2025 13:24:22 +0000 (16:24 +0300)]
metal : refactor kernel loading (#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 (#15929)
Georgi Gerganov [Sat, 13 Sep 2025 10:54:28 +0000 (13:54 +0300)]
metal : allow ops to run concurrently (#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 (#15962)
Georgi Gerganov [Sat, 13 Sep 2025 09:45:04 +0000 (12:45 +0300)]
metal : fix memory leaks (#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 (#15839)

2 months agoAdd docker protocol support for llama-server model loading (#15790)
Eric Curtin [Fri, 12 Sep 2025 15:31:50 +0000 (16:31 +0100)]
Add docker protocol support for llama-server model loading (#15790)

To pull and run models via: llama-server -dr gemma3
Add some validators and sanitizers for Docker Model urls and metadata

Signed-off-by: Eric Curtin <redacted>
2 months agocontext : remove redundant explicit casting to the same type (#15948)
Haiyue Wang [Fri, 12 Sep 2025 15:16:32 +0000 (23:16 +0800)]
context : remove redundant explicit casting to the same type (#15948)

The function 'output_reserve' return type is 'uint32_t', so need to add
explicit casting.

2 months agoserver : adjust prompt similarity thold + add logs (#15913)
Georgi Gerganov [Fri, 12 Sep 2025 14:02:55 +0000 (17:02 +0300)]
server : adjust prompt similarity thold + add logs (#15913)

ggml-ci

2 months agoVulkan iGPU device selection overhaul and PCI ID API support (#15947)
Ruben Ortlam [Fri, 12 Sep 2025 11:24:21 +0000 (13:24 +0200)]
Vulkan iGPU device selection overhaul and PCI ID API support (#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 (#15939)
Mathieu Baudier [Fri, 12 Sep 2025 07:06:20 +0000 (09:06 +0200)]
vulkan: Make device memory check more portable (#15939)

2 months agoRevert "sycl: add usage of enqueue_functions extension (#14244)" (#15910)
Neo Zhang Jianyu [Fri, 12 Sep 2025 01:15:12 +0000 (09:15 +0800)]
Revert "sycl: add usage of enqueue_functions extension (#14244)"  (#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 (#15797)
Diego Devesa [Thu, 11 Sep 2025 20:47:38 +0000 (13:47 -0700)]
ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type (#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 (#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 (#15927)

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

* fix logic for availability of v_dot2_f32_f16

2 months agonitpick : correct MB to MiB (#15934)
ddh0 [Thu, 11 Sep 2025 17:12:34 +0000 (12:12 -0500)]
nitpick : correct MB to MiB (#15934)

MB was incorrectly used for 1024 x 1024 bytes instead of MiB

2 months agoggml-cpu : add check for ARM MATMUL_INT8/i8mm support (#15922)
Daniel Bevenius [Thu, 11 Sep 2025 13:39:12 +0000 (15:39 +0200)]
ggml-cpu : add check for ARM MATMUL_INT8/i8mm support (#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 (#15614)
Charles Xu [Thu, 11 Sep 2025 10:45:40 +0000 (12:45 +0200)]
kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed (#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 (#15933)
hipudding [Thu, 11 Sep 2025 07:59:37 +0000 (15:59 +0800)]
CANN: Disable acl_graph for prefill stage (#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 (#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 (#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 agollama : support T5 models with unequal number of encoder-decoder layers (#15909)
Jie Fu (傅杰) [Wed, 10 Sep 2025 18:51:51 +0000 (02:51 +0800)]
llama : support T5 models with unequal number of encoder-decoder layers (#15909)

* Extend the support of T5 models with different encoder-decoder layers

Signed-off-by: Jie Fu <redacted>
* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update gguf-py/gguf/constants.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update gguf-py/gguf/gguf_writer.py

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

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update src/llama-arch.h

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

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

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

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

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update src/llama-hparams.h

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Co-authored-by: Sigbjørn Skjæret <redacted>
* Rename n_dec_layer --> dec_n_layer

Signed-off-by: Jie Fu <redacted>
* Adapt to cases when dec_n_layer > n_layer

Signed-off-by: Jie Fu <redacted>
---------

Signed-off-by: Jie Fu <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agograph : support non-contiguous Q in build_attn_mha (#15908)
Sigbjørn Skjæret [Wed, 10 Sep 2025 17:08:59 +0000 (19:08 +0200)]
graph : support non-contiguous Q in build_attn_mha (#15908)

* support non-contiguous Q in build_attn_mha

* Update src/llama-graph.cpp

ggml-ci

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

Co-authored-by: Georgi Gerganov <redacted>
2 months agoggml-cpu : fix padding in ggml_timestep_embedding (#15917)
Daniel Bevenius [Wed, 10 Sep 2025 15:31:40 +0000 (17:31 +0200)]
ggml-cpu : fix padding in ggml_timestep_embedding (#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 agometal : make the backend async (#15906)
Georgi Gerganov [Wed, 10 Sep 2025 14:52:35 +0000 (17:52 +0300)]
metal : make the backend async (#15906)

* metal : make the backend async

ggml-ci

* cont : add comments, extend op offload, clean up

ggml-ci

* metal : fix batch size for MUL_MAT_ID

* metal : remove deprecated ggml_backend_metal_buffer_from_ptr

* metal : create only metal buffers, no wrapping of host memory

ggml-ci

* metal : restore .alloc_buffer for buffer_from_ptr_type

ggml-ci

* metal : remove broken implementation of GGML_OP_SET

ggml-ci

* metal : clean-up loose ends, ready for tests

ggml-ci

* metal : support both private and shared buffers

ggml-ci

* metal : enable private buffers + add global device queue

* metal : disable host buffer to prevent races

ggml-ci

* metal : avoid extra copy during set_tensor

ggml-ci

* metal : use separate buffer types for shread and private Metal buffers

ggml-ci

* metal : simplify synchronization logic

ggml-ci

* metal : fix build

ggml-ci

* metal : do not implement cpy_tensor

ggml-ci

* metal : separate implementations for shared and private buffers

ggml-ci

2 months agoci : add caching for ROCm installation in release workflow (#15924)
Daniel Bevenius [Wed, 10 Sep 2025 13:39:57 +0000 (15:39 +0200)]
ci : add caching for ROCm installation in release workflow (#15924)

This commit applies the same caching to the release workflow which
currently exists for the main CI workflow that was introduced in Commit
ff02caf9eed261423289d1531a56536fbf57bfc2 ("ci : cache ROCm installation
in windows-latest-cmake-hip (#15887)").

2 months agotests : filter out no-ops from coverage report (#15900)
Daniel Bevenius [Wed, 10 Sep 2025 12:17:09 +0000 (14:17 +0200)]
tests : filter out no-ops from coverage report (#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