CUDA: fix compilation with GGML_CUDA_F16 (llama/14837)

CUDA: fix quantized KV cache + multiple sequences (llama/14822)

* CUDA: fix quantized KV cache + multiple sequences

* Update src/ggml-cuda/fattn-common.cuh

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

Co-authored-by: Georgi Gerganov <redacted>

tests : add non-cont K,V FA tests

ggml-ci

ggml: fix loongarch quantize_row_q8_1 error (llama/14827)

CANN: weight format to NZ for Ascend310P3 (llama/14407)

* weight format to nz for 310p

* remove quant weight format to nz

* clean code

* fix

* make the conditions for converting weights to NZ format consistent

* clean code

CUDA: add fused rms norm (llama/14800)

vulkan: fix rms_norm_mul to handle broadcasting dim0 (llama/14817)

cuda : implement bf16 cpy ops and enable bf16 cont (llama/14763)

* implement bf16 cpy ops and enable bf16 cont

* deduplicate copy functions

* deduplicate checks

opencl: remove unreachable `return` (llama/14806)

cuda: remove linking to cublasLt (llama/14790)

Signed-off-by: Xiaodong Ye <redacted>

opencl: fix `im2col` when `KW!=KH` (llama/14803)

opencl: add conv2d kernel (llama/14403)

* add conv2d kernel

* fix trailing whitespace

* whitespace fixe

* handle f16 input and f16 kernel, more opt

* resolve conflicts

* use enqueue_ndrange_kernel

sycl: Fix im2col (llama/14797)

kleidiai: add support for get_rows (llama/14676)

* kleidiai: add support for get_rows

* apply fixes based on code review

* apply more fixes based on code review

vulkan/cuda: Fix im2col when KW!=KH (llama/14789)

The tid is decomposed into "ow + ky*OW + kx*OW*KH". Change "ksize" to match.

ggml: adds CONV_2D op and direct GEMM Vulkan implementation (llama/14316)

* ggml/ggml-vulkan/test-backend-ops: adds CONV_2D for Vulkan

* ggml-vulkan: adds f32 scalar shader to compute 2D convolution directly
with gemm (no need for im2col),

* test-backend-ops: adds test_case_ref to check the validity/performance of ops
against reference implementations having different graphs, adds tests

* * Performance fixes: minimized branch divergence, uses collectives to
  eliminate redundant calculation, macros removed.

* Kernel shared memory size check

* Updates test-backend-ops to support graphs for performance
  measurement.

* * Apple/Win32 compile errors fixed

* Subgroup size used to determine tile size -> fixes llvmpipe errors.

* Collectives disabled by default.

* Intel support is disabled as the performance is poor.

* Conv2d enabled for Intel with disabled collectives, disabled for Apple

* test-backend-ops modifications are reverted

* Trailing spaces and missing override fixed.

* Triggering pipeline relaunch.

* Code formatted with .clang-format.

vulkan: Add logging for bf16 features to ggml_vk_print_gpu_info (#13274) (llama/14707)

Vulkan: Fix fprintf format-security warning (llama/14770)

CI: Test static build (#1307)

cmake : fix usage issues (#1257)

* CMake config: Create target only once

Fix error on repeated find_package(ggml).
For simplicity, check only for the top-level ggml::ggml.

* CMake config: Add CUDA link libs

* CMake config: Add OpenCL link libs

* CMake config: Use canonical find_dependency

Use set and append to control link lib variables.
Apply more $<LINK_ONLY...>.

* CMake config: Wire OpenMP dependency

ggml-cpu : remove stdlib include from repack.cpp (#1276)

This commit removes the inclusion of `<cstdlib>`.

The motivation for this change is that this source file does not seem to
use any functions from this header and the comment about `qsort` is a
little misleading/confusing.

sync : llama.cpp

ggml-ci

metal : fuse add, mul + add tests (llama/14596)

ggml-ci

cuda : Fix Gemma3n not executed as CUDA_GRAPH on NVGPUs (llama/14741)

* Fix Gemma3n not executed as CUDA_GRAPH on NVGPUs

Gemma3n uses Matrix-Matrix addition as part of their input processing,
wrongly triggering CUDA_GRAPH disablement on NVGPUs even when batch-size
of 1 is used.

* Exclude `project_per_layer_input` by matching node names

This ensures that all other graphs which don't exhibit this pattern do
not have their behavior changed.

* Revert unnecessary formatting changes

CUDA: set_rows + cpy.cu refactor (llama/14712)

use max work group size for device to replace the magic number (llama/14732)

ggml: Add initial WebGPU backend (llama/14521)

* Minimal setup of webgpu backend with dawn. Just prints out the adapter and segfaults

* Initialize webgpu device

* Making progress on setting up the backend

* Finish more boilerplate/utility functions

* Organize file and work on alloc buffer

* Add webgpu_context to prepare for actually running some shaders

* Work on memset and add shader loading

* Work on memset polyfill

* Implement set_tensor as webgpu WriteBuffer, remove host_buffer stubs since webgpu doesn't support it

* Implement get_tensor and buffer_clear

* Finish rest of setup

* Start work on compute graph

* Basic mat mul working

* Work on emscripten build

* Basic WebGPU backend instructions

* Use EMSCRIPTEN flag

* Work on passing ci, implement 4d tensor multiplication

* Pass thread safety test

* Implement permuting for mul_mat and cpy

* minor cleanups

* Address feedback

* Remove division by type size in cpy op

* Fix formatting and add github action workflows for vulkan and metal (m-series) webgpu backends

* Fix name

* Fix macos dawn prefix path

llama : add high-throughput mode (llama/14363)

* kv-cache : prepare K/V buffers for separation

ggml-ci

* batched-bench : fix oob write

ggml-ci

* llama : add "virtual sequences"

ggml-ci

* llama : use "stream" vs "virtual sequence"

ggml-ci

* graph : fix stream splitting when KV cache is not used

ggml-ci

* kv-cache : add multi-stream save/load support

ggml-ci

* llama : add "--attn-streams" flag

ggml-ci

* kv-cache : fix handling when find_slot fails

ggml-ci

* kv-cache : restore find_slot impl

ggml-ci

* kv-cache : add comments

* kv-cache : add bounds checks for sequence id

ggml-ci

* cont : add n_seq_max to batch allocr

ggml-ci

* kv-cache : perform stream copies lazily after llama_synchronize

ggml-ci

* kv-cache : avoid throwing exceptions across the C boundary

ggml-ci

* CUDA: 4D FlashAttention support (llama/14628)

* CUDA: 4D FlashAttention support

* CUDA: fix WMMA FA kernel

* llama : rename attn_streams -> kv_unified

ggml-ci

* common : rename kv_split -> kv_unified

ggml-ci

---------

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

ggml : add asserts (llama/14720)

* ggml : add asserts

ggml-ci

* cont : fix constant type

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

Co-authored-by: Diego Devesa <redacted>

vulkan: fix noncontig check for mat_mul_id splitting (llama/14683)

* vulkan: fix noncontig check for mat_mul_id splitting

Remove supports_op check for > 4096 (splitting fixes this)

* vulkan: fix batched matmul dequant for Q*_K

vulkan: add RTE variants for glu/add/sub/mul/div (llama/14653)

cuda: fix build warnings in set-rows.cu (unused variable) (llama/14687)

Signed-off-by: Xiaodong Ye <redacted>

sycl: Hotfix for non dnnl codepath (llama/14677)

ggml : refactor llamafile_sgemm PPC code (llama/14673)

Remove un-necessary templates from class definition and packing functions
Reduce deeply nested conditionals, if-else switching in mnapck function
Replace repetitive code with inline functions in Packing functions

2 ~ 7% improvement in Q8 Model
15 ~ 50% improvement in Q4 Model

Signed-off-by: Shalini Salomi Bodapati <redacted>

SYCL: use 1D kernel for set_rows (llama/14618)

* SYCL: Use 1D kernel for set_rows

* Remove dangling comment

* Refactor and use ceil_div

sycl: Batched mulmat rework for oneDNN dispatch (llama/14617)

cuda : add set rows for bf16 (llama/14664)

cuda : add ELU support (llama/14657)

ggml : add build-time message to remind about ggml_set_rows (llama/14661)

ggml-ci

metal : Add missing unary ops Metal support (llama/14660)

tests : cover lfm2 cases in test_ssm_conv (llama/14651)

CUDA: add set rows for f32 and f16 (llama/14551)

* CUDA: add set rows for f32 and f16

* Review: change kernel params, use strides from host

* Use 1-d kernel

* Review: use int64_t for blockDim.x, rename nb->s for clarity

sync : whispper.cpp

git : remove kompute submodule (#1300)

ggml-ci

sync : resolve conflicts (#0)

ggml-ci

sync : llama.cpp

ggml-ci

vulkan: support SET_ROWS (llama/14587)

* vulkan: support SET_ROWS

Add variants of the copy_to_quant shader that do the SET_ROWS operation.
Change these shaders to spread the work across the workgroup.
The memory access pattern is probably not great (one thread per quant block),
but should be fine for now.

* vulkan: optimize set_rows

Larger workgroups for non-quant types.
Set "norepeat" (there is manual repeat logic).
Use fastmod.

vulkan: optimizations for deepseek prompt processing (llama/14555)

* vulkan: allow unclamped loads in coopmat2 mul_mat_id shader

* vulkan: increase coopmat2 mul_mat_id tile size

* vulkan: optimize mat_mul_id row_ids search to batch loads, and port to coopmat1 path

* vulkan: use smaller FA row size when head size is large. applies to both scalar and CM2 paths (CM1 isn't used due to shared memory limits)

model : support LiquidAI LFM2 hybrid family (llama/14620)

**Important**
LFM2 was [merged ](https://github.com/huggingface/transformers/pull/39340)into transformers, but has not yet been released.
To convert into gguf, install transformers from source
```shell
pip install "transformers @ git+https://github.com/huggingface/transformers.git@main"
```

HIP : Add HIP 7.0+ compatibility for hipBLAS compute types (llama/14634)

opencl: add tiled mul_mat_f16_f32 (llama/14535)

* add tiled mul_mat_f16_f32

* fix trailing whitespace

* add insightful comments

opencl: add `set_rows` for `f16` and `f32` (llama/14547)

* opencl: add `set_rows` for `f16` and `f32`

* opencl: better choose workgroup size for `set_rows`

Docs: script to auto-generate ggml operations docs (llama/14598)

* Docs: script to auto-generate ggml operations docs

* Review: formatting changes + change github action

* Use built-in types instead of typing

* docs : add BLAS and Metal ops

---------

Co-authored-by: Georgi Gerganov <redacted>

SYCL: Initial set_rows kernel implementation (llama/14562)

* SYCL: Initial set_rows kernel implementation

* Revert max_threads to 256

* Refactor set_rows and address review comments

* Deduplicate conversion function

* Remove guard before kernel launch and refactor

* Fix and add back SFINAE

cuda : support Falcon-H1 state size for SSM_SCAN (llama/14602)

ggml : add ggml_scale_bias (llama/14417)

* ggml : add ggml_scale_bias

* ggml_vec_mad1_f32

* add more simd

* add CUDA

* sycl

* vulkan

* cann (placeholder)

* opencl

* will this fix cpu?

* fix cuda

* suggestions from coderabbit

* fix cann compile error

* vDSP_vsmsa

* rm __ARM_FEATURE_SVE

* use memcpy for op params

* make code looks more consistent

* use scalar for __ARM_FEATURE_SVE

* add x param to ggml_vec_mad1_f32

ggml : prevent integer overflow in gguf tensor size calculation (llama/14595)

vulkan: optimize flash attention split_k_reduce (llama/14554)

* vulkan: allow FA split_k with smaller KV values

* vulkan: spread split_k_reduce work across more threads

k_num can get rather large. Use the whole workgroup to reduce the M/L values.

Launch a thread for each element in the HSV dimension of the output. Helps a
lot for large HSV (like deepseek).

vulkan : fix rope with partial rotation and non-cont src (llama/14582)

cuda : fix rope with partial rotation and non-cont src (llama/14580)

* cuda : fix rope non-cont

ggml-ci

* cont : fix multi-rope + add test

ggml-ci

* sycl : try fix

ggml-ci

* cont : fix sycl + clean-up cuda

ggml-ci

CUDA: add bilinear interpolation for upscale (llama/14563)

musa: fix build warnings (unused variable) (llama/14561)

Signed-off-by: Xiaodong Ye <redacted>

CUDA: add bf16 and i32 to getrows (llama/14529)

vulkan: increase LOAD_VEC_A to 8 (IQ1/IQ2) or 4 (IQ3) (llama/14485)

Commit taken from remyoudompheng's PR https://github.com/ggml-org/llama.cpp/pull/12260

Co-authored-by: Rémy Oudompheng <redacted>

vulkan: fix rms_norm+mul fusion (llama/14545)

The fused operation was grabbing the epsilon value from the wrong place.

Add an env var to disable fusion.

Add some missing checks for supported shapes/types.

Handle fused rms_norm+mul in check_results.

vulkan: Handle updated FA dim2/3 definition (llama/14518)

* vulkan: Handle updated FA dim2/3 definition

Pack mask boolean and n_head_log2 into a single dword to keep the push
constant block under the 128B limit.

* handle null mask for gqa

* allow gqa with dim3>1

opencl: add GELU_ERF (llama/14476)

test-backend-ops: add support for specifying output format (llama/14368)

* test-backend-ops: add support for specifying output format

Signed-off-by: Xiaodong Ye <redacted>
* Address review comments

Signed-off-by: Xiaodong Ye <redacted>
* Add build_commit and build_number in test_result

Signed-off-by: Xiaodong Ye <redacted>
* Address review comments

Signed-off-by: Xiaodong Ye <redacted>
* refactor

Signed-off-by: Xiaodong Ye <redacted>
* Get build commit from ggml_commit()

Signed-off-by: Xiaodong Ye <redacted>
* Merge errors into test_operation_info && address review comments

Signed-off-by: Xiaodong Ye <redacted>
* Address review comments

Signed-off-by: Xiaodong Ye <redacted>
* Address review comments

Signed-off-by: Xiaodong Ye <redacted>
* remove visitor nonsense

* remove visitor comment

Signed-off-by: Xiaodong Ye <redacted>
* Address review comments

Signed-off-by: Xiaodong Ye <redacted>
---------

Signed-off-by: Xiaodong Ye <redacted>
Co-authored-by: slaren <redacted>

metal : disable fast math in all quantize kernels (llama/14528)

ggml-ci

CANN: Replace aclrtMemsetSync with aclnnInplaceZero operator (llama/14002)

Co-authored-by: luyuhong <redacted>

ggml : implement GEGLU_ERF and GEGLU_QUICK ops (llama/14445)

opencl : broadcast for soft_max (llama/14510)

vulkan: support mixed/deepseekR1 FA head sizes (llama/14509)

* vulkan: better parameterize FA by head sizes

* vulkan: support mixed/deepseekR1 FA head sizes

ggml: backward pass for split swiglu (llama/14483)

Fix conditional enabling following arch checks for ggml-sycl (llama/14504)

Signed-off-by: nscipione <redacted>

kv-cache : use ggml_set_rows (llama/14285)

* kv-cache : use ggml_set_rows

ggml-ci

* graph : separate k and v indices

ggml-ci

* cont : remove redundant ifs

ggml-ci

* kv-cache : improve find_slot impl

* kv-cache : bounds-check when accessing slot_info indices

* kv-cache : add comments

ggml-ci

* ggml : add TODOs for adding GGML_OP_SET_ROWS support in the backends

ggml-ci

ggml : fix FA mask dim 2 and 3 (llama/14505)

* ggml : fix FA mask dim 2 and 3

ggml-ci

* backends : unsupport batched FA in CUDA and Vulkan

ggml-ci

* vulkan : disable FA for mask->ne[2] != 1

ggml : remove kompute backend (llama/14501)

ggml-ci

CUDA: add dynamic shared mem to softmax, refactor general usage (llama/14497)

llama : initial Mamba-2 support (llama/9126)

* llama : initial Mamba-2 support

* ggml : SIMD ggml_ssm_scan for Mamba-2

* ggml : improve ggml_mul speed when masking recurrent states

* llama : support running Mamba-Codestral-7B-v0.1

* llama : fix Mamba-2 conv state saving

* ggml : make the ggml_mul fast broadcast path more consistently formatted

* llama : remove unused variable

* llama : add missing break

* convert_hf : prefer SentencePiece tokenizer for Mamba-2 when present

The tokenzier.json of Mamba-Codestral-7B-v0.1 otherwise requires
workarounds to work correctly.

* llama : avoid redundant state copy for Mamba 1 and 2

* metal : attempt to adapt SSM_SCAN for Mamba-2

* metal : fix SSM_SCAN pipeline scope

* metal : use log and exp instead of log1pf and expf in SSM_SCAN

* metal : remove unused arguments for SSM_SCAN

The max index is 31, so trimming the arguments is necessary.

* metal : add back n_seqs to SSM_SCAN args

Whoops, this is needed for the offset in the concatenated output.

* metal : fix SSM_SCAN state head offset

* metal : fix wrong number of tokens per sequence in SSM_SCAN

* ggml : remove unused fast broadcast path in GGML_MUL

This was initially added because states were masked with ggml_mul,
but this is no longer done and so this "optimisation" is no longer
necessary, or at least not worth the additional code complexity.

* ggml : avoid multiply by D in GGML_OP_SSM_SCAN

This makes the weight buft detection in src/llama.cpp simpler.

* convert : transpose Mamba-2 A, D and reshape SSM_NORM

This breaks existing conversions of Mamba-2 models
to avoid some reshapes.

Not sure if it's a good idea,
but it makes the graph slightly cleaner.

* llama : more appropriate SSM_SCAN and SSM_CONV buft support checks

* convert : fix flake8 lint

* metal : fix confusion between ; and ,

* metal : add missing args for nb references in ssm_scan_f32_group

* metal : single-user mamba2 inference works

* kv-cache : remove const_cast when setting inputs for s_copy

And also fix multi-user inference for recurrent models
by using cell_id instead of i as the kv cell index
when populating s_copy.

* convert : avoid AutoConfig for Mamba and Mamba2 hparams

* kv-cache : allow context shift for recurrent models

* graph : fix recurrent state copies when avoiding copies

Works, but using lambda functions might not be that clean.

* ggml : fix mamba2 ssm scan when compiled with SVE

* ggml-cpu : reorder SVE FMA for consistency with other SIMD arches

* cuda : implement ssm scan for Mamba2

There is still room for improvement, but it works!

* cuda : adapt Mamba1 ssm scan to shape changes from Mamba2

* mamba : fix mismatched new and delete size for llm_build_mamba

Subclasses of llm_graph_context cannot have extra fields,
because the called destructor is not the one from the subclass.
This otherwise would cause problems when runnning Mamba-(1|2) inference
when compiled -DGGML_SANITIZE_ADDRESS=ON

* cuda : graceful fallback for Mamba-1 models with weird embd size

CUDA: add softmax broadcast (llama/14475)

* CUDA: add softmax broadcast

* Pass by const ref

* Review: Use blockDims for indexing, remove designated initializers

* Add TODO for noncontigous input/output

CUDA: broadcasting for FlashAttention mask (llama/14500)

vulkan: support softmax/FA batch and broadcast (llama/14449)

sync : llama.cpp

ggml : support bcast ggml_soft_max_ext, ggml_flash_attn_ext (llama/14435)

opencl : fix possible buffer overflow in dump_tensor (llama/14490)

opencl : skip empty nodes on cgraph compute (llama/14491)

opencl : update upscale to support align corners (llama/14488)

ggml : Callback before abort (llama/14481)

* Add a callback that will be called just before abort. This allows apps without a console to display a message to the user and save data if needed.

* Return previous callback to allow callback chaining

* style fixes

---------

Co-authored-by: Diego Devesa <redacted>

ci : disable fast-math for Metal GHA CI (llama/14478)

* ci : disable fast-math for Metal GHA CI

ggml-ci

* cont : remove -g flag

ggml-ci

CANN: update aclnnGroupedMatmulV2 to aclnnGroupedMatmulV3 (llama/14411)

* [CANN]update to aclnnGroupedMatmulV2

Signed-off-by: noemotiovon <redacted>
* Support MUL_MAT_ID on 310p

Signed-off-by: noemotiovon <redacted>
* fix editorconfig

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

Signed-off-by: noemotiovon <redacted>

vulkan: Split large mul_mat_id to fit in shared memory (llama/14451)

add GELU_ERF (llama/14455)

ci : simplify, switch to ninja (#1295)

* CI: Move GGML_N_THREADS to env

* CI: Move macos-13 into matrix

* CI: Build with ninja

* CI: Remove env

examples : Test installed CMake config package (#1294)

* Add test-cmake example

* CI: Run test for installed cmake config

vulkan : implement bilinear interpolation for ggml_upscale/ggml_interpolate (#1291)

* supports GGML_SCALE_MODE_BILINEAR and GGML_SCALE_FLAG_ALIGN_CORNERS

vulkan : implement ggml_roll (#1290)

* vulkan : implement ggml_roll

* vulkan : refactor vk_op_unary_push_constants initialization

ggml : add version function to get lib version (#1286)

* ggml : add version function to get lib version

This commit adds a function `ggml_version()` to the ggml library that
returns the version of the library as a string.

The motivation for this is that it can be useful to be able to
programmatically check the version of the ggml library being used.

Usage:
```c
printf("GGML version: %s\n", ggml_version());
```
Output:
```console
GGML version: 0.0.2219
```

* ggml : add ggml_commit()

---------

Co-authored-by: Georgi Gerganov <redacted>

sync : whisper.cpp

sync : llama.cpp

ggml-ci