metal : fix loop bound in ggml_mem_ranges (llama/16412)

ggml : fix graph reallocation with multiple chunks (llama/16396)

reallocation is needed if a single chunk grows in size,
even if total allocation size stays the same or is lower

vulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE (llama/16354)

* vulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE

Replace maxMemoryAllocationSize check with maxBufferSize when creating buffers.
The maxMemoryAllocationSize limit is a "soft" limit and allocations can succeed
beyond that limit. This allows > 4GB buffers to be allocated on some
implementations (e.g. NVIDIA) and tensors this large can be used for im2col
and mul_mat.

For temporary buffers (prealloc_x/y/etc) check against maxStorageBufferRange.
I'm not sure this check is ideal, but we always use these buffers as a single
full size binding and the limit may be smaller than maxMemoryAllocationSize
or maxBufferSize, so I think this is reasonable.

Replace descriptor range uses of VK_WHOLE_SIZE with a manually computed range.
The maxStorageBufferRange may be smaller than the maxBufferSize or
maxMemoryAllocationSize (and the Vulkan spec warns about this in a note) and
it's invalid usage if VK_WHOLE_SIZE computes a range larger than
maxStorageBufferRange.

With this change, it should be possible to generate videos using wan networks
in stable-diffusion.cpp.

* vulkan: Add env var GGML_VK_FORCE_MAX_BUFFER_SIZE and use stoull

vulkan: Fix FA coopmat1 invalid array indexing (llama/16365)

When computing sinks, the cm1 shader was looping r from 0 to Br rather than
to rows_per_thread. I must have copied this from the scalar path (where it is
correct), and somehow it wasn't causing failures on current drivers.

vulkan: in flash attention, bounds check against nem1 (don't rely on GGML_KQ_MASK_PAD) (llama/16316)

ggml webgpu: add support for soft_max, optimize rms_norm (llama/16357)

* Add inplace softmax

* Move rms_norm to split row approach

* Update debug for supports_op

* clean up debug statements

* Update tests/test-backend-ops.cpp

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

Co-authored-by: Georgi Gerganov <redacted>

model : Apertus model implementation (llama/15852)

* First attempt

* No permute during convert (fixes qk tensors), proper norm application.

* RoPE = NeoX

* Coherence!

* Migrate xielu params from tensors to hyperparameters

* Simple CUDA kernel

* Revert stupid LLM refactorings

* Chat template support

* configchecker / flake8 errors

* Reorder unary.cu

* I do conclude that LLMs are, in fact, stupid.

* Fix after merge

* Final newline

* Make xIELU an UNARY_OP

* Final newline

* Correctly account for parameter shift

* Argh.

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

Co-authored-by: Georgi Gerganov <redacted>
* Refactor: remove unused methods, inline and factorize softplus, add const modifiers

* Revert CUDA changes, implement xIELU as a separate OP

* Pesky newline

* Add float2half / half2float for F16 inputs/outputs

* CUDA variants, attempt 2

* Actually, attempt 3

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

Co-authored-by: Johannes Gäßler <redacted>
* Missing convert header

* Proper formula and reference for xIELU in the comments.

* Modify unary-ops.cpp to add the functor-based logic besides the template system to retain optimizations

* Apply suggestions from code review

Co-authored-by: Sigbjørn Skjæret <redacted>
* Add tensor mappings for Apertus to global list instead

* Fix lazy on scalars

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

Co-authored-by: Johannes Gäßler <redacted>
* Add comment about the constraints on positive/negative alpha

* Change `softplus` to `ggml_softplus`

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Johannes Gäßler <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>

musa: update compile flags (llama/16265)

Signed-off-by: Xiaodong Ye <redacted>

HIP: Disable ROCWMMA fattn on CDNA when compiled against ROCWMMA 2.0.0 (llama/16221)

* HIP: Disable ROCWMMA fatt on CDNA when compiled against ROCWMMA 2.0.0

rocwmma 2.0.0 includes a bug in the code fakeing fp16 accumulation on CDNA

* CUDA: Fix volta condition in ggml_cuda_should_use_wmma_fattn

vulkan: make ggml_vk_default_dispatcher support older vulkan headers (llama/16345)

* make ggml_vk_default_dispatcher support older vulkan headers

* simpilfy with using

opencl: support pad_ext (llama/15888)

ggml webgpu: support for rope,div,sub,glu,scale,cont operators (llama/16187)

* Work on rope

* Simplify inplace operation generation and combine mul/add generation

* Work on rope variants

* implement neox rope

* rope complete

* Add sub,div,glu operators

* implement scale op

* Update cpy shader to handle cont/more types

* formatting

* Update test vars printing for rope,rms_norm

* Avoid ROPE hardcoded constants

* Add TODO to change ROPE constants to enum

Co-authored-by: Georgi Gerganov <redacted>
* fix TODO comment

---------

Co-authored-by: Georgi Gerganov <redacted>

opencl: support ne3 in get_rows (llama/15866)

ggml : bump version to 0.9.4 (#1363)

sync : whisper.cpp [no ci]

sync : llama.cpp

cuda : Enable CUDA Graph usage for Nemotron Nano v2 (NemotronH) (llama/16328)

* Fix Nemotron Nano v2 9B not executing as CUDA Graph on NVIDIA GPUs

* fix to ensure test-backend-ops check passes

metal : dynamic simdgroups for MV kernels (llama/16340)

* metal : dynamic simdgroups for MV kernels

* cont : minor

kleidiai : fix work size and threads sync for fp16 (llama/16246)

tests: override test_set_rows::max_nmse_err to allow for occasional rounding differences (llama/16295)

* tests: override test_set_rows::max_nmse_err to allow for occasional rounding differences

* apply similar error bounds to test_cpy

sync : llama.cpp

ggml: riscv: add riscv spacemit backend (llama/15288)

* ggml: add spacemit backend

Change-Id: I249bdc043485d815a9c351867137bc1e27cc2e23

* add new line at end of file

Change-Id: I889ed1c85fb45e62350ecde0c06f70450cadfbe2

* add riscv zba extension limit

Change-Id: I321eb200f859751727afe5cae13074dfce2bb0ce

* fixed for review comments, file renamed and format

Change-Id: Ia20b6ec24a36638e62e0fe07cf100916a7cce3ce

* fixed for code format, after clang-format

Change-Id: I5dc33a0412da3d3f2d77075d8939185d3009eca2

* use _Float16 instead of __fp16

Change-Id: I039fb02bb95270e641bc4442204e658735859d43

* add ci for riscv64-spacemit-ime-native

Change-Id: I711c1033061df1a289ea77891b2997599dfe8279

* update debian-13-riscv64-spacemit-ime-native ci label

Change-Id: Ifb2b891e2fca57b5da604fce2ac255f27731179a

* remove license comment for spacemit ime

Change-Id: If0dc3ca30a958631ccca0a28b62e0b825f9fb0c3

* upgrade binutils for gcc ime

Change-Id: Ibf2fa74c1064408974cb5b45f044d40987e5fb45

* add spacemit ime cross jobs

Change-Id: I80d74909941d41cb9cd09e51d8baf01c985cbfc6

* remove native compile for riscv64-spacemit-ime

Change-Id: I01920afafdc73fa7424014fd648d243f8ec9e25e

* ci : add caching for spacemit ime cross toolchain

Change-Id: Ic54a192019a2fd982bbd58225ce3bbc38f4053de

* ci: bug fixed for cache path and env

Change-Id: I28c42e10b6fff053bb6580926ca2353448cb042a

* Update .github/workflows/build-linux-cross.yml for cache path

Co-authored-by: Sigbjørn Skjæret <redacted>
* bugfixed for  build-linux-cross.yml,  syntax error

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

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

ggml-backend : add root cause in error message if loading backend library fails (llama/16172)

This PR adds additional information to an error message when loading backend library via ld_load_library() fails. This helps spotting why backend library did not load (missing library, missing dependency or unresolved symbol etc.).

sync : whisper.cpp (#1359)

* ggml : Fix MKL detection by quoting BLAS_INCLUDE_DIRS (whisper/3426)

* sync : whisper.cpp

ci : print results [no ci] (#1358)

ci : add self-hosted workflows (#1357)

* ci : add self-hosted workflows

* ci : sync env vars with llama.cpp

* cont : always install python deps

* cont : cancel ongoing runs

* cont : exclude test-backend-ops from debug build

cmake : remove metal flag (llama/0)

sync : llama.cpp

ggml : check cuda and metal argsort limits and add test (llama/16323)

* check cuda argsort limits and add test

* add metal check

ggml : fix dependencies for ggml_set_rows (llama/16318)

vulkan: Fix validation failure in quantized flash attention (llama/16292)

ggml : fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32 (llama/16307)

* fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32

* add test that fails on simd

vulkan: 64-bit im2col (llama/16135)

* vulkan: 64-bit im2col

Add variants of the im2col shaders that use buffer_device_address/buffer_reference,
and use 64-bit address calculations. This is needed for large convolutions used in
stable-diffusion.cpp.

* fix validation error for large im2col

metal : extend mat-mat multiplication support (llama/16225)

* metal : support mul_mm with src1->type == GGML_TYPE_F16

* metal : support mul_mm_id with src1->type == GGML_TYPE_F16

[no ci]

* metal : mul_mm support ne00 % 32 != 0

* metal : support mul_mm_id with ne00 % 32 != 0

* cont : remove unnecessary unrolls

* cont : simplify data loading

* metal : optimize mul_mm when output bounds checks are not needed

metal : fuse non-sequential nodes (llama/16102)

* metal : fuse non-sequential nodes

* cont : add comment

* cont : simplify bounds checks

vulkan: handle mat_mul with A matrix > 4GB (llama/16176)

* vulkan: handle mat_mul with A matrix > 4GB

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

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

* build fixes

vulkan: support arbitrary KV dimension in flash attention (llama/16160)

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

vulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)

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

CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)

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

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

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

* Review: refactor if statement

CUDA: refactor and deduplicate vector FA kernels (llama/16208)

* CUDA: refactor and deduplicate vector FA kernels

vulkan: throw system error instead of SIGABRT during init on older devices (llama/16156)

* Throw system error on old Vulkan driver rather than SIGABRT

* Optionally handle any potential error in vulkan init

vulkan: support GET_ROWS for k-quants (llama/16235)

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

devops: add s390x & ppc64le CI (llama/15925)

* devops: move s390x and ppc64le ci build

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

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

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

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

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

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

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

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

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

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

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

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

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

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

This reverts commit 3f8db04356033d6c1d7eccc75ca396bc5298250c.

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

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

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

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

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

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

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

Bitfields are allocated in different order on s390x

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

* Update URL for big-endian model

* Update .github/workflows/build.yml

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

---------

Signed-off-by: Aaron Teo <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>

metal : report OOM errors (llama/16274)

common : use cpp-httplib as a cURL alternative for downloads (llama/16185)

* vendor : update httplib

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

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

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

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

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

ggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)

* ggml-cpu: impl mxfp4 s390x

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

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

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

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

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

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

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

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

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

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

This reverts commit 1fe55724e2dc295701101bf838bdd4a512237492.

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

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

Signed-off-by: Aaron Teo <redacted>

musa: fix build warnings (llama/15611)

Signed-off-by: Xiaodong Ye <redacted>

CUDA: add a fused top-K MoE kernel (llama/16130)

* CUDA: add a fused top-K MoE kernel

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

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

* Refactor into ggml_cuda_should_use_topk_moe

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

* Review: format + micro-optimizations

* Fix bug: fix tie breakers

* Add optional norm + clean-up code

* Use smem for final write

* Add bounds check

* Use better memory pattern for writeback

ggml : fix loongarch lsx compilation error (llama/15864)

ggml : remove -dev suffix from release version (#1355)

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

pkg-config: include the new GGML_VERSION as a version (#1348)

Instead of hardcoding 0.0.0

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

The version in CMake was added in #1336

examples : fix typo mismatch in gpt (#1349)

ggml : bump version to 0.9.3 (#1353)

scripts : refactor release script into prepare and finalize stages (#1352)

This commit splits the release process into two distinct stages:

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

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

scripts : fix next dev version calculation [no ci] (#1351)

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

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

sync : llama.cpp

metal : fuse NORM + MUL + ADD, support non-multiples of 4 (llama/16220)

* metal : fuse NORM + MUL + ADD

* metal : support norms of non-multiple of 4

* cont : fix comment [no ci]

metal : relax reorder conditions (llama/16216)

metal : restore im2col perf (llama/16219)

sync : llama.cpp

rpc : use ggml logging facilities

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

ci: run the x64 and arm ci on the github machines instead (llama/16183)

* run the x64 ci on regular machines

* set up the same thing for arm

fix test-quantize-perf just like #12306

* try to disable sve

* add another sve run

llama: print memory breakdown on exit (llama/15860)

* llama: print memory breakdown on exit

ggml : split graph allocations according to backend max buffer size (llama/15815)

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

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

* fix missing newline, apple-clang warning

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

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

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

* fix dyn_tallocr_max_size and initialization

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

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

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

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

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

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

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

* fix warnings (missing static) & debug checks

ggml-cpu: Respect cpumask settings (llama/16164)

ggml : fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl (llama/15928)

* fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl

* change initialization to true

zdnn: refactor codebase + add docs (llama/16178)

* zdnn: initial matmul refactor

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

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

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

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

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

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

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

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

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

Signed-off-by: Aaron Teo <redacted>

ggml-cpu : fix typo in gemm comments [no ci] (llama/16189)

ggml : implement set_rows with i32 index (llama/16159)

* implement set_rows with i32 index

* template fix

* test quantized path

warnings--

* Apply suggestions from code review

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

* deduplicate cuda/sycl and test-fix

* indent++

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

* disable i32 index for webgpu for now

---------

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

ggml : extend ggml_can_fuse to work with non-sequential nodes (llama/16123)

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

* cont : fix wrong bounds check condition

* cont : remove unnecessary overload

ggml : add ggml_op_is_empty (llama/16122)

* ggml : add ggml_op_is_empty

* ggml : move to ggml-impl.h

Vulkan: add conv_transpose_2d operation (llama/16022)

* Vulkan: add conv_transpose_2d operation

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

* Vulkan: fix incorrect indentation in conv_transpose_2d shader

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

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

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

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

vulkan: add RTE variants of exp shader (llama/16165)

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

vulkan: vec dot matrix multiplication fix (llama/16151)

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

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

opencl: fix concat crash on win arm64 with Adreno (llama/15944)

opencl: initial `q8_0` mv support (llama/15732)

vulkan: optimize UMA buffer operations and fix driver hangs (llama/16059)

* vulkan: optimize UMA buffer operations and fix driver hangs

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

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

* vulkan: implement deferred_memset on UMA

---------

Signed-off-by: Giuseppe Scrivano <redacted>

vulkan: fix validation error about VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR (llama/16086)

ggml : prepare for development of 0.9.2-dev

ggml : bump version to 0.9.1

scripts : fix sed usage to work on Mac (#1345)

tests : adjust to new timestep_embedding operator

sync : llama.cpp

vulkan: use vec dot for matrix matrix multiplications (llama/16056)

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

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

ggml : refactor forward_dup for cpu backend (llama/16062)

* ggml : refactor forward_dup for cpu backend

* clean up a bit

* add quant/dequant perf test

ggml-amx : fix ggml_amx_init() on generic Linux (llama/16049)

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

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

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

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

cmake : fix static linking for OpenMP on Unix-like systems (llama/16031)

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

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

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

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

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

opencl: optimize mxfp4 kernels (llama/16037)

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

---------

Co-authored-by: Li He <redacted>

rename optimize_graph to graph_optimize (llama/16082)

CUDA: Optimize PAD_REFLECT_1D (llama/15957)

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

* use fast_div to improve performance

* Apply suggestion from JohannesGaessler

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

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

* use a concise expression to further speedup the cuda kernel

---------

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

CUDA: fix compilation on CC 6.0 (llama/16091)

metal : use function constants for mul_mv_ext kernels (llama/16074)

* metal : use function constants for mul_mv_ext kernels

ggml-ci

* metal : remove NW template argument

ggml-ci

* metal : adjust constants

ggml-ci

cuda : add missing F32<->I32 entries in ggml_cuda_cpy_fn (llama/16060)

metal : improve F32, F16 and BF16 mat-vec multiplication (llama/16057)

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

ggml-ci

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

ggml-ci

metal : avoid call free for non-owned buffer (llama/16067)

metal : handle nil cv during pipeline creation (llama/16065)

ggml-ci

CANN: Remove print (llama/16044)

Signed-off-by: noemotiovon <redacted>

GGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (llama/16018)

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

* Add header for linux builds

* Free staged parameter buffers at once

* Format with clang-format

* Fix thread-safe implementation

* Use device implicit synchronization

* Update workflow to use custom release

* Remove testing branch workflow

* some f32 tests passing

* Disable set_rows until it's implemented

* f32 add all tests passing

* Begin work on set_rows

* Work on set rows

* Add error buffers for reporting unsupported SET_ROWS indices

* Remove extra comments

* Add templated addition, clean up code

* Get addition and multiplication working

* Implement rms_norm

* Add get_rows implementation

* Add new get_rows files

* Refactor use of wg size entry

* Fix compilation

* Try manually unrolled q4_0 quant

* Revert "Try manually unrolled q4_0 quant"

This reverts commit 77f8b96515f7e640ae4b0e44f066321fbc4a6166.

* Move to constant max wg size

* Check for tensor size in supports_op

* Vectorize f32 and change default workgroup size

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

* fix linter errors

* Add in-place tests

---------

Co-authored-by: Neha Abbas <redacted>

metal : refactor + optimize v2 (llama/15995)

sync : llama.cpp