sync : llama.cpp

ggml-ci

cuda: make im2col a little faster (llama/15025)

cuda, sycl : fix batched gemm when ne02 == 1 && ne03 > 1 (llama/15038)

* cuda, sycl : fix batched gemm when ne02 == 1 && ne03 > 1

ggml-ci

* cont : fix cont types

ggml-ci

* cont : adopt variable names and comment from the other branch

vulkan: coopmat2 mul_mat optimizations (llama/14934)

- Increase tile size for k-quants, to match non-k-quants
- Choose more carefully between large and medium tiles, considering how it
  interacts with split_k
- Allow larger/non-power of two split_k, and make the splits a multiple of 256
- Use split_k==3 to when >1/2 and <=2/3 of the SMs would hae been used

vulkan: Support ne[3]>1 in noncontig matrix-vector multiply (llama/15015)

vulkan: optimizations for direct convolution (llama/14933)

* vulkan: optimizations for direct convolution

- Empirically choose a better tile size. Reducing BS_K/BS_NPQ helps fill
  the GPU. The new size should be amenable to using coopmat, too.
- Fix shmem bank conflicts. 16B padding should work with coopmat.
- Some explicit loop unrolling.
- Skip math/stores work for parts of the tile that are OOB.
- Apply fastdiv opt.
- Disable shuffles for NV.

* Three tiles sizes for CONV_2D, and a heuristic to choose

* reallow collectives for pre-Turing

* make SHMEM_PAD a spec constant

* fixes for intel perf - no shmem padding, placeholder shader core count

* shader variants with/without unrolling

* 0cc4m's fixes for AMD perf

Co-authored-by: 0cc4m <redacted>
---------

Co-authored-by: 0cc4m <redacted>

CUDA: fix MMQ nwarps for AMD with warp_size==32 (llama/15014)

opencl: add f16 for `add`, `sub`, `mul`, `div` (llama/14984)

ggml : Q2k interleaving implementation - x86/x64 SIMD (llama/14373)

* Initial Q2_K Block Interleaving Implementation

* Addressed review comments and clean up of the code

* Post rebase fixes

* Initial CI/CD fixes

* Update declarations in arch-fallback.h

* Changes for GEMV Q2_K in arch-fallback.h

* Enable repacking only on AVX-512 machines

* Update comments in repack.cpp

* Address q2k comments

---------

Co-authored-by: Manogna-Sree <redacted>

docker : add cann build pipline (llama/14591)

* docker: add cann build pipline

* docker: add cann build pipline

* docker: fix cann devops

* cann : fix multi card hccl

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

Co-authored-by: Xuan-Son Nguyen <redacted>
* Update ggml-cann.cpp

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Xuan-Son Nguyen <redacted>

Vulkan: Fix minor debug mode issues (llama/14899)

* vulkan: fix debug mode issues

* vulkan: remove broken check_results GGML_OP_SET_ROWS support

CANN: Improve loading efficiency after converting weights to NZ format. (llama/14985)

* CANN: Improve loading efficiency after converting weights to NZ format.

* CANN: fix typo

opencl: add `mul_mat_f32_f32_l4_lm` and `mul_mat_f16_f32_l4_lm` (llama/14809)

HIP: enable mfma mmq on gfx908 and gfx90a for select datatypes and shapes (llama/14949)

CUDA: skip masked KV slices for all FA kernels (llama/14924)

HIP: remove the use of __HIP_PLATFORM_AMD__, explicitly support only AMD targets (llama/14945)

HIP: add GGML_HIP_MMQ_MFMA option to allow disableing the MFMA path. (llama/14930)

This is useful for testing for regressions on GCN with CDNA hardware.

With GGML_HIP_MMQ_MFMA=Off and GGML_CUDA_FORCE_MMQ=On we can conveniently test the GCN code path on CDNA. As CDNA is just GCN renamed with MFMA added and limited use ACC registers, this provides a good alternative for regression testing when GCN hardware is not available.

HIP: Ignore unsupported unroll transformation in fattn-vec (llama/14931)

llvm with the amdgcn target dose not support unrolling loops with conditional break statements, when those statements can not be resolved at compile time. Similar to other places in GGML lets simply ignore this warning.

CANN: Add ggml_set_rows (llama/14943)

cuda : add softcap fusion (llama/14907)

CUDA: add roll (llama/14919)

* CUDA: add roll

* Make everything const, use __restrict__

test-backend-ops : extend test case filtering (llama/14865)

* Extend test case filtering

1. Allow passing multiple (comma-separated?) ops to test-backend-ops. This can be convenient when working on a set of ops, when you'd want to test them together (but without having to run every single op). For example:

`test-backend-ops.exe test -o "ADD,RMS_NORM,ROPE,SILU,SOFT_MAX"`

2. Support full test-case variation string in addition to basic op names. This would make it easy to select a single variation, either for testing or for benchmarking. It can be particularly useful for profiling a particular variation (ex. a CUDA kernel), for example:

`test-backend-ops.exe perf -b CUDA0 -o "MUL_MAT(type_a=f16,type_b=f32,m=4096,n=512,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=2)"`

These two can be combined. As the current `-o`, this change doesn't try to detect/report an error if an filter doesn't name existing ops (ex. misspelled)

* Updating the usage help text

* Update tests/test-backend-ops.cpp

ggml-cpu : deduplicate scalar implementations (llama/14897)

* remove redundant code in riscv

* remove redundant code in arm

* remove redundant code in loongarch

* remove redundant code in ppc

* remove redundant code in s390

* remove redundant code in wasm

* remove redundant code in x86

* remove fallback headers

* fix x86 ggml_vec_dot_q8_0_q8_0

SYCL: Add set_rows support for quantized types (llama/14883)

* SYCL: Add set_rows support for quantized types

This commit adds support for GGML_OP_SET_ROWS operation for various
quantized tensor types (Q8_0, Q5_1, Q5_0, Q4_1, Q4_0, IQ4_NL) and BF16
type in the SYCL backend.

The quantization/dequantization copy kernels were moved from cpy.cpp
to cpy.hpp to make them available for set_rows.cpp.

This addresses part of the TODOs mentioned in the code.

* Use get_global_linear_id() instead

ggml-ci

* Fix formatting

ggml-ci

* Use const for ne11 and size_t variables in set_rows_sycl_q

ggml-ci

* Increase block size for q kernel to 256

ggml-ci

* Cleanup imports

* Add float.h to cpy.hpp

CUDA: fix pointer incrementation in FA (llama/14916)

sycl: refactor quantization to q8_1 (llama/14815)

* sycl: quantization to q8_1 refactor

* Refactored src1 copy logic in op_mul_mat

ci : Move msvc to matrix (#1318)

Enable static builds and testing

simple : fix typo (#1319)

sync : whisper.cpp

cmake : Fix BLAS link interface (#1316)

vulkan : fix 32-bit builds (#1313)

The pipeline member can be cast to VkPipeline.
This is a VkPipeline_T* on 64 bit but a uint64_t on 32 bit.
Cf. VK_DEFINE_NON_DISPATCHABLE_HANDLE documentation.

sync : llama.cpp

ggml-ci

vulkan : add fp16 support for the conv_2d kernel (llama/14872)

* add f16 to conv_2d testing
* weaken conv2d test error threshold

vulkan: skip empty set_rows to avoid invalid API usage (llama/14860)

Docs: add instructions for adding backends (llama/14889)

HIP: Enable Matrix cores for MMQ Kernels, Enable stream-K for CDNA 3 (llama/14624)

This commit adds support for MFMA instructions to MMQ. CDNA1/GFX908 CDNA2/GFX90a and CDNA3/GFX942 are supported by the MFMA-enabled code path added by this commit. The code path and stream-k is only enabled on CDNA3 for now as it fails to outperform blas in all cases on the other devices.
Blas is currently only consistently outperformed on CDNA3 due to issues in the amd-provided blas libraries.
This commit also improves the awareness of MMQ towards different warp sizes and as a side effect improves the performance of all quant formats besides q4_0 and q4_1, which regress slightly, on GCN gpus.

CANN: Implement GLU ops (llama/14884)

Implement REGLU, GEGLU, SWIGLU ops according to #14158

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

Signed-off-by: Xiaodong Ye <redacted>

ggml-cpu : disable GGML_NNPA by default due to instability (llama/14880)

* docs: update s390x document for sentencepiece

Signed-off-by: Aaron Teo <redacted>
(cherry picked from commit e086c5e3a7ab3463d8e0906efcfa39352db0a48d)

* docs: update huggingface links + reword

Signed-off-by: Aaron Teo <redacted>
(cherry picked from commit 8410b085ea8c46e22be38266147a1e94757ef108)

* ggml-cpu: disable ggml-nnpa compile flag by default

fixes #14877

Signed-off-by: Aaron Teo <redacted>
(cherry picked from commit 412f4c7c88894b8f55846b4719c76892a23cfe09)

* docs: update s390x build docs to reflect nnpa disable

Signed-off-by: Aaron Teo <redacted>
(cherry picked from commit c1eeae1d0c2edc74ab9fbeff2707b0d357cf0b4d)

---------

Signed-off-by: Aaron Teo <redacted>

metal: SSM_SCAN performance (llama/14743)

* feat: Add s_off as a parameter in the args struct

This may not be necessary, but it more closely mirrors the CUDA kernel

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* perf: Parallelize mamba2 SSM_SCAN metal kernel over d_state

This is a first attempt at optimizing the metal kernel. The changes here
are:

- Launch the kernel with a thread group of size d_state
- Use simd groups and shared memory to do the summation for the y
  computation

When tested with G4 tiny preview, this shows roughly a 3x speedup on
prefill and 15% speedup on decode.

Signed-off-by: Gabe Goodhart <redacted>
* fix: Update logic to correctly do the multi-layer parallel sum

Signed-off-by: Gabe Goodhart <redacted>
* fix: Correctly size the shared memory bufer and assert expected size relationships

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* refactor: Compute block offsets once rather than once per token

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* feat: Use local variable for state recursion

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* feat: Use a secondary simd_sum instead of a for loop

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* feat: Add assertion and comment about relationship between simd size and num simd groups

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* feat: Parallelize of d_state for mamba-1

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* feat: Parallel sum in SSM_CONV

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
* Revert "feat: Parallel sum in SSM_CONV"

After discussion with @compilade, the size of the parallelism here is
not worth the cost in complexity or overhead of the parallel for.

https://github.com/ggml-org/llama.cpp/pull/14743#discussion_r2223395357

This reverts commit 16bc059660c1c59e566628201c0ca2c20c9f4bc3.

Signed-off-by: Gabe Goodhart <redacted>
* refactor: Simplify shared memory sizing

Branch: GraniteFourPerf

Signed-off-by: Gabe Goodhart <redacted>
Co-Authored-By: Georgi Gerganov <redacted>
---------

Signed-off-by: Gabe Goodhart <redacted>
Co-authored-by: Georgi Gerganov <redacted>

opencl: add fused `rms_norm_mul` (llama/14841)

* opencl: add fused `rms_norm` + `mul`

* opencl: improve workgroup size for `rms_norm_mul`

ggml : remove invalid portPos specifiers from dot files (llama/14838)

Neither "g" nor "x" are valid portPos specifiers per the official
[graphviz documents](https://graphviz.org/docs/attr-types/portPos/):

> If a compass point is used, it must have the form "n","ne","e","se","s","sw","w","nw","c","_".

I tested locally for it to fall back to default portPos specifier if an
invalid portPos is specified. As a consequence, we can remove associated
code.

rpc : check for null buffers in get/set/copy tensor endpoints (llama/14868)

sched : fix multiple evaluations of the same graph with pipeline parallelism (llama/14855)

ggml-ci

musa: upgrade musa sdk to rc4.2.0 (llama/14498)

* musa: apply mublas API changes

Signed-off-by: Xiaodong Ye <redacted>
* musa: update musa version to 4.2.0

Signed-off-by: Xiaodong Ye <redacted>
* musa: restore MUSA graph settings in CMakeLists.txt

Signed-off-by: Xiaodong Ye <redacted>
* musa: disable mudnnMemcpyAsync by default

Signed-off-by: Xiaodong Ye <redacted>
* musa: switch back to non-mudnn images

Signed-off-by: Xiaodong Ye <redacted>
* minor changes

Signed-off-by: Xiaodong Ye <redacted>
* musa: restore rc in docker image tag

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

Signed-off-by: Xiaodong Ye <redacted>

contrib : recommend PRs to llama.cpp (#1312)

* contrib : recommend PRs to llama.cpp

* cont : wording

cmake : Indent ggml-config.cmake (#1310)

sync : llama.cpp

ggml-ci

sycl: fixed semantics of block offset calculation (llama/14814)

metal : fix fusion across different encoders (llama/14849)

* metal : fix fusion across different encoders

ggml-ci

* cont : add assertion

ggml-ci

sycl: fix undefined variable in work group size check (llama/14843)

CUDA: fix overflow in FA, tune performance (llama/14840)

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)