metal : refactor kernel loading code (llama/4794)

* metal : detect more GPU families

* metal : refactor kernel loading

* metal : set kernel family requirements

* metal : fix kernel init + fix compile options

* metal : take into account simdgroup reduction support

* metal : print only skipped kernels

* metal : fix check for simdgroup reduction support

* metal : check for Metal 3

* metal : free allocations

* metal : normalize encoder:setComputePipelineStatus calls

ggml-ci

* metal : fix Metal3 family check

ggml-ci

* metal : check for simdgroup matrix mul. feature

ggml-ci

CUDA: faster q8_0 -> f16 dequantization (llama/4895)

sync : whisper.cpp

ggml : fix 32-bit ARM compat for IQ2_XS (whisper/1758)

* ggml : fix 32-bit ARM compat

* ggml : fix fix

* ggml : fix fix fix

whisper : fix segment length with params.no_timestamps == true

params : don't compute timestamps when not printing them (whisper/1755)

backend_sched : fix assignments

ggml-ci

examples : remove obsolete starcoder mmap example

ggml-ci

update gpt-2 example

sync : llama.cpp

ggml-ci

llama : ggml-backend integration (llama/4766)

* llama : ggml-backend integration

* ggml-backend : add names to buffers

* fix unmap after loading

* batched-bench : add tensor_split param

* llama : check for null tensor_split

* ggml-backend : increase GGML_MAX_BACKENDS

* improve graph splitting, partial fix for --no-kv-offload

* cuda : add ggml-backend split buffer support

* cuda : do not create buffer types for devices that don't exist (fixes usage without CUDA devices available)

* ggml : fix null backend dereference (llama/4807)

* ggml : fix null backend dereference

* ggml : also check ggml_backend_is_cpu

* test-backend-ops : check buffer allocation failures

* llama : add cparam (split_mode) and command line argument (--split-mode, -sm) to configure the split mode (none, layer or row)

* ggml : fix mul_mat_id work size

* llama : rewrite session kv load/set without graphs

* minor

* llama : only initialize used backends, free backends on context free

* llama : abort ctx if cuda backend init fails

* llama : rewrite lora with ggml-backend and compute on CPU

ggml-ci

* llama : only map to a backend buffer the region of the file mapping containing the tensors used in the buffer

* opencl : add ggml-backend buffer type

* cuda : only use batched_cublas with batched mat muls (fixes fp16 tg perf)

* llama : on Metal, by default offload the full model

ggml-ci

* metal : page align the data ptr (llama/4854)

* Apply suggestions from code review

Co-authored-by: Johannes Gäßler <redacted>
* cuda : fix split buffer free

* address review comments

* llama-bench : add split-mode parameter

* fix whitespace

* opencl : fix double initialization

* server : add --split-mode parameter

* use async copy and compute to improve multi-gpu performance

ggml-ci

* use async memcpys to copy the graph outputs to the CPU

* fix opencl

* use a host buffer for the cpu compute buffer for faster copies to the gpu

---------

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

CUDA: fix softmax compile for old CUDA versions (llama/4862)

Importance Matrix calculation (llama/4861)

* imatrix: 1st version

* imatrix: WIP

* Cleanup

* Update examples/imatrix/imatrix.cpp

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

Co-authored-by: Iwan Kawrakow <redacted>
Co-authored-by: Georgi Gerganov <redacted>

sync : whisper.cpp

sync : ggml

main : add cli option to disable system prints (whisper/1740)

sync : llama.cpp

ggml : SOTA 2-bit quants (add IQ2_XS) (llama/4856)

* iq2_xs: basics

* iq2_xs: this should have been in the basics

* iq2_xs: CUDA and scalar CPU works

* iq2_xs: WIP Metal

* iq2_xs: Metal now works

* iq2_xs: working, but dog slow, ARM_NEON dot product

* iq2_xs: better ARM_NEON dot product

We are now at 19.5 t/s for TG-128 and 61 t/s for PP-512 when
running on the CPU.

* iq2_xs: AVX2 dot product - 19.5 t/s

* iq2_xs: faster AVX2 dit product

21.4 t/s for TG-128, 59.2 t/s for PP-512.
The latter is 2x compared to the previous version.

* iq2_xs: had forgotten to delete iq2-data.h

* Add llama enum for IQ2_XS

---------

Co-authored-by: Iwan Kawrakow <redacted>

metal : put encoder debug group behind a define (llama/4873)

metal : improve dequantize precision to match CPU (llama/4836)

ggml-ci

ggml : fix vld1q_s8_x4 32-bit compat (llama/4828)

* ggml : fix vld1q_s8_x4 32-bit compat

ggml-ci

* ggml : fix 32-bit ARM compat (cont)

ggml-ci

CUDA: faster softmax via shared memory + fp16 math (llama/4742)

metal : fix deprecation warning (#690)

ggml : remove ggml_cpy_inplace and ggml_cont_inplace (#693)

Update README.md (#692)

metal : wrap each operation in debug group (#690)

ggml : change GGML_MAX_NAME at compile time (#682)

* change GGML_MAX_NAME to 128

* allow controlling the value of GGML_MAX_NAME through external macro definitions

Fix execlp call (#689)

NULL can be an integer constant expression with the value zero, in this case the behavior would be undefined because of an incorrect type being passed to the variable arguments.

sync : llama.cpp

SOTA 2-bit quants (llama/4773)

* iq2_xxs: basics

* iq2_xxs: scalar and AVX2 dot products

Needed to change Q8_K to have quants in the -127...127 range,
else the IQ2_XXS AVX implementation becomes very awkward.
The alternative would have been to use Q8_0 instead. Perhaps
I'll change later, for now this is what we have.

* iq2_xxs: ARM_NEON dot product

Somehow strangely slow (112 ms/token).

* iq2_xxs: WIP Metal

Dequantize works, something is still wrong with the
dot product.

* iq2_xxs: Metal dot product now works

We have
PP-512 = 475 t/s
TG-128 = 47.3 t/s

Not the greatest performance, but not complete garbage either.

* iq2_xxs: slighty faster dot product

TG-128 is now 48.4 t/s

* iq2_xxs: slighty faster dot product

TG-128 is now 50.9 t/s

* iq2_xxs: even faster Metal dot product

TG-128 is now 54.1 t/s.

Strangely enough, putting the signs lookup table
into shared memory has a bigger impact than the
grid values being in shared memory.

* iq2_xxs: dequantize CUDA kernel - fix conflict with master

* iq2_xxs: quantized CUDA dot product (MMVQ)

We get TG-128 = 153.1 t/s

* iq2_xxs: slightly faster CUDA dot product

TG-128 is now at 155.1 t/s.

* iq2_xxs: add to llama ftype enum

* iq2_xxs: fix MoE on Metal

* Fix missing MMQ ops when on hipBLAS

I had put the ggml_supports_mmq call at the wrong place.

* Fix bug in qequantize_row_iq2_xxs

The 0.25f factor was missing.
Great detective work by @ggerganov!

* Fixing tests

* PR suggestion

---------

Co-authored-by: Iwan Kawrakow <redacted>

CUDA: fixed redundant value dequantization (llama/4809)

ggml : use __builtin_amdgcn_sudot4 in __dp4a for gfx11 (llama/4787)

ggml : do not sched_yield when calling BLAS (llama/4761)

* ggml : do not sched_yield when calling BLAS

ggml-ci

* ggml : fix do_yield logic

ggml-ci

* ggml : simplify do_yield logic

ggml-ci

Print backend name on test-backend-ops failure (llama/4751)

ggml : include stdlib.h before intrin.h (llama/4736)

fix ROCm on Windows (#683)

* fix ROCm on windows

* keep windows does not link m.lib

* add hip missed compile definitions

* never link m.lib

* use find_library for m

---------

Co-authored-by: slaren <redacted>

gguf : add keys for kv sizes to spec (#676)

* Add keys for kv sizes to GGUF spec

* Fix types of key_length and value_length

fix : cuda order of synchronization when setting a buffer (#679)

* fix : cuda order of synchronization when setting a buffer

* also sync before memcpy

---------

Co-authored-by: slaren <redacted>

metal : switch back to default.metallib (#681)

ggml-ci

ggml : fix q2_k bpw in comments (#680)

whisper : reset the "batched" timings (whisper/1721)

ggml : add error handling to graph_compute (whisper/1714)

scripts : allow to skip commits during sync (#678)

Update src/ggml-cuda.cu

Co-authored-by: slaren <redacted>

cuda : mark I16 and I32 ops as unsupported

ggml-ci

metal : add kernel_get_rows_i32

ggml-ci

sync : llama.cpp

ggml-ci

metal : optimize ggml_mul_mat_id (faster Mixtral PP) (llama/4725)

* ggml : disable fast-math for Metal (cmake build only)

ggml-ci

* metal : fix Metal API debug warnings

* cmake : add -fno-inline for Metal build (llama/4545)

* metal : fix API debug warnings

* metal : fix compile warnings

* metal : use uint64_t for strides

* cmake : rename option to LLAMA_METAL_SHADER_DEBUG

* metal : fix mat-vec Q8_0 kernel for BS > 1

* metal : normalize mat-vec kernel signatures

* cmake : respect LLAMA_QKK_64 option

* metal : fix mat-vec Q4_K kernel for QK_K == 64

* metal : optimizing ggml_mul_mat_id (wip)

* metal : minor fix

* metal : opt mul_mm_id

metal : enable shader debugging (cmake option) (llama/4705)

* ggml : disable fast-math for Metal (cmake build only)

ggml-ci

* metal : fix Metal API debug warnings

* cmake : add -fno-inline for Metal build (llama/4545)

* metal : fix API debug warnings

* metal : fix compile warnings

* metal : use uint64_t for strides

* cmake : rename option to LLAMA_METAL_SHADER_DEBUG

* metal : fix mat-vec Q8_0 kernel for BS > 1

* metal : normalize mat-vec kernel signatures

* cmake : respect LLAMA_QKK_64 option

* metal : fix mat-vec Q4_K kernel for QK_K == 64

ggml-ci

ggml : add ggml_vdotq_s32 alias (llama/4715)

ggml-ci

CUDA: fixed tensor cores not being used on RDNA3 (llama/4697)

ggml : add ggml_cpu_has_avx_vnni() (llama/4589)

* feat: add avx_vnni based on intel documents

* ggml: add avx vnni based on intel document

* llama: add avx vnni information display

* docs: add more details about using oneMKL and oneAPI for intel processors

* docs: add more details about using oneMKL and oneAPI for intel processors

* docs: add more details about using oneMKL and oneAPI for intel processors

* docs: add more details about using oneMKL and oneAPI for intel processors

* docs: add more details about using oneMKL and oneAPI for intel processors

* Update ggml.c

Fix indentation upgate

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

Co-authored-by: Georgi Gerganov <redacted>

CUDA: fix tensor core logic for Pascal and HIP (llama/4682)

cuda: fix vmm oom issue on NVIDIA AGX Orin (llama/4687)

Signed-off-by: hydai <redacted>

scripts : fix sync order + metal sed

swift : add Swift Package declaration (#674)

* feat: adds swift package declaration
- allows importing ggml as package dependency in llama.cpp and whisper.cpp
- resolves issues with duplicate symbol error when importing llama.cpp and whisper.cpp as package dependency

* fixes the src paths in package.swift

ggml : extend ggml_get_rows, ggml_repeat, ggml_concat (#639)

* add more int ops

* ggml_compute_forward_dup_bytes

* add tests

* PR comments

* tests : minor indentations

---------

Co-authored-by: Georgi Gerganov <redacted>

scripts : do not sync synced commits

sync : whisper.cpp

ci : build with CLBlast + ggml-opencl use GGML_API (whisper/1576)

* Build with CLBlast

* Declare GGML_API

After rebasing, examples/talk-llama failed:

"D:\a\whisper.cpp\whisper.cpp\build\ALL_BUILD.vcxproj" (build target) (1) ->
"D:\a\whisper.cpp\whisper.cpp\build\examples\talk-llama\talk-llama.vcxproj" (default target) (14) ->
(Link target) ->
  llama.obj : error LNK2019: unresolved external symbol ggml_cl_free_data referenced in function "public: __cdecl llama_model::~llama_model(void)" (??1llama_model@@QEAA@XZ) [D:\a\whisper.cpp\whisper.cpp\build\examples\talk-llama\talk-llama.vcxproj]
  llama.obj : error LNK2019: unresolved external symbol ggml_cl_transform_tensor referenced in function "public: void __cdecl llama_model_loader::load_all_data(struct ggml_context *,void (__cdecl*)(float,void *),void *,struct llama_mlock *)" (?load_all_data@llama_model_loader@@QEAAXPEAUggml_context@@P6AXMPEAX@Z1PEAUllama_mlock@@@Z) [D:\a\whisper.cpp\whisper.cpp\build\examples\talk-llama\talk-llama.vcxproj]
  D:\a\whisper.cpp\whisper.cpp\build\bin\Release\talk-llama.exe : fatal error LNK1120: 2 unresolved externals [D:\a\whisper.cpp\whisper.cpp\build\examples\talk-llama\talk-llama.vcxproj]

scripts : print list of sync commits

scripts : fix format-patch range

scripts : do not sync commits from this repo

sync : whisper.cpp

ggml-ci

whisper : Replace WHISPER_PRINT_DEBUG with WHISPER_LOG_DEBUG (whisper/1681)

scripts : add sync-whisper-am.sh

ggml : fix some mul mat cases + add tests for src1 F16 (#669)

* fixed mul-mat error for old GPUs

* style fixes

* add mul mat src1 f16 test cases, fix more cases

ggml-ci

---------

Co-authored-by: bssrdf <redacted>
Co-authored-by: slaren <redacted>

gpt-2 : update README.md (#671)

readme : add steps for Android compilation and inference (#664)

* Update README.md for android instructions

* fix minor typo

scripts : fix sed in sync-llama.am.sh

sync : llama.cpp

ggml-ci

ggml : fix dot product for ARM (llama/4630)

ggml-ci

cuda : fix vmm pool with multi GPU (llama/4620)

* cuda : fix vmm pool with multi GPU

* hip

* use recommended granularity instead of minimum

* better error checking

* fix mixtral

* use cudaMemcpy3DPeerAsync

* use cuda_pool_alloc in ggml_cuda_op_mul_mat

* consolidate error checking in ggml_cuda_set_device

* remove unnecessary inlines

ggml-ci

* style fixes

* only use vmm for the main device

* fix scratch buffer size, re-enable vmm pool for all devices

* remove unnecessary check id != g_main_device

Update comment for AdamW implementation reference. (llama/4604)

Co-authored-by: Will Findley <redacted>

Fix new CUDA10 compilation errors (llama/4635)

sync : llama.cpp

ggml-ci

cmake : update CUDA build to support VMM

cuda : improve cuda pool efficiency using virtual memory (llama/4606)

* cuda : improve cuda pool efficiency using virtual memory

* fix mixtral

* fix cmake build

* check for vmm support, disable for hip

ggml-ci

* fix hip build

* clarify granularity

* move all caps to g_device_caps

* refactor error checking

* add cuda_pool_alloc, refactor most pool allocations

ggml-ci

* fix hip build

* CUBLAS_TF32_TENSOR_OP_MATH is not a macro

* more hip crap

* llama : fix msvc warnings

* ggml : fix msvc warnings

* minor

* minor

* cuda : fallback to CPU on host buffer alloc fail

* Update ggml-cuda.cu

Co-authored-by: Johannes Gäßler <redacted>
* Update ggml-cuda.cu

Co-authored-by: Johannes Gäßler <redacted>
* ensure allocations are always aligned

* act_size -> actual_size

---------

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

fallback to CPU buffer if host buffer alloc fails (llama/4610)

CUDA: fixed row rounding for 0 tensor splits (llama/4594)

scripts : fix PR number parsing during sync

scripts : improve llama sync patch

scripts : sync tests / headers

scripts : remove exit

scripts : fix PR number sed

scripts : add sync-llama-am.sh

sync : llama.cpp (ggml_scale, ggml_row_size, ggml_mul_mat_set_prec) (#662)

* sync : llama.cpp (ggml_scale, ggml_row_size, ggml_mul_mat_set_prec)

ggml-ci

* ggml : add comment about backward GGML_OP_DIAG_MASK_INF (#4203)

* llama : fix platforms without mmap (#4578)

* llama : fix platforms without mmap

* win32 : limit prefetch size to the file size

* fix win32 error clobber, unnecessary std::string in std::runtime_error

* ggml-alloc : fix ggml_tallocr_is_own

* whisper : minor

* ggml : cuda jetson + arm quants warnings

ggml-ci

---------

Co-authored-by: Herman Semenov <redacted>
Co-authored-by: slaren <redacted>

cuda : fix synchronization with tensor get/set (#659)

cuda : fix im2col_f32_f16 (#658)

cmake : change installation path of static libraries to standard directory (#645)

Co-authored-by: Cheng-Chi Wang <redacted>

sync : llama (mul_mat_id + get_rows kernels, typos) (#649)

* sync : llama (mul_mat_id + get_rows kernels, typos)

ggml-ci

* cuda : restore im2col

ggml-ci

* metal : fix accuracy of dequantization kernels

* cuda : restore correct im2col kernel

ggml-ci

* metal : fix moe test by reducing the expert size

ggml-ci

* cuda : fix bin bcast when src1 and dst have different types

---------

Co-authored-by: slaren <redacted>

ggml: new gpu kernels + extends ggml_leaky_relu + ggml_pad (#621)

* add new cuda kernels and new op ggml_pad

* add ggml_tanh cuda kernel

* remove old broadcast impl

* restore some changes

* cuda: optimized im2col + group_norm kernels

* extent ggml_leaky -> ggml_leaky_relu

* fix some code issues

* cuda: concat support 4 dims

* cuda: fix ggml_acc + add backends ops test

* restore ggml_pad + add backend op test

* metal : implement GGML_OP_ACC

* ggml : fix bug in ggml_upscale

* metal : add ggml_upscale

* metal : add ggml_tanh

* metal : add ggml_gelu_quick

* ggml : make ggml_pad more general purpose

* metal : add ggml_pad

* ggml_leaky_relu as regular op + fix identation

* cuda: ggml_acc admit all op_parms

* negative_slope better pass param

* metal : add ggml_leaky_relu

* metal : add ggml_group_norm

* cuda : minor

* ggml : add GGML_OP_LEAKY_RELU to ggml_compute_backward

* metal : soft max, tanh, supports_op fixes

* test-backend-ops : add sentinels between tensors to detect overflows

---------

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

gguf : document Mixtral changes in spec (#646)

* add new tensor names

* add new keys

* fix tensor names

* gguf : change wording a bit

---------

Co-authored-by: Georgi Gerganov <redacted>

sync : whisper.cpp (metal soft max fix + example prints)

sync : llama.cpp (fused soft max, gpu cpy ops, etc.) (#640)

* sync : llama.cpp (fused soft max, gpu cpy ops, etc.)

ggml-ci

* cuda : restore accidentally deleted changes

ggml-ci

* cuda : fix rope + disable device-side dequantize

ggml-ci

* test-backend-ops : enable stablelm rope test

* cuda : remove rope assert

* sync.sh : add test-backend-ops

* ggml : fix ggml_concat + ggml_get_n_tasks logic

* sync : whisper.cpp

ggml-ci

* metal : fix assert

* ci : fix Metal path to shaders

ggml-ci

* whisper : fix bug if metal init fails

---------

Co-authored-by: slaren <redacted>

ggml-backend : remove backend self-registration (#641)

test-backend-ops : add performance eval mode + improve CUDA repeat and binary broadcast ops performance (#636)

* ggml-cuda : implement repeat with bin_bcast

* ggml-cuda : change supports_op for mul_mat to match compute_forward

* test-backend-ops : add performance eval mode

* improve formatting

* add sd test cases

* fix test case

* ggml-cuda : bin_bcast: better block sizes, two elements per thread

* metal : add dim3 broadcast support for mul mat

* cleanup

* typo

* metal : enable mul mat-vec for dim2 > 1

* metal : mul mat-vec support dim3 broadcasts

ggml-ci

* ggml-cuda : fix bin_bcast for ne0=1
ggml-ci

* ggml-cuda : limit block size z dim to 64

* test-backend-ops : add test cases

* test-backend-ops : add warmup run, print test type before trying to compute

* ggml-cuda : bin_bcast: collapse dimensions when possible, add fallback kernel for large tensors
ggml-ci

* test-backend-ops : avoid division by zero

---------

Co-authored-by: Georgi Gerganov <redacted>

test-backend-ops : initialize ggml_argsort test with unique values to avoid ties (#634)

ggml-ci

metal : check supported ops at runtime (#632)

* metal : check supported ops at runtime

* metal : remove TODOs

ggml : full broadcast in mul, add, div + ggml_mul_mat_id, ggml_argsort, ggml_top_k (#625)

* ggml : support broadcasting in dim 0 in add and mul

* add cuda add/mul broadcast impl
add configurable eps to cuda norm

* add metal impl
ggml-ci

* deduplicate code in cuda impl

* try to optimize cuda impl

* ggml : support broadcasting in ggml_div

* test-backend-ops : allow filtering by op and backend

* ggml-cuda : add ggml_div impl

* ggml : add ggml_mul_mat_id, ggml_sort, ggml_top_k (CPU only)

* fix ggml_div threads

* fix ggml_div with accelerate

* ggml_sort -> ggml_argsort

* whatever

* actually fix accelerate div

* disable opencl ci

* ci : disable ctest error check temporarily until we fix backend ops test

* cmake : propagete GGML_USE_xxx compile flags with ggml target

* whisper : utlize new ggml_add broadcast for dim 0

* cmake : adendum to ee666ae9

* ggml_backend_graph_copy : fix leak

* ggml_cuda : add ggml_sum_rows impl

* metal : add ggml_div

* metal : add ggml_sum_rows

* ggml_cuda : add ggml_argsort impl

* move kernel

* metal : add ggml_argsort

* mul_mat_id : fix missing init task

* cuda/metal: fix argsort synchronization

* metal : add ggml_mul_mat_id

* ggml-cuda : add mul_mat_id for f16 + tensor cores

* test-backend-ops : add tests for quants mat mul

* ggml : fix q5_0 and q5_1 hist stats

* test-backend-ops : use smaller matrices to avoid automatic offloading, add mat-vec tests

* metal : fix alibi to match the CPU behavior

* metal : check dimensions in supports_op

* test-backend-ops : reduce error threshold for mat muls

* ggml-cuda : simplify dequantize funs, add supports_op by type for mul_mat_id

* ggml-cuda : support quantized types in mul_mat_id with cublas

* ggml-cuda : add fallback over CPU for mul_mat_id

* test-backend-ops : increase mul mat error threshold

* cleanup
ggml-ci

* test-backend-ops : fix usage

* cleanup

* ci : re-enable tests

* metal : fix compile warnings

---------

Co-authored-by: Georgi Gerganov <redacted>