sycl : support to malloc memory on device more than 4GB, update the doc and script (llama/17566)

Co-authored-by: Neo Zhang Jianyu <redacted>

ggml: replace hwcap with riscv_hwprobe for RVV detection (llama/17567)

Signed-off-by: Wang Yang <redacted>

Vulkan: MMVQ Integer Dot K-Quant and MUL_MAT_ID support (llama/16900)

* vulkan: split mul_mmq_funcs for mul_mat_vecq use

* add mxfp4 mmvq

* add q2_k mmvq

* add q3_k mmvq

* add q4_k and q5_k mmvq

* add q6_k mmvq

* handle 4x4 quants per mmvq thread

* enable MUL_MAT_ID mmvq support

* enable subgroup optimizations for mul_mat_vec_id shaders

* device tuning

* request prealloc_y sync after quantization

* fix indentation

* fix llvmpipe test failures

* fix mul_mat_id mmvq condition

* fix unused variable warning

vulkan: improve topk perf for large k, fix overflow in unit tests (llama/17582)

ggml : add GGML_SCHED_NO_REALLOC option to disable reallocations in ggml_backend_sched (llama/17276)

* ggml : add GGML_SCHED_NO_REALLOC option to disable reallocations in ggml_backend_sched
Enabled in ggml-ci for testing.

* llama : update worst-case graph for unified cache

* ci : disable op offload in some tests

* fix spelling

---------

Co-authored-by: Georgi Gerganov <redacted>

enable fp16/fast_fp16/bf16_mma on PH1 (llama/17551)

* [MUSA] enable fp16/fast_fp16/bf16_mma on PH1

Signed-off-by: Xiaodong Ye <redacted>
* Update ggml/src/ggml-cuda/fattn-vec.cuh

Co-authored-by: Johannes Gäßler <redacted>
* Update ggml/src/ggml-cuda/fattn-vec.cuh

Co-authored-by: Johannes Gäßler <redacted>
* Update ggml/src/ggml-cuda/fattn-tile.cuh

Co-authored-by: Johannes Gäßler <redacted>
* Address review comments

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

Signed-off-by: Xiaodong Ye <redacted>
Co-authored-by: Johannes Gäßler <redacted>

ggml-cuda: add stricter checking for fusion (llama/17568)

* ggml-cuda: make conditions for fusion more explicit

* ggml-cuda: remove size check as std::equal already does it

model : Qwen3 Next (llama/16095)

* Qwen3 Next - cleaned up version

* Whitespaces and stuff

* Correct minor errors

* Update src/llama-model.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
* Misc. fixes.

* Clean up code, add missing hybrid qualifier

* Did someone transpose the SOLVE_TRI result matrix? Perhaps...

* Whitespace

* Proper tensors for cb calls

* Use llama-graph.h vertical alignment

* BROKEN: chunking

* Set new tensors as inputs.

* Proper chunk logic

* It's the circle of life...

* More shenanigans for n_seq > 1

* Nail in the coffin?

* Fix Windows build

* Eh, one fails on Windows, the other fails on Mac... just use general capture.

* quant : cleanup

* model : cleanup

* qwen3 : cleanup

* cont : cleanup

* cont : cleanup

* ggml : revert change

* qwen3 : cleanup

* cont : cleanup

* Readd cmath

* qwen3 : fix typo

* Update convert_hf_to_gguf.py

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

* fix my bad suggestion

---------

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

CUDA: no FP16 arithmetic for vector FA kernel (llama/17558)

vulkan: Implement GGML_OP_TRI (llama/17503)

* vulkan: Implement GGML_OP_TRI

* check types match

rpc : cache and reuse compute graphs (llama/15405)

Store the last computed graph and reuse it when possible.
Also do not return response from GRAPH_COMPUTE and assume it always
completes successfully. If this this is not the case, the server closes
the connection. This saves us a network round trip to the server.

HIP: enable mul_mat_f for RDNA4 (llama/17437)

* enable mmf for rdna4

* move some mmvf to mmf

* revert lds128 for wmma loading

* Revert "revert lds128 for wmma loading"

This reverts commit db9ae8b6b4738a5def5b393caa1611d52133e9b5.

* Revert "enable mmf for rdna4"

This reverts commit 698c9f24187b990e35c3b73a8067e5387e6ddbd4.

* Revert "move some mmvf to mmf"

This reverts commit 99b92bd6653cc8593607f641e44606391691792f.

* enable mul_mat for rdna4

---------

Co-authored-by: zhang hui <redacted>

SOLVE_TRI CUDA kernel for small matrices (llama/17457)

refactor pad_reflect_1d to make the UT case pass (llama/17204)

Co-authored-by: Zhang Jianyu <redacted>

vulkan: Implement SOLVE_TRI (llama/17486)

* vulkan: Implement SOLVE_TRI

* load B matrix through shared memory

* use FLOAT_TYPE

cuda : fix UMA detection on discrete GPUs. (llama/17537)

ggml-cpu: aarm64: q4_K repack gemm and gemv implementations (dotprod only) (llama/17494)

* Enabled q4_K_4x8 path

* Fixed generic Q4_K 8x4 implementation

* wip: dotprod gemm

* Working arm q4_K dotprod gemm

Signed-off-by: Alberto Cabrera <redacted>
* Undo acc rename

Signed-off-by: Alberto Cabrera <redacted>
* Q4_K arm dotprod gemm

Signed-off-by: Alberto Cabrera <redacted>
* Fix: q4_qs reinterpret from uint to int

Signed-off-by: Alberto Cabrera <redacted>
* Removed comments

* Fixed macro guards

* Fixed unused vars in generic implementation

* Fixed unused vars in 8x4 repack

* Fixed unused vars in generic implementation, unneeded comment

* Missing arch fallback for x86

* minor : style

---------

Signed-off-by: Alberto Cabrera <redacted>
Co-authored-by: Georgi Gerganov <redacted>

vulkan : move contiguous checks to device_supports_op (llama/17490)

* vulkan : remove op_supports_incontiguous and add missing constraints in device_supports_op

* im2col: remove contraints on src0 (kernel input)

vulkan: use a fixed 1KB buffer for the add_rms_fusion opt (llama/17514)

opencl: add sqr, sqrt, mean and ssm_conv (llama/17476)

* opencl: add sqr

* opencl: add sqrt

* opencl: add mean

* opencl: add ssm_conv

* opencl: add missing cl_khr_fp16

* opencl: do sqrt in f32 then convert to f16 for better precision

Fix chunks being too small with small matrix sizes (llama/17526)

vulkan: allow graph_optimize for prompt processing workloads (llama/17475)

vulkan: Implement top-k (llama/17418)

* vulkan: Implement top-k

Each pass launches workgroups that each sort 2^N elements (where N is usually 7-10)
and discards all but the top K. Repeat until only K are left. And there's a fast
path when K==1 to just find the max value rather than sorting.

* fix pipeline selection

* vulkan: Add N-ary search algorithm for topk

* microoptimizations

ggml-cpu : add RISC-V Zvfh impl for ggml_vec_mad_f16 (llama/17448)

* ggml-cpu : add RISC-V Zvfh impl for ggml_vec_mad_f16

* ggml-cpu : dedup scalar impl

* Update ggml/src/ggml-cpu/vec.h

---------

Co-authored-by: Georgi Gerganov <redacted>

ggml : fix ARM feature verification (llama/17519)

On arm64 with `cmake` version 3.31.6, the final feature verification fails:

    -- ARM detected flags: -mcpu=neoverse-v2+crc+sve2-aes+sve2-sha3+nossbs
    -- Performing Test GGML_MACHINE_SUPPORTS_dotprod
    -- Performing Test GGML_MACHINE_SUPPORTS_dotprod - Success
    -- Performing Test GGML_MACHINE_SUPPORTS_i8mm
    -- Performing Test GGML_MACHINE_SUPPORTS_i8mm - Success
    -- Performing Test GGML_MACHINE_SUPPORTS_sve
    -- Performing Test GGML_MACHINE_SUPPORTS_sve - Success
    -- Performing Test GGML_MACHINE_SUPPORTS_sme
    -- Performing Test GGML_MACHINE_SUPPORTS_sme - Failed
    -- Performing Test GGML_MACHINE_SUPPORTS_nosme
    -- Performing Test GGML_MACHINE_SUPPORTS_nosme - Success
    -- Checking for ARM features using flags:
    --   -U__ARM_FEATURE_SME
    --   -mcpu=neoverse-v2+crc+sve2-aes+sve2-sha3+nossbs+dotprod+i8mm+sve+nosme
    -- Performing Test HAVE_DOTPROD
    -- Performing Test HAVE_DOTPROD - Failed
    -- Performing Test HAVE_SVE
    -- Performing Test HAVE_SVE - Failed
    -- Performing Test HAVE_MATMUL_INT8
    -- Performing Test HAVE_MATMUL_INT8 - Failed
    -- Performing Test HAVE_FMA
    -- Performing Test HAVE_FMA - Success
    -- Performing Test HAVE_FP16_VECTOR_ARITHMETIC
    -- Performing Test HAVE_FP16_VECTOR_ARITHMETIC - Failed
    -- Performing Test HAVE_SME
    -- Performing Test HAVE_SME - Failed
    -- Adding CPU backend variant ggml-cpu: -U__ARM_FEATURE_SME;-mcpu=neoverse-v2+crc+sve2-aes+sve2-sha3+nossbs+dotprod+i8mm+sve+nosme

We need to explicitly replace `;` with spaces from the list to make
`CMAKE_REQUIRED_FLAGS` work correctly...

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

HIP: Patch failed testcase in WMMA-MMQ kernels for RDNA 4 (llama/17502)

* patch failed test case MUL_MAT(type_a=q4_0,type_b=f32,m=576,n=512,k=576,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1) for enabling WMMA on RDNA4

* Quick clean up on mma.cuh to add ggml_cuda_memcpy_1 back in for half2 and bfloat162

CANN: Add MROPE and IMROPE support (llama/17401)

* CANN: ROPE supports both MROPE and IMROPE.

1. Optimize the caching logic of rope_cache_init.
2. Add support for mRoPE and i-mRoPE.

Note that on Ascend 910B devices, it is necessary to disable FA
in CLIP and disable NZ-format conversion. These two issues are
still under investigation.

* Resolve review comments

vulkan: Implement GGML_OP_CUMSUM (llama/17479)

ggml : add ggml_top_k (llama/17365)

* ggml : add ggml_top_k

* cont : add ggml_argsort_top_k

* metal : add top_k support

* ggml : cleanup

* tests : add virtual err() function for test_case

* ggml : add comments

CANN: supports out_prod operator for F32 and F16 (llama/17406)

Co-authored-by: tianhao <redacted>

vulkan: Use fewer rows for scalar FA when HS is not a multiple of 16 (llama/17455)

vulkan: more FA details in vk_perf_logger (llama/17443)

HIP: WMMA-MMQ kernels for RDNA 4 (llama/17156)

* first commit naive test to enable mmq for RDNA4

* adding appropriate WMMA instructions

* git rebase on top of master: fixing the correctness of the mat mul operations, updating layout mappings for RDNA4

* clean up merge conflicts

* add comments and code clean up

* PR clean up, addressed comments

* enable MMQ fallback on RDNA4

* addressed comments: add guards in load generic, separate wmma branch for use_mmq function

* Revert build-xcframework.sh

* Formating: remove trailing whitespace

* revert CMake files

* clean up after rebase: remove duplicated change, revert cmake files

* clean up after rebase: revert changes from build-xcframework.sh

* clean up: remove extra space line in mma.cuh

* Revert "clean up: remove extra space line in mma.cuh"

This reverts commit b39ed57c4529906466bd0bc7c2a86e08fc2f8bee.

ggml-cpu: arm64: q4_K repack gemm and gemv implementations (i8mm) (llama/16739)

* Enabled q4_K_8x8_q8_K path on ARM

* wip: I8mm qs multiplication, pending bias

* cpu : arm : REPACK gemm q4_K8x8 implementation

Signed-off-by: Alberto Cabrera <redacted>
* Guard gemm with proper features, improved superblock scale and min calc

Signed-off-by: Alberto Cabrera <redacted>
* cpu: arm: Implemented REPACK gemv for Q4_K

Signed-off-by: Alberto Cabrera <redacted>
* Removed completed TODO

* Fixed missing guards when selecting optimal repack type for Q4_K

Signed-off-by: Alberto Cabrera <redacted>
* Fixed macro guard for gemv

* Fixed wrong comment in GEMV

* Fixed warning for unused variable

* vdotq_s32 -> ggml_vdotq_s32

Signed-off-by: Alberto Cabrera <redacted>
* Clang-format issues

* Apply suggestions from code review

Co-authored-by: Diego Devesa <redacted>
* Removed unnecessary GGML_UNUSED

* Fixed guards in q4_k gemm and gemv (repack)

---------

Signed-off-by: Alberto Cabrera <redacted>
Co-authored-by: Diego Devesa <redacted>

ggml: add RISC-V cpu-feats (llama/17461)

* ggml: add RISC-V cpu-feats

Signed-off-by: Wang Yang <redacted>
* fix comment[1]

---------

Signed-off-by: Wang Yang <redacted>

hexagon: add support for ROPE_NEOX (llama/17458)

CANN: Define `cann_graph_update_required` before macro (llama/17434)

**Description of the problem**

`cann_graph_update_required` is redundantly defined and
initialized as `false` inside two mutually exclusive macro branches.

**Proposed solution**

Define it right before the macro so that it could serve both
branches.

ggml-hexagon: Initial Hexagon v68/v69 support (llama/17394)

* ggml-hexagon: fix build error with GCC

Add stdexcept include to fix GCC build errors

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: check VTCM acquire failures

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: disable destination bypass on older than v73

v68 errors out if having bypass enabled when the VTCM is the destination.

At least on v68 this made things actually work... not a proper fix though, so to look at later...

Signed-off-by: Mohamed Mediouni <redacted>
* ggml-hexagon: add initial v68/v69 support

v68 is the Hexagon revision notably used on the Snapdragon 8cx
Gen 3 and the QCM6490.

Also add support for v69.

8MB isn't a supported page size, so relax asked for page size constraint
for HAP_compute_res_attr_set_vtcm_param_v2 to optimal.

Signed-off-by: Mohamed Mediouni <redacted>
---------

Signed-off-by: Mohamed Mediouni <redacted>

ggml-hexagon: add `hex_supported_buffer` for better buffer supported check (llama/17212)

* hexagon: add buffer support checks for hexagon sessions

* refactor: simplify buffer support checks in hexagon operations

* hexagon: update buffer support checks to use tensor structure

* refactor: streamline buffer initialization for DSP queue in hexagon operations

* refactor: simplify buffer initialization in DSP queue for hexagon operations

* refactor: optimize hex_supported_buffer function by fold expression

* wip

* refactor: simplify dspqueue_buffers_init function and its usage in hexagon operations

* fix: improve nan handling at hvx_vec_fast_sigmoid_fp32_guard

* refactor: optimize hvx_vec_inverse_fp32_guard for better nan handling

* refactor: update hvx_vec_fast_sigmoid_fp32_guard to use adjusted exponent limits

* refactor: modify hvx_vec_fast_sigmoid_fp32_guard to accept parameters for improved flexibility

* refactor: update hvx_vec_exp_fp32_guard to accept max_exp and inf parameters to save some instructions

* refactor: move hvx_vec_inverse_fp32_guard implementation to hvx-inverse.c for better perf

cuda : support non-contiguous i32 to i32 copy (llama/17326)

* support non-contiguous i32 to i32 copy

* add tests

* rename cpy_flt to cpy_scalar and reindent params

vulkan: remove a couple unnecessary switches (llama/17419)

Revive MUL_MAT_ID to perf testing (llama/17397)

HIP: RDNA4 tensor core support for MMF (llama/17077)

* mmf for rdna4

* align the padding for rdna4

* forbit mul_mat_f for rdna4

* fix as comment

* remove device kernels

* add constexpr for early return

* update based on review comment

* change based on the review comment

* pass compile error

* keep code consistency

---------

Co-authored-by: zhang hui <redacted>

opencl: refine condition for kqv mm (llama/17392)

vulkan: disable async for older Intel devices (llama/17369)

* vulkan: disable async for older Intel devices

* update detection logic

* use name string for detection

CANN: Refactor `evaluate_and_capture_cann_graph` (llama/17333)

* CANN: Refactor `evaluate_and_capture_cann_graph`

**Description of the problem**

* `matched_graph` is obtained even if graph mode is disabled.
* End of graph capture and graph replay are unnecessarily placed in different `if` blocks.

**Proposed solution**

* Obtain `matched_graph` only if graph mode is enabled.
* Place end of graph capture and graph reply inside the same `if` block.
* Unify graph related comments.

* Remove trailing whitespace

ggml-hexagon: fix swiglu failure at `test-backend-ops` (llama/17344)

* refactor: use hvx_vec_exp_fp32_guard_inf for overflow handling in hvx_exp_f32

* feat: add fast sigmoid function with overflow guard for fp32

* refactor: replace hvx_vec_inverse_fp32 with hvx_vec_inverse_fp32_guard_inf for improved overflow handling

* feat: enhance hvx_add_scalar_f32 with overflow handling using infinity guard

* wip

* add HVX_Vector_Alias

wip

* wip

* fix: improve handling of src1 tensor in glu_swiglu_fp32_per_thread function

* fix nc

* wip

* wip

* handle nan at inverse

* wip

* fix neg

* wip

* rename

* fix hvx_vec_inverse_fp32_guard_inf to handle infinity and NaN cases correctly

* wip

* fix hvx_vec_inverse_fp32_guard_inf to handle NaN cases correctly

* wip

* wip

* wip

* fix output sign

ggml : Fix transposed SOLVE_TRI result (llama/17323)

* Did someone transpose the SOLVE_TRI result matrix? Perhaps...

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

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

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

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

DGX Spark: UMA support (llama/17368)

* DGX Spark: UMA support

* Updates from PR feedback

* More PR feedback cleanup

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

Co-authored-by: Georgi Gerganov <redacted>
* Remove trailing whitespace

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

---------

Co-authored-by: Georgi Gerganov <redacted>

ggml : remove useless and error-prone variadic macros (llama/17399)

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

kleidiai: fix zero-size array declaration (llama/17240)

ggml-cpu:add RISC-V RVV (Zvfh) optimization for FP16 vector scaling (llama/17314)

* ggml-cpu:add RISC-V RVV (Zvfh) optimization for FP16 vector scaling

Signed-off-by: Wang Yang <redacted>
* fix comment

* fix comment 2

---------

Signed-off-by: Wang Yang <redacted>

vulkan: implement ADD1, ARANGE, FILL, SOFTPLUS, STEP, ROUND, CEIL, FLOOR, TRUNC (llama/17319)

* vulkan: initialize array

* vulkan: implement ADD1

* vulkan: implement ARANGE

* vulkan: implement FILL

* vulkan: implement SOFTPLUS

* vulkan: implement STEP

* vulkan: implement ROUND

* vulkan: implement CEIL

* vulkan: implement FLOOR

* vulkan: implement TRUNC

* docs: update Vulkan ops

Signed-off-by: Giuseppe Scrivano <redacted>

vulkan: support larger argsort (llama/17313)

* vulkan: support larger argsort

This is an extension of the original bitonic sorting shader that puts the
temporary values in global memory and when more than 1024 threads are needed
it runs multiple workgroups and synchronizes through a pipelinebarrier.

To improve the memory access pattern, a copy of the float value is kept with
the index value. I've applied this same change to the original shared memory
version of the shader, which is still used when ncols <= 1024.

* Reduce the number of shader variants. Use smaller workgroups when doing a single pass, for a modest perf boost

* reduce loop overhead

* run multiple cols per invocation, to reduce barrier overhead

vulkan: Add copy_transpose shader (llama/17371)

cuda: fix rope fusion for gemma3 (llama/17378)

Fix too relaxed check on CUDA "fast copy" (can_be_transposed) condition (llama/17332)

* Fix too relaxed check on CUDA "fast copy" (can_be_transposed) condition

* Argh.

* Making CISC happy ;)

* Integrate CONT tests

* Use loopy loop

* Skip new tests for (B)F16 for now.

vulkan: force full subgroups for flash attention to fix intel subgroup crash (llama/17356)

ggml-cpu: Don't pass -mpowerpc64 when -mcpu already implies it (llama/17308)

CANN: fix acl_tensor_ptr usage in ASCEND_310P ROPE (llama/17347)

* cann: fix acl_tensor_ptr usage in ASCEND_310P ROPE implementation

Fix compilation errors in the ASCEND_310P-specific ROPE operation code
by adding .get() calls when passing acl_tensor_ptr smart pointers to
functions expecting raw aclTensor* pointers.

This fixes the code that was missed in the previous refactoring commit
(8981848) which changed ggml_cann_create_tensor() return type from
aclTensor* to acl_tensor_ptr.

* cann: format code

vulkan: support noncontig i32 copy (llama/17328)

vulkan: add log RTE support to fix Nvidia CI (llama/17320)

* vulkan: add log RTE support to fix Nvidia CI

* actually use the rte shader

cmake : fix ARM feature verification (llama/17170)

* cmake : fix ARM feature verification

Use check_cxx_source_compiles to prevent conflicts with
the existing GGML_NATIVE detection code.

Signed-off-by: Adrien Gallouët <redacted>
* cmake : unset __ARM_FEATURE when feature is disabled

Signed-off-by: Adrien Gallouët <redacted>
* cmake : fix scope, this is really a macro

Signed-off-by: Adrien Gallouët <redacted>
* arm_neon.h is useless

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

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

ggml : add missing AVX512 feature checks (llama/17270)

_mm512_cvtepu8_epi16        requires  __AVX512BW__
_mm512_srli_epi16           requires  __AVX512BW__
__builtin_ia32_inserti32x8  requires  __AVX512DQ__

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

ggml : remove dirty flag from version string (#1391)

This commit removes the "-dirty" suffix from the GGML version string.

The motivation for this change is to ensure that the version string
works with different ways of checking out ggml and using it in projects.
By removing the dirty flag from the version string, we avoid potential
artifacts like shared libraries getting a -dirty suffix in their names.

Instead, if the project is built from a dirty git state, the dirty flag
will be appended to the commit hash in the GGML_BUILD_COMMIT variable.
This will enable users to still identify that the build was made from
from a modified/dirty state even though the version might match a "real"
version.

For example, the commit can be produces as follows:
```c++
    printf("commit: %s\n", ggml_commit());
```
Which would print the following for a dirty build:
```console
commit: 781baf2a-dirty
```

Refs: https://github.com/ggml-org/ggml/pull/1363#issuecomment-3569691546

sync : whisper.cpp

metal : fix compile on macos 11 (whisper/3533)

sync : llama.cpp

metal : support I32 -> I32 copy (llama/17317)

metal : faster argsort (llama/17315)

* metal : faster argsort

* cont : keep data in registers

metal : add cumsum (llama/17305)

CANN: Use smart pointers to manage ACL objects (llama/17238)

* CANN: Use smart pointers to manage ACL objects

Previously, ACL objects were managed via manual destruction, which
led to multiple memory-leak issues during runtime. This patch replaces
manual memory management with smart pointers so that ACL objects
are properly released and ownership is clearly defined.

Note that the ownership of an ACL object belongs to the function
that creates it. Other internal functions should operate on these ACL
objects using raw pointers to avoid unintended ownership transfers.

Additionally, since aclTensorList automatically frees its contained
aclTensor objects, any aclTensor added to a tensor list must release
ownership to avoid double free operations.

This PR also removes the asynchronous task submission mechanism.
Due to changes in recent CANN versions, tiling time has significantly
decreased. Even with a dual-thread submission model, the dispatch
overhead still falls on the critical path, making async submission
less beneficial. Moreover, aclGraph support provides a much better
path to reducing operator dispatch latency.

* CANN: resolve review comments

vulkan: add LOG operation support for F32 and F16 (llama/17183)

* vulkan: add LOG operation support for F32 and F16

Part of #14909.

* vulkan: Fix LOG operation types

* docs: Update operation support documentation for Vulkan LOG operation

* vulkan: fix log_f16 shader

* docs: restore missing LOG test cases and regenerate ops.md

vulkan: fix MMQ quantize_y condition (llama/17301)

metal : remove obosolete asserts (llama/17295)

opencl: fix rms_norm_mul (llama/17250)

* opencl: use subgrroup reduce for reduction in rms_norm_mul

* opencl: add comment about workgroup size

opencl: add kernel to handle mat mul in attention to improve encoding speed (llama/17181)

* Add mul_mm_f16_f32_kq_kqv kernel

* Add ggml_cl_mul_mat_kq_kqv_adreno func

* fix whitespace

* remove unused variable

* remove redundant

* refactor and clean up

* remove trailing whitespace

sycl : unify unary kernels with a generic implementation and enable wide operator support (llama/17213)

* SYCL: add generic unary op implementation for multiple ops (ABS/SGN/…); unify non-contiguous access

* SYCL: update documentation and sycl.csv to reflect new unary op support

* update ops.md after syncing SYCL.csv changes

* Fix SYCL.csv merge conflict

* Update ops.md after fixing SYCL.csv conflicts

* Fix SYCL.csv tail after merge conflict and regenerate ops.md

* Fix line endings and final newline in SYCL.csv

* Remove TOPK_MOE entries from SYCL.csv as requested

* Update ops.md after removing TOPK_MOE from SYCL.csv

* Regenerated SYCL.csv and synced ops.md with upstream

* Update ops.md using create_ops_docs.py

vulkan: Fuse mul_mat_id+add_id+mul and mul_mat+add+add. (llama/17287)

These both show up in gpt-oss. Also, cleanup the mul_mat_vec fusion code a bit.

vulkan: Replace 16-bit unpack8 calls to work around legacy Windows AMD driver bug (llama/17285)

vulkan: implement ABS and NEG (llama/17245)

* docs: update Vulkan ops

* vulkan: add NEG op

* vulkan: add ABS op

---------

Signed-off-by: Giuseppe Scrivano <redacted>

vulkan: Use ggml_vk_tensor_subbuffer in mul_mat_vec(id) paths (llama/17244)

* vulkan: Use ggml_vk_tensor_subbuffer in mul_mat_vec(id) paths

* set allow_misalign

vulkan: skip all-negative-inf blocks in FA (llama/17186)

vulkan: change graph_compute to be async and enable get_tensor_async (llama/17158)

* vulkan: change graph_compute to be async and enable get_tensor_async

This allows some additional CPU/GPU overlap for large pp workloads. Also seems
to help a bit for token gen, maybe getting rid of a small bubble between
graph_compute and get_tensor.

Async set and copy functions seem to be very rarely used, so I didn't enable
them because I didn't have a good way to test them.

The async commands need to be ordered against each other, so put them all on
the compute queue. The non-async commands still use the transfer queue.

The fence for graph_compute/get_tensor_async is submitted and waited on in
ggml_vk_synchronize.

* fix thread safety errors

* teardown context cleanly

* Handle async read to non-pinned dst

metal : support argsort for ne00 > 1024 (llama/17247)

* metal : refactor argsort

* cont : sort chunks

* cont : merge sorted buckets

* cont : cleanup

metal : make the FA extra sizes consistent (llama/17143)

ggml-cpu: handle 3d tensors in repack mat_mul (llama/17241)

* ggml-cpu: handle 3d tensors in repack mul_mat

* Removed unnecessary branch, removed need for <algorithm>

* Fixed dst_ptr pointer in chunk + clang_format

* GGML_ASSERT to check wdata within bounds

* Accidental ggml.h inclusion

* Improved GGML_ASSERT on wdata boundaries

* Address performance regression in Qwen and llama.cpp due to chunking

ggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063)

* Add ops needed for new hybrid models: SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM

* Update ggml/include/ggml.h

Co-authored-by: Georgi Gerganov <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Code review

* Whitespace

* Update tests/test-backend-ops.cpp

Co-authored-by: Diego Devesa <redacted>
* This is actually sigmoid, duh.

* Add CONST, remove TRI_KEEP, other changes from review

* Update tests/test-backend-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml.c

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml.c

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml-cuda/unary.cu

Co-authored-by: Aman Gupta <redacted>
* Remove extra script

* Update ggml/src/ggml.c

Co-authored-by: Diego Devesa <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Diego Devesa <redacted>
* moving changes from laptop [no ci]

* pre-rebase

* Update tests/test-backend-ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update tests/test-backend-ops.cpp

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

* ggml : cleanup

* cont : fix ggml_fill srcs

* tests : add note

* ggml : add ggml_fill_inplace

* ggml : add asserts

* ggml : fix ggml_fill constant cast

* cont : ggml_tri minor

* Use TENSOR_LOCALS

* Fix regression from #14596, regenerate

* Don't make commits at night...

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Diego Devesa <redacted>
Co-authored-by: Aman Gupta <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>

vulkan: remove shell call from vulkan-shaders-gen tool, revert file check (llama/17219)

* vulkan: remove shell call from vulkan-shaders-gen tool

* use string vector for command execution

* Fix condition

* use string, remove const_cast

* Fix dependency file quotation on Windows

---------

Co-authored-by: Jeff Bolz <redacted>

sched : fix reserve ignoring user tensor assignments (llama/17232)

ggml-cpu : add RISC-V vector intrinsic support for silu and cvar operations (llama/17227)

Signed-off-by: Wang Yang <redacted>

metal: accelerated conv2d (llama/17175)

* metal: accelerated conv2d

* cont : cleanup

---------

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

Revert "ggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)" (llama/17233)

This reverts commit 1c398dc9eca9c366ce98deb0e6f3538e444ebc8a.

ggml-cpu : use template for argsort (llama/17222)

CANN: Add cross_entropy_loss op support (llama/16886)

* update L2_NORM op support

* update L2_NORM op support

* remove extra whitespace

* cann: update cross_entropy_loss op support

* remove trailing whitespaces

* rebase the latest code in the main repository and remove the l2_norm operator that already exists in another pull request.

* undo the l2_norm operator deletion

CUDA: fuse rope + set_rows (llama/16884)

* CUDA: add fused rope

* move k forward_expand up

* create helper function instead of re-using params

* make assert statement more in line with comment

* rope_norm: coalesced writes to global mem

CUDA: static assert to prevent misuse of memcpy_1 (llama/17198)

ggml : use std::sort in ggml_argsort CPU implementation (llama/17211)

* ggml : use std::sort in ggml_argsort CPU implementation

* cont : add missing header

ggml-cpu: handle 3d tensors in repack mat_mul (llama/17030)

* ggml-cpu: handle 3d tensors in repack mul_mat

* Removed unnecessary branch, removed need for <algorithm>

* Fixed dst_ptr pointer in chunk + clang_format

* GGML_ASSERT to check wdata within bounds

* Accidental ggml.h inclusion

* Improved GGML_ASSERT on wdata boundaries

CANN: Add L2_NORM op support (llama/16856)

* update L2_NORM op support

* update L2_NORM op support

* remove extra whitespace