ggml-cuda: check for srcs outside the cgraph (llama/18583)

* ggml-cuda: check for srcs outside the cgraph

* review: use leafs instead

vulkan: fix topk_moe_sigmoid_norm_bias failures in GLM-4.6 (llama/18582)

vulkan: handle quantize_q8_1 overflowing the max workgroup count (llama/18515)

* vulkan: handle quantize_q8_1 overflowing the max workgroup count

* vulkan: Fix small tile size matmul on lavapipe

* fix mul_mat_id failures

CANN: add operator fusion support for ADD + RMS_NORM (llama/17512)

This commit implements operator fusion for ADD + RMS_NORM operations
in the CANN backend to reduce memory access overhead and improve
performance. The fusion is controlled by the GGML_CANN_OPERATOR_FUSION
environment variable (default: false).

Changes:
- Implement ggml_cann_op_add_rms_norm_fused() using ACLNN AddRmsNorm
- Add ggml_cann_can_fuse() to check fusion eligibility
- Integrate fusion logic into computation graph evaluation
- Add test cases for ADD + RMS_NORM fusion
- Update documentation with new environment variable

The fusion combines ADD and RMS_NORM into a single kernel call,
which is more efficient than executing them separately.

sampling : add support for backend sampling (llama/17004)

* sampling : add support for backend sampling

This commit adds support for performing sampling operations on the
backend (e.g. GPU) as part of the model computation graph.

The motivation for this feature is to enable sampling to be performed
directly on the backend as part of the computation graph being executed,
allowing for some or all of the sampling to be done on the backend.

For example, the backend sampler chain might select/sample a token
directly in which case only the sampled token needs to be transferred
from device memory to host memory.

It is also possible for the backend samplers to perform filtering of
the logits, or compute and filter the probability distribution, in
which case only the filtered logits or probabilites need to be
transferred back to system memory for further processing by CPU
samplers.

Currently the backend sampling works in a similar manner to how
pooling works, it is a function that is called by build_graph and the
sampler operations become part of the models computation graph.

* llama-cli : add backend sampler configuration

* server : add backend sampling options/configuration

* webui : add backend sampling options

* ggml : add initial cumsum implementation for CUDA

* sampling : enable all backend sampler tests

This commit enables all exisiting backend sampler tests in the
test-backend-sampler. Previously, some tests were disabled because
there were missing ggml operation implementations.

* graph : do not include llama-model.h

* sampling : always expose sampled_ids

This commit precomputes and caches the full-vocab token id list in
llama_context's constructor, so llama_get_backend_sampled_token_ids_ith
always returns a valid pointer.

The motivation for this is that this enables both common/sampling.cpp
and src/llama-sampling.cpp can simplify their logic.

Not all backends samplers that process logits need to set the
sampled_tokens_id as they may not change the order of the logits, for
example the temperature sampler only scales the logits but does not
change their order. Simliar the logit bias sampler only adds bias to
specific token ids but does not change the order of the logits. In
these cases there will not be a device to host copy of the sampled
token ids, and this is the use case where having this precomputed
list is useful.

* sampling : ensure at most one output token per seq

This commit adds a check in the batch allocator to ensure that when
backend sampling is enabled, at most one output token is specified per
sequence.

* CUDA: Optimize argsort for gpu-based token sampling

Argsort is used for top-k currently. WE optimize argsort by 2 things:

1. Use `DeviceRadixSort` for single-row/sequence to parallelize it
   across our SMs
2. Use `DeviceSegmentedSort` for multi-row/sequence as this is the
   correct entrypoint (the function chooses different execution paths,
   it contains `DeviceSegmentedRadixSort` as one of the paths and will
   choose the best one according to heuristics.
   https://nvidia.github.io/cccl/cub/api/structcub_1_1DeviceSegmentedSort.html#overview

Some perf numbers for a RTX PRO 6000:

On the kernel level, tested with
`GGML_CUDA_DISABLE_GRAPHS=1 ./test-backend-ops -o ARGSORT perf`
Before:
```
  ARGSORT(type=f32,ne=[65000,16,1,1],order=0):                  4130 runs -   359.24 us/run
  ARGSORT(type=f32,ne=[200000,1,1,1],order=0):                  8192 runs -   861.34 us/run
  ARGSORT(type=f32,ne=[200000,16,1,1],order=0):                 1343 runs -  1020.01 us/run
```

After:
```
  ARGSORT(type=f32,ne=[65000,16,1,1],order=0):                  4130 runs -   312.41 us/run
  ARGSORT(type=f32,ne=[200000,1,1,1],order=0):                 16384 runs -    63.48 us/run
  ARGSORT(type=f32,ne=[200000,16,1,1],order=0):                 1343 runs -   874.36 us/run
```

CUDA: disable cuda graph when using n-cpu-moe (llama/18593)

* CUDA: disable cuda graph when using n-cpu-moe

* call ggml_cuda_set_device

ggml-cuda: remove unused params in ggml_cuda_graph (llama/18579)

ggml-cuda: fixes for concurrent streams (llama/18496)

CUDA: only allocate FA tmp buffer if needed (llama/18564)

(Bugfix, ggml-cuda) Pool alloc count fix + small size computation type adjustment (llama/18559)

* CUDA: Fixed obj byte size instead of obj count being passed to pool alloc (fattn-common, dst_tmp_meta)

* CUDA: Explicitly casted some of the int alloc counts before multiplication in argsort

---------

Co-authored-by: pl752 <redacted>

ggml-hexagon: optimize activation function (llama/18393)

* refactor: refactor silu

* refactor: optimize swiglu

* refactor: remove unncessary if in swiglu

* refactor: refactor swiglu_oai

* chore: fix formatting issue

vulkan: Optimize GGML_OP_CUMSUM (llama/18417)

* vulkan: Optimize GGML_OP_CUMSUM

There are two paths: The preexisting one that does a whole row per workgroup
in a single shader, and one that splits each row into multiple blocks and does
two passes. The first pass computes partials within a block, the second adds
the block partials to compute the final result. The multipass shader is used
when there are a small number of large rows.

In the whole-row shader, handle multiple elements per invocation.

* use 2 ELEM_PER_THREAD for AMD/Intel

* address feedback

vulkan: Implement mmvq for iq1_s/iq1_m (llama/18450)

metal : adjust extra size for FA buffer to avoid reallocations (llama/18545)

rpc : use unordered_map::reserve and emplace (llama/18513)

cuda : fix copy of large tensors (ggml_nbytes <= INT_MAX assertion) (llama/18433)

* ggml-cuda: fixed assertion in ggml_cuda_cpy (llama/18140)

* ggml-cuda: changes in data types to int64_t

* ggml-cuda: added asserts for CUDA block numbers

* ggml-cuda: changed the condition for y and z dimension

ggml-cuda: remove unneccesary prints on ggml_cuda_init (llama/18502)

vulkan: extend topk_moe to handle sigmoid w/exp_probs_b for nemotron (llama/18295)

* vulkan: extend topk_moe to handle sigmoid w/exp_probs_b for nemotron

Also handle GGML_OP_SCALE at the end (nemotron, deepseek2).

Fewer pipeline variants and spec constants, just use push constants.

In test_topk_moe, change exp_probs_b to be 1D, matching real networks.

Update test-backend-ops and ggml-backend to allow verifying multiple outputs
in a fusion test (topk_moe has two outputs). Previously only the final node
was verified.

* change test_topk_moe to allow results in arbitrary order

* disable sigmoid fusion for moltenvk

Update upstream

Merge tag 'upstream/0.9.5' into debian/latest

Upstream release

ggml : bump version to 0.9.5 (#1410)

sync : whisper.cpp

sync : llama.cpp

metal : add count_equal op (llama/18314)

* add count equal for metal

* remove trailing whitespace

* updated doc ops table

* changed shmem to i32

* added multi tg and templating

* removed BLAS support from Metal docs

* Apply suggestions from code review

Co-authored-by: Georgi Gerganov <redacted>
* add memset to set dst to 0

* metal : cleanup

---------

Co-authored-by: Georgi Gerganov <redacted>

CUDA: fix KQ max calculation (llama/18487)

metal : remove BF16 x F16 kernels (llama/18456)

sync : llama.cpp

sycl: add newline at the end of CMakeLists.txt (llama/18503)

Work around broken IntelSYCLConfig.cmake in Intel oneAPI 2025.x (llama/18345)

* cmake: work around broken IntelSYCLConfig.cmake in oneAPI 2025.x

* [AI] sycl: auto-detect and skip incompatible IntelSYCL package

Automatically detect compiler versions with incompatible IntelSYCL
CMake configuration files and fall back to manual SYCL flags instead
of requiring users to set options manually.

Fixes build failures with oneAPI 2025.x where IntelSYCLConfig.cmake
has SYCL_FEATURE_TEST_EXTRACT invocation errors.

* refactor: improve SYCL provider handling and error messages in CMake configuration

* refactor: enhance SYCL provider validation and error handling in CMake configuration

* ggml-sycl: wrap find_package(IntelSYCL) to prevent build crashes

kleidiai: add and integrate SVE 256-bit vector-length kernel (llama/18458)

* kleidiai: add and integrate SVE 256-bit vector-length kernel

* updated for review comments

CUDA: add log line when mxfp4 acceleration is used (llama/18483)

* CUDA: add log line when mxfp4 acceleration is used

* add in backend_get_features

CUDA: fix replacment of bad archs in CMake (llama/18457)

CUDA: Blackwell features for non-native builds (llama/18436)

cuda: fix race condition in cumsum (llama/18448)

* ggml-cuda: fix race condition in cumsum

* remove unneccesary sync_threads

HIP: Use mmq on MFMA devices for MUL_MAT_ID in cases where a lot of splits would be generated (llama/18202)

Revert "ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON (#18413)" (llama/18426)

rpc: fix segfault on invalid endpoint format (llama/18387)

* rpc: fix segfault on invalid endpoint format

* rpc: add error log for failed endpoint connection

cmake: Added more x86_64 CPU backends when building with `GGML_CPU_ALL_VARIANTS=On` (llama/18186)

* minor: Consolidated `#include <immintrin.h>` under `ggml-cpu-impl.h`

* cmake: Added more x86-64 CPU backends when building with `GGML_CPU_ALL_VARIANTS=On`

- `ivybridge`
- `piledriver`
- `cannonlake`
- `cascadelake`
- `cooperlake`
- `zen4`

Resolves: #17966

ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON (llama/18413)

opencl: allow resizing transpose buffers (llama/18384)

* opencl: allow resizing transpose buffers instead of using fixed sizes

* opencl: remove commented code

ggml-cuda: Use same regex for GGML_NATIVE=OFF (llama/18407)

vulkan: preprocess mul_mat_id experts and discard workgroups more quickly (llama/18352)

Run a preprocess to count how many times each expert is used, and use this to
quickly discard workgroups that aren't needed.

vulkan: optimize decodeFuncB in coopmat2 mul_mat_id shader (llama/18349)

* vulkan: Use BK=32 for coopmat2 mul_mat_id

* vulkan: optimize decodeFuncB in coopmat2 mul_mat_id shader

Disable robustness, remove the OOB check in decodeFuncB, and initialize the
row_ids to zero to avoid OOB access.

Don't slice/offset the B matrix to ic * BN, only to adjust the coord back down
to the range [0, BN) in decodeFuncB. Instead just slice with a row offset of
zero and remove the '& (BN - 1)'. This allows the compiler to common some of
the shared memory loads.

vulkan: Use BK=32 for coopmat2 mul_mat_id (llama/18332)

vulkan: small dequantization improvements (llama/18380)

* iq4_xs

* quants

vulkan: Support UPSCALE w/antialias (llama/18327)

vulkan: handle rope with large number of rows (llama/18306)

CANN: implement the SSM_CONV operator (llama/17737)

* CANN: implement SSM_CONV operator

Co-authored-by: Aleksei Lobanov, <redacted>
Co-authored-by: Sujin Kang, <redacted>
* CANN: remove custom error limit for SSM_CONV

* CANN: merge SSM_CONV tensor shape/strides into one line

---------

Co-authored-by: Sujin Kang, <redacted>

ggml-cuda: fix regex for arch list (llama/18371)

* ggml-cuda: fix regex for arch list

* make regex exact

cuda: optimize cumsum cub path (llama/18362)

* cuda: optimize cumsum cub path

* remove heavy perf test

ggml-cuda: fix blackwell native builds (llama/18361)

* ggml-cuda: fix blackwell native builds

Replace 12x in native architectures by 12xa

* replace for GGML_NATIVE=OFF too

* only replace for native

* remove 120f-virtual for default compilation

---------

Co-authored-by: Aman Gupta <aman>

CANN: Add support for CONV_TRANSPOSE_1D when kernel size > 255 (llama/17934)

* CONV_TRANSPOSE_1D kernel_size>255

* remove condition check

* fix the bug of type conversion

* removing trailing whitespaces

* fix: return true in the switch case

ggml : optimize cuda cumsum fallback kernel (llama/18343)

CUDA: experimental native mxfp4 support for blackwell (llama/17906)

* CUDA: experimental native mxfp4 support for blackwell

* optimize load_tiles

* optimize quantize_mxfp4

* cleanup

* first pass review: formatting

* use interleaved layout for mma

* mmq: add assert for size

* use __nv_fp4x4_e2m1

* use iter_k as 512, cleanup

* Use 1200 as blackwell instead of 1000

* address review comments

* mmq: fix stride

* quantize.cu: use reference impl of e8m0 scale

* address review comments

* add 120f-virtual + minor fixes

---------

Co-authored-by: Aman Gupta <aman>

vulkan: fix command buffer corruption in ggml_backend_vk_event_wait (llama/18302)

CANN : refactor ACL graph cache (llama/17752)

Move the graph property checking code into methods of LRU cache.

Signed-off-by: Wang Weixuan <redacted>

vulkan: use fewer FA rows for small cache runs (llama/18280)

CANN: Uses yarn_ramp cache in ROPE (llama/17725)

rpc : add check for rpc buffer type (llama/18242)

ggml-hexagon: create generalized functions for cpu side op (llama/17500)

* refactor: replace ggml_hexagon_mul_mat with template-based binary operation for improved flexibility

* refactor: replace ggml_hexagon_mul_mat_id with template-based binary operation for improved flexibility

* refactor: initialize buffer types and streamline dspqueue_buffers_init calls for clarity

* add comment

* refactor: remove redundant buffer checks in hexagon supported operations

* wip

* add missing include to fix weak symbol warning

* add ggml_hexagon_op_generic

* refactor: simplify tensor operation initialization and buffer management in hexagon implementation

* refactor: streamline hexagon operation initialization and buffer management

* refactor: update function signatures and streamline request handling in hexagon operations

* wip

* ggml-hexagon: clean up code formatting and improve unary operation handling

* wip

* rename

* fix: add support for permuted F16 tensors and enhance quantization checks in matrix operations

* refactor: replace ggml_hexagon_mul_mat with template-based binary operation for improved flexibility

refactor: replace ggml_hexagon_mul_mat_id with template-based binary operation for improved flexibility

refactor: initialize buffer types and streamline dspqueue_buffers_init calls for clarity

refactor: remove redundant buffer checks in hexagon supported operations

add missing include to fix weak symbol warning

add ggml_hexagon_op_generic

refactor: simplify tensor operation initialization and buffer management in hexagon implementation

refactor: streamline hexagon operation initialization and buffer management

refactor: update function signatures and streamline request handling in hexagon operations

ggml-hexagon: clean up code formatting and improve unary operation handling

fix: add support for permuted F16 tensors and enhance quantization checks in matrix operations

# Conflicts:
# ggml/src/ggml-hexagon/ggml-hexagon.cpp

* hexagon: fix merge conflicts

* hexagon: minor cleanup for buffer support checks

* hexagon: factor out op_desc and the overal op logging

* hexagon: further simplify and cleanup op dispatch logic

* snapdragon: update adb scripts to use llama-cli and llama-completion

* fix pipeline failure

---------

Co-authored-by: Max Krasnyansky <redacted>

ggml-hexagon: gelu optimization (llama/18151)

* feat: working gelu with src0 put on vtcm

* feat: gelu ping-pong for both in and out

* fix: fixu compile error

* break: distinguish dma ddr->vtcm and vtcm->ddr operation

* fix: fix dma queue size

* break: update dma api to either pop src or dst ptr

* fix: fix activation vtcm allocation issue for src1 when swapperd

* refactor: ping-pong gelu logic to avoid unnecessary if else

* dma: improved queue interface and prefetch handling

* gelu: fix N+2 block prefetch

---------

Co-authored-by: Max Krasnyansky <redacted>

llamafile: add rvv support for sgemm kernels (llama/18199)

Co-authored-by: Rehan Qasim <redacted>

opencl: unpack q4_0 for adreno in get_tensor (llama/18278)

vulkan: Extend rope fusions to allow mrope (llama/18264)

Extend the test-backend-ops tests as well.

vulkan: Implement set_tensor_async and the event interfaces (llama/18047)

The goal is to enable the async loading code paths in
llama_model_loader::load_all_data, originally from #7896. This works and the
loads themselves are faster, but with host visible vidmem I think the cost of
allocating/mapping vidmem moves and becomes more expensive, and I don't see a
benefit by default. But with GGML_VK_DISABLE_HOST_VISIBLE_VIDMEM=1 I do see a
significant improvement in model loading time.

llama: fix RPC for -fit on (llama/18233)

vulkan: fix im2col overflowing maxworkgroupcount (llama/18180)

vulkan/cuda: fix topk_moe with exp_probs_b (llama/18071)

I updated test_topk_moe to more closely match llm_graph_context::build_moe_ffn
and added coverage for exp_probs_b and some other missing combinations. This
exposed a bug in both CUDA and Vulkan backends where they were assuming the
input to argsort and the input to get_rows are the same. I'd like to optimize
this graph in another change, but for now just get it functional.

CUDA also had a bug where it got n_experts from the wrong place, leading to
GGML_ASSERT failures in some of the new tests.

vulkan: support GGML_UNARY_OP_XIELU (llama/18062)

vulkan: in graph_optimize, try to group ADD operations (llama/18060)

I saw the adds not staying together in the new nemotron 3 nano model.

Vulkan: some improvement on mul_mat_iq2_xs (llama/18031)

* Some improvement on mul_mat_iq2_xs

Refactor calculations for db values and grid data to optimize performance and reduce redundancy.

* Fix trailing whitespace

tests: Avoid floating point precision false positives in SUM (llama/17471)

* tests: Avoid floating point precision false positives in SUM

* also apply to test_mean

test-backend-ops: improve msvc build time (llama/18209)

Added comments explaining thread block size selection logic based on row count and column size, derived from historical commit context (llama/18212)

ggml-hexagon: Implement true Q8_0 quantization on Hexagon NPU for more accurate mixed-precision matmul operations (llama/17977)

* feat: implement real Q8_0

* feat: adding cmake option for configuring FP32 quantize group size

* typo: set() shall be used

---------

Co-authored-by: ngdxzy <redacted>

vulkan: Add perf logger mode with concurrency (llama/17944)

This implements a variation of the perf logger where rather than timing each
operation individually with effectively a barrier in between, we put the
timing boundaries where we already synchronize and time the groups of work
that normally overlap. This can be useful to help understand whether
individual operations need to be optimized, or if the group is already running
efficiently.

GGML_VK_PERF_LOGGER_CONCURRENT=1 enables the new mode (when
GGML_VK_PERF_LOGGER is also set).

GGML_VK_SYNC_LOGGER=1 replaces the ENABLE_SYNC_LOGGING compile time switch.

model : add ASR support for LFM2-Audio-1.5B (conformer) (llama/18106)

* ASR with LFM2-Audio-1.5B

* Set rope_theta

* Fix comment

* Remove rope_theta setting

* Address PR feedback

* rename functions to conformer

* remove some redundant ggml_cont

* fix missing tensor

* add prefix "a." for conv tensors

* remove redundant reshape

* clean up

* add test model

---------

Co-authored-by: Tarek Dakhran <redacted>

ggml-cpu: extend support for RVV floating-point kernels (llama/17318)

* cmake: add BF16 RVV flag for ggml-cpu

* ggml-cpu: add floating-point conversion kernels

* ggml: add floating-point kernels

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: fix lmul in vec_dot_bf16

* ggml-cpu: change redsum to lmul 4, fix leftover

---------

Co-authored-by: Rehan Qasim <redacted>

remove i_major_dual (llama/18157)

Co-authored-by: zhang hui <redacted>

ggml-hexagon: swiglu_oai operation (llama/18114)

* snapshot: debug ggml-hexagon swiglu-oai

* fix: fix hvx_min_scalar_f32

* feat: working swiglu-oai

* chore: fix formating isue

ggml-hexagon: gelu operation (llama/17921)

* feat: inital support for gelu using sigmoid approximation

* snapshot: faster gelu using polynomial approximation

* test: disable l2-block prefetch in polynomail approximation

* Revert "test: disable l2-block prefetch in polynomail approximation"

This reverts commit 72339994d45b2bed887e79994403c378d90b62b5.

* Revert "snapshot: faster gelu using polynomial approximation"

This reverts commit 2a787a61d11f9e63e5943a2e6d134b2f0c402ace.

* debug: temporarily disable unnecessary log message for debug purpose

* Feat: optiized unaligned sigmoid_f32

* Feat: larger l2prefetch block

* feat: apply unaligned-load optimization on mul and mul_scalar

* Revert "debug: temporarily disable unnecessary log message for debug purpose"

This reverts commit 84f2f23aa9f17e2fa826db969cd825d0ab192995.

* refactor: cleanup commented unused code

* chore: reformat code with clang-formatter to pass cli test

* Revert "chore: reformat code with clang-formatter to pass cli test"

This reverts commit 952877ec24732b12010c7fa7ed3fc8de4b74e718.

* fix: fix loop overflow

* chore: fix formating ci error

ggml-cpu: ARM64: repack version of q8_0 (dotprod and i8mm) (llama/18096)

* wip: skeleton for q8_0 repack

* q8_0 repack GEMV implementations

* GEMM implementations

* Formatting

* Fixed format consistency of repack gemm and gemv declarations

* gemv and gemm generic location consistent with declarations

* Removed non-correct unused variables statements

* Cleanup, consistent style

* Missing generic fallbacks for x86 and powerpc

HIP: Refactor mma for RDNA and CDNA (llama/17990)

* mma.cuh for rdna4

* mma for rdna3

* mmq for rdna4

* mmq for rdna3

* align i-major and j-major

* cdna

* fix cuda error

* add missing tile of mfma

* fix j-major wrong ne on CDNA

* fix gramma and empty spaces

---------

Co-authored-by: zhang hui <redacted>

Improve arm64 optimization

Improve arm64 optimization

Disable arm64 optimizations

Disable hardening

Re-enable ARM optimizations

Reduce optimizations on arm64

Reduce optimizations for arm64

Disable hardening

Revert to GGML_NATIVE=ON on Ubuntu arm64

Use default CUDA flags

Build Vulkan backend only on amd64

Fix ARM flags

Build Vulkan backend only on amd64

Fix ARM flags

Explicitly specify ARM optimizations

Update upstream

Merge tag 'upstream/0.9.4.395' into debian/latest

Upstream pinned commit ac0c8be4