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

ggml : remove trailing whitespace (llama/0)

opencl : add GEGLU, REGLU, SWIGLU (llama/14456)

Add Conv2d for CPU (llama/14388)

* Conv2D: Add CPU version

* Half decent

* Tiled approach for F32

* remove file

* Fix tests

* Support F16 operations

* add assert about size

* Review: further formatting fixes, add assert and use CPU version of fp32->fp16

metal : disable fast-math for some cpy kernels (llama/14460)

* metal : disable fast-math for some cpy kernels

ggml-ci

* cont : disable for q4_1

ggml-ci

* cont : disable for iq4_nl

ggml-ci

ggml-cpu: sycl: Re-enable exp f16 (llama/14462)

test-backend-ops : disable llama test (llama/14461)

cmake : Remove redundant include path in CMakeLists.txt (llama/14452)

* Update docker.yml

修改docker.yml文件中的内容使其停止周期性的运行该workflow，如果想要运行该workflow可以手动启动

* Remove redundant include path in CMakeLists.txt

The parent directory '..' was removed from the include directories for the ggml-cpu-feats target, to avoid unnecessary include paths.

* Enable scheduled Docker image builds

Uncomments the workflow schedule to trigger daily Docker image rebuilds at 04:12 UTC, improving automation and keeping images up to date.

scripts : make the shell scripts cross-platform (llama/14341)

SYCL: disable faulty fp16 exp kernel (llama/14395)

* SYCL: disable faulty fp16 CPU exponent for now

* Revert "SYCL: disable faulty fp16 CPU exponent for now"

This reverts commit ed0aab1ec31b4eb4b0f275dd7acd41d96a375202.

* SYCL: disable faulty fp16 CPU exponent for now

* Fix logic of disabling exponent kernel

ggml : fix unmerged GGML_FPxx_TO_FPxx refactoring (llama/14443)

ggml : implement REGLU/GEGLU/SWIGLU ops (llama/14158)

* implement unary REGLU/GEGLU/SWIGLU cpu ops

* relax constraints

* duplicate shape of source

* fix ggml_vec_geglu_f16

* special case gated ops

* implement unary REGLU/GEGLU/SWIGLU cuda ops

* tighten constraints again

* refactor into GGML_GLU_OP

* metal : add glu kernels

ggml-ci

* add CUDA_GLU_BLOCK_SIZE [no ci]

* more constraints and use 64bit ints

ggml-ci

* 64bit multiplication [no ci]

* implement swapped variants (cpu/cuda)

* update comment [no ci]

ggml-ci

* Vulkan: Add GLU ops and shaders

* SYCL: Implement fused kernel GEGLU, SWIGLU and REGLU for single up+gate

* ggml : implement GLU for split up/gate (llama/14181)

* implement GLU for split up/gate

* add tests for ggml_glu_split

* Vulkan: Implement glu_split logic and shader support

* add split to logging [no ci]

* SYCL: refactor element_size ops and add split up and gate support to gated kernels

* SYCL: switch GEGLU to use tanh approximation

---------

Co-authored-by: 0cc4m <redacted>
Co-authored-by: Akarshan <redacted>
* GGML: increase OP count in assertion

* Refactor: Optimize SYCL element-wise operations with unary function inlining

This commit refactors the SYCL element-wise operations to improve performance by:

- Inlining unary operations (sgn, abs, elu, gelu, silu, etc.) to reduce kernel launch overhead.
- Introducing helper functions `op_xxx` for each unary operation to encapsulate the logic.
- Replacing direct kernel calls with calls to these inlined functions.
- Using `__dpct_inline__` to encourage compiler inlining.
- Minor code cleanup and consistency improvements.

The changes aim to reduce kernel launch overhead and improve the overall efficiency of element-wise operations on SYCL devices.

* vulkan: Increase workgroup size for GLU, for performance (llama/14345)

* vulkan: Increase workgroup size for GLU, for performance

* vulkan: change GLU shaders to do one element per invocation rather than one row per workgroup

* merge fix

* metal : add support for split and swap

ggml-ci

---------

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

vulkan: Add fusion support for RMS_NORM+MUL (llama/14366)

* vulkan: Add fusion support for RMS_NORM+MUL

- Add a use_count to ggml_tensor, so we can detect if an output is used more than once.
- Change the ggml-vulkan rms_norm shader to optionally multiply by another tensor.
- Add detection logic and basic fusion logic in ggml-vulkan.
- Add some testing support for fusion. Rather than computing one node at a time, allow
for computing the whole graph and just testing one node's results. Add rms_norm_mul tests
and enable a llama test.

* extract some common fusion logic

* fix -Winconsistent-missing-override

* move ggml_can_fuse to a common function

* build fix

* C and C++ versions of can_fuse

* move use count to the graph to avoid data races and double increments when used in multiple threads

* use hash table lookup to find node index

* change use_counts to be indexed by hash table slot

* minimize hash lookups

style fixes

* last node doesn't need single use.
fix type.
handle mul operands being swapped.

* remove redundant parameter

---------

Co-authored-by: slaren <redacted>

CUDA: add bf16 and f32 support to cublas_mul_mat_batched (llama/14361)

* CUDA: add bf16 and f32 support to cublas_mul_mat_batched

* Review: add type traits and make function more generic

* Review: make check more explicit, add back comments, and fix formatting

* Review: fix formatting, remove useless type conversion, fix naming for bools

vulkan: handle noncontig in the final case of ggml_vk_get_cpy_pipeline (llama/14378)

vulkan: lock accesses of pinned_memory vector (llama/14333)

fix async_mode bug (llama/14432)

vulkan: Fix GGML_VULKAN_SHADER_DEBUG_INFO (llama/14427)

This setting needs to be passed through to vulkan-shaders-gen

ggml : add ggml_set_rows (llama/14274)

* ggml : add ggml_set_rows

Add ggml_set_rows(a, b, c) which copies rows from 'b' into 'a' using
indices from 'c'.

ref: #8366

* use I64 for indices

* ggml : add repeat impl for i64

* ggml : add ggml_is_contiguous_rows

* ggml : ggml_set_rows support broadcast

* ggml : ggml_set_rows support quantized dst

ggml-ci

* ggml : support GGML_TYPE_F32 ".from_float" trait

* ggml : ggml_set_rows update comment + better index name

* tests : add ggml_set_rows

* metal : add ggml_set_rows implementation

ggml-ci

* ggml : simplify forward_dup_f32

* ggml : fix supports_op

* tests : add comment to set_rows

* ggml : leave the repeat_i64 for a separate PR

ggml-ci

* ggml : set_rows use std::min instead of MIN

* ggml : better error message for set_rows unsupported type

* metal : perform op->type check only once

* tests : more consistent implementation + more tests

ggml-ci

---------

Co-authored-by: Georgi Gerganov <redacted>

cmake: regen vulkan shaders when shaders-gen sources change (llama/14398)

* Add shaders-gen sources as target deps

metal : add special-case mat-vec mul for ne00 == 4 (llama/14385)

ggml-ci

metal : batch rows copy in a single threadgroup (llama/14384)

* metal : batch rows copy in a single threadgroup

ggml-ci

* metal : handle some edge cases when threadgroup size is not a power of 2

ggml-ci

musa: enable fp16 mma (all) and cublas on qy2 (llama/13842)

* musa: enable fp16 mma (all) and cublas on qy2

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

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

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

Signed-off-by: Xiaodong Ye <redacted>
* musa: disable MUL_MAT_ID (q2_k × f32) due to precision issues

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

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

ggml-cpu: enable IBM NNPA Vector Intrinsics (llama/14317)

* ggml-cpu: add nnpa compile flag

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

* ggml-cpu: add fp16->fp32 nnpa first

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

* ggml-cpu: add fp32->fp16

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

* ggml-cpu: better variable names

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

* docs: update s390x docs

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

* ggml-cpu: add debugging prints to see if dlf16 is correct

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix print vs printf

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

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: ensure fp16 and fp32 load and stores are called

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fp16 load ensured to hit

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove sigint from fp16 store

for some reason, the function is not getting a hit when debugged with
    gdb. we will need to investigate further

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: activate nnpa for ggml_cpu_fp16_to_fp32

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: nnpa activate ggml_cpu_fp16_to_fp32 for 8 elements

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: nnpa switch to vec_xst test

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: switch to vec_xst for 4 element loops also

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

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove noop, general code cleanup

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: clarify variable naming

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: activate nnpa for ggml_cpu_fp32_to_fp16

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add breakpoint for debugging

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: test fix for conversion failure

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: disable fp32->fp16 nnpa conversions for now

there are some conversion failures in nnpa that requires the eyes of an
ibm stsm. will create a separate pr to introduce the fp32->fp16 change.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: switch to elif macro

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: reattempt fp32->fp16

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

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: reattempt fp32->fp16

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

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: change to typedef vector types

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add 4 element loops for fp32->fp16

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: clarified vector naming

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: bring back fp32->fp16 store nnpa

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: activate nnpa fp32->fp16 or fp16->fp32 compute

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add nnpa macro check in ggml-impl

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add missing __func__

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: diagnose why __NNPA__ macro is not being defined

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: import vecintrin.h to fix compiler errors

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: update macro tests

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move s390x typedef to own header file

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: move s390x typedef to own header file"

This reverts commit 157f856c34589566151630e294563a420702db39.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: switch to importing ggml-cpu-impl instead

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

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: test more macros

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add debug prints

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: bruteforce macro definitions

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move macro definitions

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add ggml-impl.h to cmakelists

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: switch to private macros

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move s390x typedef to own header file

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

* ggml-cpu: move things around

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: bring back compile macros

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: switch to quotes for import

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add compiler error macro

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add s390x detection in ggml-src

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: bring back compile definitions

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: undo cmakelists work

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: move s390x typedef to own header file"

This reverts commit 18d79e1a30b39d9aaa0bd58400c5cf2c32135c9a.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove typedefs.h

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove typedef from cmakelists

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add ggml-impl.h future notes

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add todo comment for future reference

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: clarify naming of dlf16

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove unnecessary target compile definitions

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move nnpa fp16->fp32 and fp32->fp16 to simd-mappings

Signed-off-by: Aaron Teo <redacted>
* ggml: refactor fp32->fp16 and fp16->fp32 simd to ggml-cpu

Signed-off-by: Aaron Teo <redacted>
* docs: update broken huggingface link for s390x

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix duplicate func names during compile

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: fix duplicate func names during compile"

This reverts commit fbb733451f27677063b914d4f6c9a9841d45b38d.

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml: refactor fp32->fp16 and fp16->fp32 simd to ggml-cpu"

This reverts commit bd288e8fa52b5244f65cee21cb61062f1a9e0ca5.

Signed-off-by: Aaron Teo <redacted>
* ggml: refactor fp16<->fp32 simd to ggml-cpu

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix missing simd-mappings.h import in quants.c

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix missing simd-mappings.h within repack

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix amx mmq missing simd-mappings.h

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: attempt at fixing loongarch failing build

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move nnpa together with other fp16<->fp32 simd

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix wrong refactor of ggml-base

ref: https://github.com/ggml-org/llama.cpp/pull/14317#discussion_r2164176555

Signed-off-by: Aaron Teo <redacted>
* ggml: remove dependency on ggml-cpu from ggml-base

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: rename all fp16<->fp32 macros to prefix with ggml_cpu

ref: https://github.com/ggml-org/llama.cpp/pull/14317#discussion_r2164449406

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: remove mistaken fallback macro

fallback logic was already implemented but i was too sleepy to realise

Signed-off-by: Aaron Teo <redacted>
* ggml: move ggml_table_f32_f16 to ggml-cpu

ref: https://github.com/ggml-org/llama.cpp/pull/14317#discussion_r2164775006

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: move ggml_table_f32_f16 back to ggml-base due to ci failures

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: move ggml_table_f32_f16 back to ggml-base due to ci failures"

This reverts commit 32a3533564bdb7902cefb9c89b1c9e956a81ce29.

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml: move ggml_table_f32_f16 to ggml-cpu"

This reverts commit 9e40d984ad27d7b60392fb2b7548885201864fe4.

Signed-off-by: Aaron Teo <redacted>
* ggml: move ggml_table_f32_f16 to ggml-cpu

ref: https://github.com/ggml-org/llama.cpp/pull/14317#discussion_r2164775006

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

* ggml: move ggml_table_f32_f16 to ggml-cpu.c

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: extern c ggml_table_f32_f16 + chore docs

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: dedup ggml_table_f32_f16 from simd-mappings.h

we rely on the variable declaration in ggml-cpu.c instead

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: dedup ggml_table_f32_f16 from simd-mappings.h"

This reverts commit f71b21d2f74f5e03ec0c2b4fefd3cbf395aecf16.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: bring back ggml_table_f32_f16

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: bring back ggml_table_f32_f16"

This reverts commit 2dce119178bed5ef5c8398c4230ddd14fef80e49.

Signed-off-by: Aaron Teo <redacted>
* fix ggml time initialization

* fix f32_f16 table init

* remove extra line

---------

Signed-off-by: Aaron Teo <redacted>
Co-authored-by: slaren <redacted>

ggml : do not output unprintable characters on GGUF load failure (llama/14381)

sycl: GGML_SYCL_DISABLE_OPT on by default for all Intel Devices (llama/13973)

opencl: ref count `ggml_backend_opencl_context` and refactor profiling (llama/14254)

* Move profiling info into `ggml_backend_opencl_context`
* Add `enqueue_ndrange_kernel` to launch kernel

CUDA/HIP: optimize mmv paths taken for HIP devices (llama/14324)

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

CUDA: mul_mat_v support for batch sizes > 1 (llama/14262)

* CUDA: mul_mat_v support for batch sizes > 1

* use 64 bit math for initial offset calculation

HIP: enable vec fattn on RDNA4 (llama/14323)

CUDA: add mean operation (llama/14313)

* CUDA: add mean operation

* add back sum_rows_f32_cuda

* Review: early exit if col!=0

Add support for VK_EXT_debug_utils to add labels to Vulkan objects. (llama/13792)

* Add support for VK_EXT_debug_utils to add labels to Vulkan objects. In step 1 compute pipelines are getting labeled.

* remove #ifdef for debug utils and add queue marker.

metal : fix thread-safety (llama/14300)

ggml-ci

ggml-cpu : "align corners" for bilinear upscale/downscale (#1285)

* add "align corners" mode for bilinear upscale, and allow downscaling
* add ggml_interpolate, deprecate ggml_upscale_ext, pass in align-corners as bit-flag
* test-backend-ops: replace ggml_upscale_ext with ggml_interpolate, add test cases for downscale and align-corners

build : fix build with clang-cl on Windows (#1284)

* build : fix building tests with clang-cl on Windows

- clang-cl.exe (clang with MSVC CLI) doesn't like the space in /STACK option
- cl.exe (MSVC) works either way

* build : fix MSVC compiler warnings in test-roll.cpp

ggml-quants : rename best_mad to best_error (#1283)

This commit renames the variable `best_mad` to `best_error` in the
`make_qkx2_quants` function.

The motivation for this is that the name `best_mad` can be somewhat
confusing if mean absolute deviation (MAD) is not in use.

tests : cleanup old tests (#1282)

ggml-ci

sync : llama.cpp

ggml-ci

CUDA: add conv_2d_transpose (llama/14287)

* CUDA: add conv_2d_transpose

* remove direct include of cuda_fp16

* Review: add brackets for readability, remove ggml_set_param and add asserts

sycl: add usage of enqueue_functions extension (llama/14244)

* Add header and namespace to use enqueue_functions extension

* Convert submit and parallel_for to use new extension in convert.cpp

* Convert submit and parallel_for to use extension in ggml-sycl.cpp

* Convert submit and parallel_for to use extension in gla.cpp

* Convert submit and parallel_for in mmq.cpp

* Convert submit and parallel_for in mmvq.cpp

* Convert submit and parallel_for in remaining files

* Convert all simple parallel_for to nd_launch from enqueue_functions
extension

* Wrapping extension in general function

Create a general function that enable the enqueue_functions extension if
it is enable in the compiler, otherwise call the general SYCL function
to launch kernels.

---------

Signed-off-by: nscipione <redacted>

Implement GGML_CPU_ALL_VARIANTS for PowerPC (llama/14286)

* Add PowerPC feature detection and scoring

* ggml-cpu: Implement GGML_CPU_ALL_VARIANTS for PowerPC

* ggml-cpu: Delay some initializations until function is called

When using GGML_BACKEND_DL=ON, these initializations might use
instructions that are not supported by the current CPU.

---------

Co-authored-by: Diego Devesa <redacted>

cuda : synchronize graph capture and cublas handle destruction (llama/14288)

Workarounds an issue that may cause CUDA graph capture to fail when a cuBLAS handle is destroyed in a different thread

ggml : fix repack work size for mul_mat_id (llama/14292)

ggml-ci

ggml: Update KleidiAI to v1.9.0 (llama/14277)

CUDA: add conv_2d_dw (llama/14265)

* CUDA: add conv_2d_dw

* better naming

* simplify using template

* Review: fix operation ordering in ggml-cuda, use __forceinline__, use more const

ggml-cpu : remove unnecesary arm feature detection (llama/14281)

Support for Arm runtime feature detection has now been added to GGML_CPU_ALL_VARIANTS. This removes the old and not very functional code.

build : suppress gcc15 compile warnings (llama/14261)

* Change _contains_any() substrs to std::string_view and fix the find comparison logic.

sycl: Cleanup codepaths in Get Rows in sycl backend (llama/14215)

Addresses unused reorder path

llamafile : support s390x SIMD instruction set (llama/14273)

Vulkan: Set device max size for host memory to avoid OOM warning and fallback to CPU buffer (llama/14249)

metal : add mean kernel (llama/14267)

* metal : add mean kernel

ggml-ci

* cont : dedup implementation

ggml-ci

ggml-cpu: reduce asm calls for hsum (llama/14037)

Signed-off-by: Aaron Teo <redacted>

ggml-cpu: fix uncaught underscore terminators (llama/14023)

Signed-off-by: Aaron Teo <redacted>

ggml: Add Apple support for GGML_CPU_ALL_VARIANTS (llama/14258)

Add `ggml_roll` (#1274)

* ggml : add ggml_roll

* use set/get_op_params & std::min

sync : whisper.cpp

sync : llama.cpp

ggml-ci

cmake: remove shader-gen step-targets from ggml-vulkan (llama/14226)

* Remove step-targets from vulkan-shaders-gen

* Unset DESTDIR when building vulkan-shaders-gen

ggml-cpu : remove the weak alias trick (llama/14221)

musa: fix build warning (unused variable) (llama/14231)

Signed-off-by: Xiaodong Ye <redacted>

llama : add thread safety test (llama/14035)

* llama : add thread safety test

* llamafile : remove global state

* llama : better LLAMA_SPLIT_MODE_NONE logic

when main_gpu < 0 GPU devices are not used

---------

Co-authored-by: Georgi Gerganov <redacted>

cmake: clean up external project logic for vulkan-shaders-gen (llama/14179)

* Remove install step for vulkan-shaders-gen

* Add install step to normalize msvc with make

* Regenerate modified shaders at build-time

HIP: disable rocwmma on gfx12 by default until rocm 7.0 (llama/14202)

ggml: Add Android support for GGML_CPU_ALL_VARIANTS (llama/14206)

vulkan: mutex around vkQueueSubmit (llama/14127)

This fixes the remaining crash in test-thread-safety on my system.

ggml-cpu : rework weak alias on apple targets (llama/14146)

* ggml-cpu : rework weak alias on apple targets

* fix powerpc detection

* fix ppc detection

* fix powerpc detection on darwin

CUDA/HIP: fix ssm_scan on devices where warp size is not 32 (llama/14196)

HIP: Replace usage of depricated preprocessor macro __AMDGCN_WAVEFRONT_SIZE__ (llama/14183)

sycl: Adding additional cpy dbg print output (llama/14034)

SYCL: Bump oneMath commit (llama/14152)

Update oneMath commit to merged PR https://github.com/uxlfoundation/oneMath/pull/669
which adds SYCL-Graph support for recording CUDA BLAS commands.

With this change the `MUL_MAT` tests now pass on DPC++ CUDA backends with SYCL-Graph
enabled. Prior to this change, an error would be thrown.

```
$ GGML_SYCL_DISABLE_GRAPH=0 ./bin/test-backend-ops -b SYCL0 -o MUL_MAT -p type_a=f16,type_b=f32,m=16,n=1,k=256,bs=\\[1,1\\],nr=\\[2

UR CUDA ERROR:
        Value:           700
        Name:            CUDA_ERROR_ILLEGAL_ADDRESS
        Description:     an illegal memory access was encountered
        Function:        operator()
        Source Location: $HOME/dpcpp/unified-runtime/source/adapters/cuda/queue.cpp:154

Native API failed. Native API returns: 2147483646 (UR_RESULT_ERROR_UNKNOWN)
Exception caught at file:$HOME/llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp, line:3598, func:operator()
SYCL error: CHECK_TRY_ERROR((stream)->wait()): Meet error in this line code!
  in function ggml_backend_sycl_synchronize at $HOME/llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp:3598
$HOME/llama.cpp/ggml/src/ggml-sycl/../ggml-sycl/common.hpp:118: SYCL error
Could not attach to process.  If your uid matches the uid of the target
process, check the setting of /proc/sys/kernel/yama/ptrace_scope, or try
again as the root user.  For more details, see /etc/sysctl.d/10-ptrace.conf
ptrace: Operation not permitted.
No stack.
The program is not being run.
```

sycl: Remove not needed copy f16->f32 for dnnl mul mat (llama/14125)

cmake : handle whitepsaces in path during metal build (llama/14126)

* cmake : handle whitepsaces in path during metal build

ggml-ci

* cont : proper fix

ggml-ci

---------

Co-authored-by: Daniel Bevenius <redacted>

Implement GGML_CPU_ALL_VARIANTS for ARM (llama/14080)

* ggml-cpu: Factor out feature detection build from x86

* ggml-cpu: Add ARM feature detection and scoring

This is analogous to cpu-feats-x86.cpp. However, to detect compile-time
activation of features, we rely on GGML_USE_<FEAT> which need to be set
in cmake, instead of GGML_<FEAT> that users would set for x86.

This is because on ARM, users specify features with GGML_CPU_ARM_ARCH,
rather than with individual flags.

* ggml-cpu: Implement GGML_CPU_ALL_VARIANTS for ARM

Like x86, however to pass around arch flags within cmake, we use
GGML_INTERNAL_<FEAT> as we don't have GGML_<FEAT>.

Some features are optional, so we may need to build multiple backends
per arch version (armv8.2_1, armv8.2_2, ...), and let the scoring
function sort out which one can be used.

* ggml-cpu: Limit ARM GGML_CPU_ALL_VARIANTS to Linux for now

The other platforms will need their own specific variants.

This also fixes the bug that the the variant-building branch was always
being executed as the else-branch of GGML_NATIVE=OFF. The branch is
moved to an elseif-branch which restores the previous behavior.

vulkan: Better thread-safety for command pools/buffers (llama/14116)

This change moves the command pool/buffer tracking into a vk_command_pool
structure. There are two instances per context (for compute+transfer) and
two instances per device for operations that don't go through a context.
This should prevent separate contexts from stomping on each other.