rpc : pack only RPC structs (llama/9959)

fix mul_mat_vec_q and *_vec_q error (llama/9939)

Co-authored-by: arthw <redacted>

rpc : backend refactoring (llama/9912)

* rpc : refactor backend

Use structs for RPC request/response messages

* rpc : refactor server

Add SYCL Backend registry, device and Event Interfaces (llama/9705)

* implemented missing SYCL event APIs

* sycl : Added device and backend reg interfaces

* Restructured ggml-sycl.cpp

add amx kernel for gemm (llama/8998)

add intel amx isa detection

add vnni kernel for gemv cases

add vnni and amx kernel support for block_q8_0

code cleanup

fix packing B issue

enable openmp

fine tune amx kernel

switch to aten parallel pattern

add error message for nested parallelism

code cleanup

add f16 support in ggml-amx

add amx kernels for QK_K quant formats: Q4_K, Q5_K, Q6_K and IQ4_XS

update CMakeList

update README

fix some compilation warning

fix compiler warning when amx is not enabled

minor change

ggml-ci

move ggml_amx_init from ggml.c to ggml-amx/mmq.cpp

ggml-ci

update CMakeLists with -mamx-tile, -mamx-int8 and -mamx-bf16

ggml-ci

add amx as an ggml-backend

update header file, the old path for immintrin.h has changed to ggml-cpu-impl.h

minor change

update CMakeLists.txt

minor change

apply weight prepacking in set_tensor method in ggml-backend

fix compile error

ggml-ci

minor change

ggml-ci

update CMakeLists.txt

ggml-ci

add march dependency

minor change

ggml-ci

change ggml_backend_buffer_is_host to return false for amx backend

ggml-ci

fix supports_op

use device reg for AMX backend

ggml-ci

minor change

ggml-ci

minor change

fix rebase

set .buffer_from_host_ptr to be false for AMX backend

vulkan : add backend registry / device interfaces (llama/9721)

* vulkan : add backend registry / device interfaces

* llama : print devices used on model load

fix: allocating CPU buffer with size `0` (llama/9917)

fix: use `vm_allocate` to allocate CPU backend buffer on macOS (llama/9875)

* fix: use `vm_allocate` to allocate CPU backend buffer on macOS

* fix: switch to `posix_memalign` to keep existing `free()` usages work

* feat: move `GGML_ALIGNED_MALLOC` to `ggml-backend-impl.h`, add support for `vm_allocate` on macOS

* style: formatting

* fix: move const outside of `#ifndef`

* style: formatting

* fix: unused var

* fix: transform `GGML_ALIGNED_MALLOC` and `GGML_ALIGNED_FREE` into functions and add them to `ggml-impl.h`

* fix: unused var

* fix: page align to `GGUF_DEFAULT_ALIGNMENT`

* fix: page align to `TENSOR_ALIGNMENT`

* fix: convert `TENSOR_ALIGNMENT` to a macro

* fix: increase page size to `32` on iOS

* fix: iOS page size

* fix: `hbw_posix_memalign` alignment

CUDA: fix 1D im2col, add tests (#993)

ggml : remove redundant set of contexts used field (#978)

This commit removes the setting of the `used` field of the contexts in
the global state (g_state) in `ggml_init`.

The motivation for this change is that I believe that this additional
initialization might not be required after the changes in Commit
45fc4fed0b9fb5b1af4a8525cbebb95e11208732 ("sync : latest changes from
whisper.cpp"), which changed the initialization of the contexts field
from `{ 0 }` to `{ { 0 } }`:

```console
             g_state = (struct ggml_state) {
-                /*.contexts =*/ { 0 },
+                /*.contexts =*/ { { 0 } },
             };
```
My understanding is that the `{0}` initialization might not have
zero-initialized all the nested fields in every array element because of
compiler differences, and might have been the reason for having the
explicit setting of the `used` fields to false.

tests : update type traits call (#0)

ggml-ci

sync : llama.cpp

Fix cann compilation error (llama/9891)

Fix cann compilation error after merging llama.cpp supports dynamically loadable backends.

Vectorize load instructions in dmmv f16 CUDA kernel (llama/9816)

* Vectorize load instructions in dmmv f16 CUDA kernel

Replaces scalar with vector load instructions, which substantially
improves performance on NVIDIA HBM GPUs, e.g. gives a 1.27X overall
speedup for Meta-Llama-3-8B-Instruct-F16 BS1 inference evaluation on
H100 SXM 80GB HBM3. On GDDR GPUs, there is a slight (1.01X) speedup.

* addressed comment

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

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

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

ggml : move more prints to the ggml log system (llama/9839)

* ggml : move more prints to the ggml log system

* show BLAS OpenMP warnings in all builds using debug print

rpc : add backend registry / device interfaces (llama/9812)

* rpc : add backend registry / device interfaces

* llama : add llama_supports_rpc API

* ggml_backend_rpc_start_rpc_server -> ggml_backend_rpc_start_server

musa: add docker image support (llama/9685)

* mtgpu: add docker image support

Signed-off-by: Xiaodong Ye <redacted>
* mtgpu: enable docker workflow

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

Signed-off-by: Xiaodong Ye <redacted>

ggml : fix BLAS with unsupported types (llama/9775)

* ggml : do not use BLAS with types without to_float

* ggml : return pointer from ggml_internal_get_type_traits to avoid unnecessary copies

* ggml : rename ggml_internal_get_type_traits -> ggml_get_type_traits

it's not really internal if everybody uses it

ggml : add backend registry / device interfaces to BLAS backend (llama/9752)

* ggml : add backend registry / device interfaces to BLAS backend

* fix mmap usage when using host buffers

Update building for Android (llama/9672)

* docs : clarify building Android on Termux

* docs : update building Android on Termux

* docs : add cross-compiling for Android

* cmake : link dl explicitly for Android

ggml : add metal backend registry / device (llama/9713)

* ggml : add metal backend registry / device

ggml-ci

* metal : fix names [no ci]

* metal : global registry and device instances

ggml-ci

* cont : alternative initialization of global objects

ggml-ci

* llama : adapt to backend changes

ggml-ci

* fixes

* metal : fix indent

* metal : fix build when MTLGPUFamilyApple3 is not available

ggml-ci

* fix merge

* metal : avoid unnecessary singleton accesses

ggml-ci

* metal : minor fix [no ci]

* metal : g_state -> g_ggml_ctx_dev_main [no ci]

* metal : avoid reference of device context in the backend context

ggml-ci

* metal : minor [no ci]

* metal : fix maxTransferRate check

* metal : remove transfer rate stuff

---------

Co-authored-by: slaren <redacted>

metal : single allocation of encode_async block (llama/9747)

* Single allocation of encode_async block with non-ARC capture in ggml-metal.m

* Moving Block_release to the deallocation code

* Release encode block when re-setting encoding buffer count if needed

* Update ggml/src/ggml-metal.m

---------

Co-authored-by: Georgi Gerganov <redacted>

ggml-alloc : remove buffer_id from leaf_alloc (#987)

This commit removes the buffer_id field from the leaf_alloc struct.

The motivation for is that this field is only written to and never
read/used as far as I can tell. Each tensor_alloc has a buffer_id field
and this is what caused me to look into this more closely, to
understand what the buffer_id in leaf_alloc was used for.

zig : remove obsolete build script

sync : whisper.cpp

vulkan : retry allocation with fallback flags (whisper/2451)

Co-authored-by: Samuel Morris <redacted>

spm : update backend.c -> backend.cpp

examples: add dataset, data shuffling to MNIST (#982)

sync : whisper.cpp

metal : zero-init buffer contexts (whisper/0)

sync : llama.cpp

ggml : fix typo in example usage ggml_gallocr_new (#984)

ggml : fixes after sync (#983)

ggml : remove test-backend-buffer

ggml : fix CUDA build warnings

sync : whisper.cpp

ggml : remove old file (skip) (#0)

cont : fixes

examples : adapt to new ggml backend interfaces

ggml-ci

ggml-backend : add device and backend reg interfaces (llama/9707)

Also:

- metal : fix compute pass descriptor autorelease crash
- ggml-backend : add device description to CPU backend
- ggml: unify backend logging mechanism

sync : llama.cpp

Fixed dequant precision issues in Q4_1 and Q5_1 (llama/9711)

ggml-backend : add device and backend reg interfaces (llama/9707)

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

Initial cmake support of SYCL for AMD GPUs (llama/9658)

sycl: initial cmake support of SYCL for AMD GPUs

vulkan : do not use tensor->extra (llama/9407)

* vulkan : do not use tensor->extra

This patch allows using the Vulkan backend with the RPC backend as
tensor->extra is no longer used.

Ref: #8536

* Adapt GGML_VULKAN_CHECK_RESULTS to extra removal (llama/2)

---------

Co-authored-by: 0cc4m <redacted>

ggml/ex: calculate accuracy in graph, adapt MNIST (#980)

ggml: refactor cross entropy loss CPU impl. (#976)

readme : refresh

metal : add perf-metal tool + fix build

metal : reduce command encoding overhead (llama/9698)

ggml-ci

test: fix OPT_STEP_ADAMW for test-backend-ops (#974)

vulkan : mul_mat: fix UB with small warps (#952)

When the device's warp size is less than 16,
it is possible for loadstride_a (mul_mm.comp:114)
and loadstride_b (mul_mm.comp:115) to be set to 0.
Because they are calculated as: the workgroup size,
multiplied by LOAD_VEC_* (which can be 1) and divided by 16.
And the workgroup size is set to be the same as the
warp/subgroup size.

The loadstride_* variables are used as increments in the
loops that populate the buffers used for the multiplication.

When they are 0 they cause an infinite loop.
But infinite loops without side-effects are UB and the
values of loadstride_* are known at compile time.
So, the compiler quietly optimizes all the loops away.
As a consequence, the buffers are not populated and
the multiplication result is just a matrix with all elements
set to 0.

We prevent the UB by making sure that the workgroup size
will never be less than 16, even if our device has a
smaller warp size (e.g. 8).

Signed-off-by: Salvatore Mesoraca <redacted>

ggml : fix ggml_cast (#973)

ggml: fix gradient allocation logic (#966)

* ggml: fix gradient allocation logic

* gradient allocation in ggml_build_backward_expand

* fixup

* fix test-backend-ops grad

* suggestions by slaren

* fix test1.c

* fix legacy opt API

* fix test-grad0

* remove keep arg

sync : llama.cpp

ggml : define missing HWCAP flags (llama/9684)

ggml-ci

Co-authored-by: Willy Tarreau <redacted>

test-backend-ops : use flops for some performance tests (llama/9657)

* test-backend-ops : use flops for some performance tests

- parallelize tensor quantization

- use a different set of cases for performance and correctness tests

- run each test for at least one second

ggml : add run-time detection of neon, i8mm and sve (llama/9331)

* ggml: Added run-time detection of neon, i8mm and sve

Adds run-time detection of the Arm instructions set features
neon, i8mm and sve for Linux and Apple build targets.

* ggml: Extend feature detection to include non aarch64 Arm arch

* ggml: Move definition of ggml_arm_arch_features to the global data section

Enable use to the rebar feature to upload buffers to the device. (llama/9251)

mtgpu: enable VMM (llama/9597)

Signed-off-by: Xiaodong Ye <redacted>

ggml : remove assert for AArch64 GEMV and GEMM Q4 kernels (llama/9217)

* ggml : remove assert for AArch64 GEMV and GEMM Q4 kernels

* added fallback mechanism when the offline re-quantized model is not
optimized for the underlying target.

* fix for build errors

* remove prints from the low-level code

* Rebase to the latest upstream

cann: fix crash when llama-bench is running on multiple cann devices (llama/9627)

CUDA: remove bad assert (#972)

vulkan : multithread pipeline creation (#963)

vulkan : fix build for GGML_VULKAN_RUN_TESTS, add TFLOPS to log (#961)

vulkan : argsort barriers must be under uniform control flow (#951)

a return before a barrier (that happens only in some threads in
a workgroup) leads to UB.
While the old code actually works on some devices,
it fails on some others (i.e. "smaller" GPUs).

BTW, I think it would be better to set specialization constants
when the graph is built, in that way the local workgroup
could be sized appropriately.
But it would take a lot of work.

Signed-off-by: Salvatore Mesoraca <redacted>

ggml : fix GGML_MAX_N_THREADS + improve formatting (#969)

sync : llama.cpp

ggml-ci

ggml : add AVX512DQ requirement for AVX512 builds (llama/9622)

log : add CONT level for continuing previous log entry (llama/9610)

threads: fix msvc build without openmp (llama/9615)

We're missing atomic_thread_fence() in MSVC builds when openmp is disabled.

cuda: add q8_0->f32 cpy operation (llama/9571)

llama: enable K-shift for quantized KV cache
It will fail on unsupported backends or quant types.

threads: improve ggml_barrier scaling with large number of threads (llama/9598)

Make sure n_barrier and n_barrier_passed do not share the cache line to avoid cache line bouncing.
This optimization shows performance improvements even for n_threads <= 8 cases.

Resurect TSAN (Thread Sanitizer) check so that we can avoid doing expensive read-modify-write
in the normal case and just use thread-fence as originally intended.

ggml : AVX512 gemm for Q4_0_8_8 (llama/9532)

* AVX512 version of ggml_gemm_q4_0_8x8_q8_0

* Remove zero vector parameter passing

* Rename functions and rearrange order of macros

* Edit commments

* style : minor adjustments

* Update x to start from 0

---------

Co-authored-by: Georgi Gerganov <redacted>

metal : use F32 prec for K*Q in vec FA (llama/9595)

ggml-ci

Revert "[SYCL] fallback mmvq (#9088)" (llama/9579)

This reverts commit 50addec9a532a6518146ab837a85504850627316.

musa: enable building fat binaries, enable unified memory, and disable Flash Attention on QY1 (MTT S80) (llama/9526)

* mtgpu: add mp_21 support

Signed-off-by: Xiaodong Ye <redacted>
* mtgpu: disable flash attention on qy1 (MTT S80); disable q3_k and mul_mat_batched_cublas

Signed-off-by: Xiaodong Ye <redacted>
* mtgpu: enable unified memory

Signed-off-by: Xiaodong Ye <redacted>
* mtgpu: map cublasOperation_t to mublasOperation_t (sync code to latest)

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

Signed-off-by: Xiaodong Ye <redacted>

Fix merge error in #9454 (llama/9589)

Signed-off-by: Molly Sophia <redacted>

CUDA: enable Gemma FA for HIP/Pascal (llama/9581)

RWKV v6: RWKV_WKV op CUDA implementation (llama/9454)

* ggml: CUDA unary op EXP

Signed-off-by: Molly Sophia <redacted>
* ggml: rwkv_wkv op CUDA impl

Signed-off-by: Molly Sophia <redacted>
---------

Signed-off-by: Molly Sophia <redacted>

ggml-alloc : fix list of allocated tensors with GGML_ALLOCATOR_DEBUG (llama/9573)

Update CUDA graph on scale change plus clear nodes/params (llama/9550)

* Avoid using saved CUDA graph if scale changes and reset nodes/params on update

Fixes https://github.com/ggerganov/llama.cpp/issues/9451

* clear before resize

examples : adapt to ggml.h changes (#0)

ggml-ci

sync : llama.cpp

ggml : refactoring (llama/#0)

- d6a04f87
- 23e0d70b

ggml : fix builds (llama/0)

ggml-ci

ggml : fix trailing whitespace (llama/0)

ggml-ci

CUDA: fix sum.cu compilation for CUDA < 11.7 (llama/9562)

ggml : fix n_threads_cur initialization with one thread (llama/9538)

* ggml : fix n_threads_cur initialization with one thread

* Update ggml/src/ggml.c

---------

Co-authored-by: Max Krasnyansky <redacted>

threadpool : skip polling for unused threads (llama/9461)

* threadpool: skip polling for unused threads

Currently all threads do N polling rounds even if only 1 thread is active (n_threads_cur == 1).
This commit adds a check to skip the polling for unused threads (ith >= n_threads_cur).

n_threads_cur is now an atomic_int to explicitly tell thread sanitizer that it is written
from one thread and read from other threads (not a race conditions).

* threadpool: further simplify and improve ggml_barrier

Avoid using strict memory order while polling, yet make sure that all threads go through
full memory barrier (memory fence) on ggml_barrier entrace and exit.

* threads: add simple barrier test

This test does lots of small, parallel matmul ops where the barriers in between dominate the overhead.

* threadpool: improve thread sync for new-graphs

Using the same tricks as ggml_barrier. All the polling is done with relaxed memory order
to keep it efficient, once the new graph is detected we do full fence using read-modify-write
with strict memory order.

* threadpool: improve abort handling

Do not use threadpool->ec (exit code) to decide whether to exit the compute loop.
threadpool->ec is not atomic which makes thread-sanitizer rightfully unhappy about it.

Instead introduce atomic threadpool->abort flag used for this. This is consistent with
how we handle threadpool->stop or pause.

While at it add an explicit atomic_load for n_threads_cur for consistency.

* test-barrier: release threadpool before releasing the context

fixes use-after-free detected by gcc thread-sanitizer on x86-64
for some reason llvm sanitizer is not detecting this issue.

ggml : link MATH_LIBRARY not by its full path (llama/9339)

cmake : do not hide GGML options + rename option (llama/9465)

* cmake : do not hide GGML options

ggml-ci

* build : rename flag GGML_CUDA_USE_GRAPHS -> GGML_CUDA_GRAPHS

for consistency

ggml-ci

ggml : IQ4_NL sgemm + Q4_0 AVX optimization (llama/9422)

* squashed

readd my iq4_nl sgemm PR https://github.com/ggerganov/llama.cpp/pull/8049

have ggml_vec_dot_q4_0 do two blocks per loop for avx

try out f16c ggml_vec_dot_iq4_nl, but it's not really faster. as per https://github.com/ggerganov/llama.cpp/pull/8549 we can calculate several blocks at a time with no issue

* shuffle

* remove f16c iq4_nl as i cant make it faster than before

metal : handle zero-sized allocs (llama/9466)

common : reimplement logging (llama/9418)

https://github.com/ggerganov/llama.cpp/pull/9418

cmake : correct order of sycl flags (llama/9497)

cmake : try to fix sycl+intel build (llama/9487)

ggml : ggml_type_name return "NONE" for invalid values (llama/9458)

When running on Windows, the quantization utility attempts to print the types that are not set which leads to a crash.

cmake : use list(APPEND ...) instead of set() + dedup linker (llama/9463)

* cmake : use list(APPEND ...) instead of set() + dedup linker

ggml-ci

* cmake : try fix sycl

* cmake : try to fix sycl 2

* cmake : fix sycl build (llama/9469)

* try fix sycl build

* use CMAKE_CXX_FLAGS as a string variable

---------

Co-authored-by: Georgi Gerganov <redacted>
* one more CMAKE_CXX_FLAGS fix (llama/9471)

---------

Co-authored-by: Michael Podvitskiy <redacted>

cann: Add host buffer type for Ascend NPU (llama/9406)

* feat: Add host buffer type for Ascend NPU(CANN backend)

* fix some checking errors

* Add a few comments

riscv : modify Makefile and add a RISCV_VECT to print log info (llama/9442)

- Added ggml_cpu_has_riscv_v() in GGML to print system info in log
- Modified Makefile to only use flag when cross compiling for RISC-V

cann: Fix error when running a non-exist op (llama/9424)