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>

readme : add link to seamless_comm repo

ggml : disable `fprintf` when building with NDEBUG (#631)

Co-authored-by: Judd <redacted>

ggml-cuda : fix usage without CUDA devices (#627)

cmake : add ROCm config (#626)

ggml-backend update: buffer types, backend registry, graph compare, tests (#620)

* ggml-backend update

* update metal backend

* show metal logs with ggml-backend

* move buffer types to functions

* cuda: add per-device backends

* cuda: add host buffer type

* fix metal build

* ggml_backend_alloc_ctx_tensors : ignore allocated tensors

* ggml_backend_compare_graph_backend fixes

* ci : try to fix metal build

* metal : first print device info, then build kernels

* ci : disable GGML_METAL on Github Actions

* test-backend-ops initial impl (unary and get_rows)

* more op tests

* cleanup

* print test params, add more tests cases for add and mul

* add tests for im2col

* better f16 init

* metal : add basic impl of supports_op

* add test for ggml_concat

* update im2col test params, show callstack with GGML_ASSERT on CUDA failures

* add more rope tests

* add more rope and mul_mat test cases

* add more get_rows test cases
ggml-ci

* add more norm and rms_norm test cases with different eps

* ci : fix metal resource path

ggml-ci

* tests : silence warning

* add ggml_backend_tensor_alloc and ggml_backend_view_init for initializing tensors without ggml-alloc

* add mul_mat test cases without dims 3 and 4
ggml-ci

* check for nans and infs
ggml-ci

* add diag_mask_inf test cases without dims 3 and 4
ggml-ci

* fix cuda leak while backend reg

* fix msvc issues

* remove backend_sched debug causes by default

* gpt-2 : increase graph size

ggml-ci

---------

Co-authored-by: Georgi Gerganov <redacted>

tests : update test-vec0.c for mingw (#619)

For correct building under mingw64

readme : add vit.cpp (#618)

gguf : add tokenizer.chat_template documentation (#616)

ggml : fix ggml_set_2d_inplace (#611)

gguf : prevent out-of-bounds-access on invalid magic (close #614)

sync : whisper.cpp (update whisper example + minor) (#613)

ggml-ci

sync : llama.cpp (cuda, gguf and linker fixes)

update examples and tests to use ggml_allocr_new_measure_from_backend (#608)

* update examples and tests to use ggml_allocr_new_measure_from_backend

* update comments

sync : llama.cpp (CUDA ReLU, CPU-only with CUDA, bloom fix, etc) (#607)

ggml-ci

sync : whisper.cpp (whisper full GPU, fix warnings) (#606)

* sync : whisper.cpp (whisper full GPU, fix warnings)

ggml-ci

* ci : enable CUDA / Metal

ggml-ci

* cuda : fallback to CPU for mul mat ne03 != ne13 (fix SAM + CUDA)

ggml-ci

ggml : replace conv 1D - 2D stage_0 and stage_1 with im2col and mul_mat (#564)

* added conv2d stage 0 - 1 cuda kernels

* add im2col + refactor conv1d and conv2d

* fix params invalid index

* add conv1d and conv2d unit tests

* resolving wrong values and fix mul_mat validation

* improve tests + reduce code duplication

* add cuda kernels

* more data test

* fix ggml_op_count to 70

* add temp test - gemm != mul_mat

* tests : fix test-mul-mat matrix multiplication

* test-mul-mat match gemm == ggml_mul_mat with conv2d op

* replaced gemm by ggml_mul_mat

* ggml_mul_mat cpu backend support fp16 src1

* ggml_mul_mat cuda backend fp16 fixed

* remove unnecessary ggml_cont and removed conv1d-2d functions deprecated

* some fixes

* explain conv1d reshapes

* ggml : fix tests on Arm + do not use BLAS for F16 data

* tests : fix FP16 handling on Arm

* ggml : avoid ggml_cont and ggml_transpose in ggml_conv_xd

* ci : switch back to release

* cuda : fix wrong pointer usage

* ggml : add metal support for im2col and f16xf16 mul mat

* ggml : im2col opts

* Update src/ggml-cuda.cu

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

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

sync : whisper.cpp (ARM 32-bit, abort callback, wav_writer, etc.) (#602)

sync : llama.cpp (CUDA opts, ggml-quants, YARN, etc.) (#601)

ggml-ci

sam : passing parameters and simple prompt (#598)

- most of the model hyperparameters can now be set on CLI
- user can define their own mask prefix
- user can define their own point prompt, although just one

Signed-off-by: Jiri Podivin <redacted>

sam : update documentation to provide executable example (#596)

Also adds the example sample image to the repo to simplify replication.

Signed-off-by: Jiri Podivin <redacted>