]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
5 months agoDebian development package only depends on base libraries
Mathieu Baudier [Tue, 21 Jan 2025 11:03:04 +0000 (12:03 +0100)]
Debian development package only depends on base libraries

5 months agoImprove debian build based on lintian
Mathieu Baudier [Tue, 21 Jan 2025 10:57:53 +0000 (11:57 +0100)]
Improve debian build based on lintian

5 months agoImprove dependencies
Mathieu Baudier [Tue, 21 Jan 2025 10:17:46 +0000 (11:17 +0100)]
Improve dependencies

5 months agoForce build type Release
Mathieu Baudier [Tue, 21 Jan 2025 10:17:35 +0000 (11:17 +0100)]
Force build type Release

5 months agoMake sure that multiarch is well supported
Mathieu Baudier [Tue, 21 Jan 2025 06:49:29 +0000 (07:49 +0100)]
Make sure that multiarch is well supported

5 months agoImprove debian build based on lintian feedback
Mathieu Baudier [Mon, 20 Jan 2025 19:11:46 +0000 (20:11 +0100)]
Improve debian build based on lintian feedback

5 months agoInitial Debian packaging
Mathieu Baudier [Mon, 20 Jan 2025 14:18:52 +0000 (15:18 +0100)]
Initial Debian packaging

5 months agosync : whisper.cpp upstream/0.0.1642
Georgi Gerganov [Tue, 14 Jan 2025 08:39:25 +0000 (10:39 +0200)]
sync : whisper.cpp

5 months agoscripts : sync gguf (cont)
Georgi Gerganov [Tue, 14 Jan 2025 07:41:40 +0000 (09:41 +0200)]
scripts : sync gguf (cont)

5 months agoGGUF: C++ refactor, backend support, misc fixes (skip) (llama/11030)
Johannes Gäßler [Tue, 14 Jan 2025 07:31:07 +0000 (09:31 +0200)]
GGUF: C++ refactor, backend support, misc fixes (skip) (llama/11030)

ggml-ci

5 months agoscripts : sync gguf
Georgi Gerganov [Tue, 14 Jan 2025 07:29:27 +0000 (09:29 +0200)]
scripts : sync gguf

5 months agoggml : add opencl backend (skip) (llama/10693)
lhez [Tue, 14 Jan 2025 07:24:03 +0000 (09:24 +0200)]
ggml : add opencl backend (skip) (llama/10693)

---------

Co-authored-by: Skyler Szot <redacted>
Co-authored-by: Shangqing Gu <redacted>
Co-authored-by: Alexander Angus <redacted>
Co-authored-by: Hongqiang Wang <redacted>
Co-authored-by: Max Krasnyansky <redacted>
5 months agoscripts : sync opencl
Georgi Gerganov [Tue, 14 Jan 2025 07:20:27 +0000 (09:20 +0200)]
scripts : sync opencl

5 months agosync : llama.cpp
Georgi Gerganov [Tue, 14 Jan 2025 07:17:22 +0000 (09:17 +0200)]
sync : llama.cpp

ggml-ci

5 months agocuda : CUDA Graph Compute Function Refactor (precursor for performance improvements...
Andreas Kieslinger [Mon, 13 Jan 2025 15:45:53 +0000 (16:45 +0100)]
cuda : CUDA Graph Compute Function Refactor (precursor for performance improvements) (llama/11042)

* Refactor: Moves cuda graph executable update step to separate function.

* Refactor: Moves cuda graph update check to separate function.

* Refactor: Moves cuda graph maintenance (update or adjusting copy parameters) to separate function for improved readability.

* Fix: Adds missing reference to maintain_cuda_graph() definition.

* Refactor: Improves structure and abstractions by moving CUDA graph evaluation and capture to its own function.

* Refactor: Moves node graph checks and copy ops into individual function for improved readability.

* Refactor: Removes code permanently excluded from compilation to increase readability.

* Style: Adds missing newline

* Style: Consolidates several neighboring '#ifdef USE_CUDA_GRAPH' into a single one

* Refactor: Makes 'cuda_graph_update_required' a local variable

* remove double lines between functions

---------

Co-authored-by: slaren <redacted>
5 months agoggml : do not define GGML_USE_CUDA when building with GGML_BACKEND_DL (llama/11211)
Radoslav Gerganov [Mon, 13 Jan 2025 11:31:41 +0000 (13:31 +0200)]
ggml : do not define GGML_USE_CUDA when building with GGML_BACKEND_DL (llama/11211)

Build fails when using HIP and GGML_BACKEND_DL:
```
/usr/bin/ld: ../ggml/src/libggml.so: undefined reference to `ggml_backend_cuda_reg'
collect2: error: ld returned 1 exit status
```
This patch fixes this.

5 months agoVulkan: Fix float16 use on devices without float16 support + fix subgroup_size_contro...
0cc4m [Fri, 10 Jan 2025 05:39:33 +0000 (06:39 +0100)]
Vulkan: Fix float16 use on devices without float16 support + fix subgroup_size_control validation error (llama/11161)

* Vulkan: Remove float16 use in shaders

* Fix validation error about subgroup_size_control extension

5 months agollama: add support for QRWKV6 model architecture (llama/11001)
Molly Sophia [Fri, 10 Jan 2025 01:58:08 +0000 (09:58 +0800)]
llama: add support for QRWKV6 model architecture (llama/11001)

llama: add support for QRWKV6 model architecture (llama/11001)

* WIP: Add support for RWKV6Qwen2

Signed-off-by: Molly Sophia <redacted>
* RWKV: Some graph simplification

Signed-off-by: Molly Sophia <redacted>
* Add support for RWKV6Qwen2 with cpu and cuda GLA

Signed-off-by: Molly Sophia <redacted>
* RWKV6[QWEN2]: Concat lerp weights together to reduce cpu overhead

Signed-off-by: Molly Sophia <redacted>
* Fix some typos

Signed-off-by: Molly Sophia <redacted>
* code format changes

Signed-off-by: Molly Sophia <redacted>
* Fix wkv test & add gla test

Signed-off-by: Molly Sophia <redacted>
* Fix cuda warning

Signed-off-by: Molly Sophia <redacted>
* Update README.md

Signed-off-by: Molly Sophia <redacted>
* Update ggml/src/ggml-cuda/gla.cu

Co-authored-by: Georgi Gerganov <redacted>
* Fix fused lerp weights loading with RWKV6

Signed-off-by: Molly Sophia <redacted>
* better sanity check skipping for QRWKV6 in llama-quant

thanks @compilade

Signed-off-by: Molly Sophia <redacted>
Co-authored-by: compilade <redacted>
---------

Signed-off-by: Molly Sophia <redacted>
Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: compilade <redacted>
5 months agoSYCL: Refactor ggml_sycl_compute_forward (llama/11121)
Akarshan Biswas [Fri, 10 Jan 2025 00:13:03 +0000 (05:43 +0530)]
SYCL: Refactor ggml_sycl_compute_forward (llama/11121)

* SYCL: refactor ggml_sycl_compute_forward

* SYCL: add back GGML_USED(dst) to ggml_sycl_cpy

* SYCL: add function name to noop debug

* SYCL: Some device info print refactoring and add details of XMX availability

5 months agofix: add missing msg in static_assert (llama/11143)
hydai [Wed, 8 Jan 2025 20:03:28 +0000 (04:03 +0800)]
fix: add missing msg in static_assert (llama/11143)

Signed-off-by: hydai <redacted>
5 months agollamafile : ppc64le MMA INT8 implementation (llama/10912)
amritahs-ibm [Wed, 8 Jan 2025 10:54:19 +0000 (16:24 +0530)]
llamafile : ppc64le MMA INT8 implementation (llama/10912)

This change upstreams llamafile's cpu matrix
multiplication kernels for ppc64le using MMA
builtins for quantised int8 datatype.

This change results in 10% - 70% improvement
in total speed(ie all tokens/total time), across
various batch sizes.

The patch is tested with Meta-Lllama-3-8B,
Mistral-7B, Llama-2-7B-chat-hf models on a
IBM POWER10 machine.

Signed-off-by: Amrita H S <redacted>
5 months agoDisable GL_KHR_cooperative_matrix Vulkan extension if not available. (llama/11117)
Mathieu Baudier [Wed, 8 Jan 2025 08:18:13 +0000 (09:18 +0100)]
Disable GL_KHR_cooperative_matrix Vulkan extension if not available. (llama/11117)

* Disable GL_KHR_cooperative_matrix Vulkan extension if not available.

* Perform Vulkan extensions checks in a more sensible order

* Remove unnecessary #ifdef directive

5 months agofix: Vulkan shader gen binary path when Cross-compiling (llama/11096)
ag2s20150909 [Wed, 8 Jan 2025 08:17:29 +0000 (16:17 +0800)]
fix: Vulkan shader gen binary path when Cross-compiling (llama/11096)

* fix: Vulkan shader gen binary path when cross compiling

5 months agoGGUF: C++ refactor, backend support, misc fixes (llama/11030)
Johannes Gäßler [Tue, 7 Jan 2025 17:01:58 +0000 (18:01 +0100)]
GGUF: C++ refactor, backend support, misc fixes (llama/11030)

* GGUF: C++ refactor, backend support, misc fixes

remove ggml_tensor.backend

update CODEOWNERS [no ci]

remove gguf_get_data from API

revise GGUF API data types

5 months agoggml-backend : only offload from host buffers (fix) (llama/11124)
Diego Devesa [Tue, 7 Jan 2025 15:11:57 +0000 (16:11 +0100)]
ggml-backend : only offload from host buffers (fix) (llama/11124)

5 months agoggml-backend : only offload from host buffers (llama/11120)
Diego Devesa [Tue, 7 Jan 2025 11:38:05 +0000 (12:38 +0100)]
ggml-backend : only offload from host buffers (llama/11120)

5 months agorpc : code cleanup (llama/11107)
Radoslav Gerganov [Tue, 7 Jan 2025 06:37:02 +0000 (08:37 +0200)]
rpc : code cleanup (llama/11107)

Remove duplicated macros, use GGML_LOG_ERROR for errors

5 months agoSYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6 (llama/11087)
Akarshan Biswas [Tue, 7 Jan 2025 06:26:07 +0000 (11:56 +0530)]
SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6 (llama/11087)

* SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6

* Revert "SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6"

This reverts commit f62dc45f318e48d375e7734b34cbddee81deed52.

* Reland: Use get_multi_ptr instead of deprecated get_pointer in wkv6

5 months agoCUDA: add BF16 support (llama/11093)
Johannes Gäßler [Mon, 6 Jan 2025 01:33:52 +0000 (02:33 +0100)]
CUDA: add BF16 support (llama/11093)

* CUDA: add BF16 support

5 months agoVulkan: Add device-specific blacklist for coopmat for the AMD proprietary driver...
0cc4m [Sat, 4 Jan 2025 20:09:59 +0000 (21:09 +0100)]
Vulkan: Add device-specific blacklist for coopmat for the AMD proprietary driver (llama/11074)

* Vulkan: Add device-specific blacklist for coopmat for the AMD proprietary driver

* Add (TM) to AMD name check

5 months agoSupport for models with non-512-aligned tensors over RPC. (llama/11047)
matt23654 [Sat, 4 Jan 2025 16:10:30 +0000 (16:10 +0000)]
Support for models with non-512-aligned tensors over RPC. (llama/11047)

* Added init tensor calling code

* Added get_alloc_size forwarding

* Cleaned up and improved type/error handling.

* fix: remove trailing whitespaces.

* Cleanup and use GGML error logging functions.

* Handle potentially dangerous edge cases.

* Apply suggestions from code review

Co-authored-by: Diego Devesa <redacted>
---------

Co-authored-by: Diego Devesa <redacted>
5 months agofix: Vulkan shader gen binary path (llama/11037)
Gilad S. [Sat, 4 Jan 2025 08:17:31 +0000 (10:17 +0200)]
fix: Vulkan shader gen binary path (llama/11037)

5 months agoggml : allow loading backend with env variable (#1059)
Radoslav Gerganov [Sun, 5 Jan 2025 07:50:37 +0000 (09:50 +0200)]
ggml : allow loading backend with env variable (#1059)

ref: #1058

5 months agosync : whisper.cpp
Georgi Gerganov [Sat, 4 Jan 2025 08:51:09 +0000 (10:51 +0200)]
sync : whisper.cpp

5 months agosync : llama.cpp
Georgi Gerganov [Fri, 3 Jan 2025 11:33:25 +0000 (13:33 +0200)]
sync : llama.cpp

ggml-ci

5 months agometal : avoid uint (llama/11019)
Georgi Gerganov [Fri, 3 Jan 2025 09:26:14 +0000 (11:26 +0200)]
metal : avoid uint (llama/11019)

5 months agoggml : fixes for AVXVNNI instruction set with MSVC and Clang (llama/11027)
Srihari-mcw [Tue, 31 Dec 2024 14:23:33 +0000 (19:53 +0530)]
ggml : fixes for AVXVNNI instruction set with MSVC and Clang (llama/11027)

* Fixes for clang AVX VNNI

* enable AVX VNNI and alder lake build for MSVC

* Apply suggestions from code review

---------

Co-authored-by: slaren <redacted>
5 months agovulkan: optimize mul_mat for small values of N (llama/10991)
Jeff Bolz [Mon, 30 Dec 2024 17:27:11 +0000 (11:27 -0600)]
vulkan: optimize mul_mat for small values of N (llama/10991)

Make the mul_mat_vec shaders support N>1 (as a spec constant, NUM_COLS) where
the batch_strides are overloaded to hold the row strides. Put the loads from the
B matrix in the innermost loop because it should cache better.

Share some code for reducing the result values to memory in mul_mat_vec_base.

5 months agovulkan: im2col and matmul optimizations for stable diffusion (llama/10942)
Jeff Bolz [Sun, 29 Dec 2024 09:16:34 +0000 (03:16 -0600)]
vulkan: im2col and matmul optimizations for stable diffusion (llama/10942)

* tests: Add im2col perf tests

* vulkan: optimize im2col, more elements per thread

* vulkan: increase small tile size for NV_coopmat2

* vulkan: change im2col to 512 elements per workgroup

5 months agovulkan: Use push constant offset to handle misaligned descriptors (llama/10987)
Jeff Bolz [Sun, 29 Dec 2024 08:35:11 +0000 (02:35 -0600)]
vulkan: Use push constant offset to handle misaligned descriptors (llama/10987)

5 months agovulkan: multi-row k quants (llama/10846)
Eve [Thu, 26 Dec 2024 15:54:44 +0000 (10:54 -0500)]
vulkan: multi-row k quants (llama/10846)

* multi row k quant shaders!

* better row selection

* more row choices

* readjust row selection

* rm_kq=2 by default

5 months agoexamples, ggml : fix GCC compiler warnings (llama/10983)
Peter [Thu, 26 Dec 2024 13:59:11 +0000 (00:59 +1100)]
examples, ggml : fix GCC compiler warnings (llama/10983)

Warning types fixed (observed under MSYS2 GCC 14.2.0):
* format '%ld' expects argument of type 'long int', but argument has type 'size_t'
* llama.cpp/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp:81:46: warning: missing initializer for member '_STARTUPINFOA::lpDesktop' [-Wmissing-field-initializers]  (emitted for all struct field except first)

5 months agoggml : more perfo with llamafile tinyblas on x86_64 (llama/10714)
Djip007 [Tue, 24 Dec 2024 17:54:49 +0000 (18:54 +0100)]
ggml : more perfo with llamafile tinyblas on x86_64 (llama/10714)

* more perfo with llamafile tinyblas on x86_64.

- add bf16 suport
- change dispache strategie (thanks:
https://github.com/ikawrakow/ik_llama.cpp/pull/71 )
- reduce memory bandwidth

simple tinyblas dispache and more cache freindly

* tinyblas dynamic dispaching

* sgemm: add M blocs.

* - git 2.47 use short id of len 9.
- show-progress is not part of GNU Wget2

* remove not stable test

5 months agoggml : use wstring for backend search paths (llama/10960)
Diego Devesa [Tue, 24 Dec 2024 03:05:27 +0000 (04:05 +0100)]
ggml : use wstring for backend search paths (llama/10960)

ggml-ci

5 months agoggml : fix arm enabled features check (llama/10961)
Diego Devesa [Tue, 24 Dec 2024 03:05:17 +0000 (04:05 +0100)]
ggml : fix arm enabled features check (llama/10961)

5 months agoggml : fix const usage in SSE path (llama/10962)
Diego Devesa [Mon, 23 Dec 2024 19:25:52 +0000 (20:25 +0100)]
ggml : fix const usage in SSE path (llama/10962)

5 months agoggml : fix run-time on FreeBSD in get_executable_path() (llama/10948)
yuri@FreeBSD [Mon, 23 Dec 2024 00:20:11 +0000 (16:20 -0800)]
ggml : fix run-time on FreeBSD in get_executable_path() (llama/10948)

5 months agovulkan: build fixes for 32b (llama/10927)
Jeff Bolz [Sun, 22 Dec 2024 09:44:01 +0000 (03:44 -0600)]
vulkan: build fixes for 32b (llama/10927)

* vulkan: build fixes for 32b

Should fix #10923

* vulkan: initialize some buffer/offset variables

5 months agovulkan: optimize coopmat2 dequant functions (llama/10855)
Jeff Bolz [Sat, 21 Dec 2024 07:04:45 +0000 (01:04 -0600)]
vulkan: optimize coopmat2 dequant functions (llama/10855)

Change the code to do 16b loads when possible and extract the appropriate
component late, so the code is effectively decoding a pair of elements and
then selecting one. This can allow more commoning to happen in the compiler
when neighboring elements are loaded.

5 months agoggml-cpu: replace NEON asm with intrinsics in ggml_gemv_q4_0_4x8_q8_0() (llama/10874)
Adrien Gallouët [Fri, 20 Dec 2024 23:33:37 +0000 (00:33 +0100)]
ggml-cpu: replace NEON asm with intrinsics in ggml_gemv_q4_0_4x8_q8_0() (llama/10874)

* ggml-cpu: replace NEON asm with intrinsics in ggml_gemv_q4_0_4x8_q8_0()

Signed-off-by: Adrien Gallouët <redacted>
* ggml-cpu: format code

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

Signed-off-by: Adrien Gallouët <redacted>
5 months agoSYCL: Migrate away from deprecated ggml_tensor->backend (llama/10840)
Akarshan Biswas [Fri, 20 Dec 2024 15:31:28 +0000 (21:01 +0530)]
SYCL: Migrate away from deprecated ggml_tensor->backend (llama/10840)

* Migrate to tensor->buffer for checking backend buffer type: 1

* SYCL: common.cpp try to migrate away from tensor->backend

* SYCL: fix assertions and add proper comments

* SYCL: remove extra space

* SYCL: Add back static to ggml_backend_buffer_is_sycl_split function

* SYCL: Add pragma directive to suppress warning spam

* SYCL: Integrate debug logs with GGML_LOG and other fixes

* Revert "SYCL: Integrate debug logs with GGML_LOG and other fixes"

This reverts commit 2607b7de0f0d2f4f1f690226f86fa861aa39cb97.
Let's keep the current SYCL specific logging mechanism for now

* SYCL: Use GGML_SYCL_DEBUG after reverting

* SYCL: reg_get_proc_address func, update to the current func signature

* SYCL: Refactor SYCL buffer checks in ggml_sycl_cpy_tensor_2d

5 months agoggml : add test for SVE and disable when it fails (llama/10906)
Diego Devesa [Fri, 20 Dec 2024 12:31:28 +0000 (13:31 +0100)]
ggml : add test for SVE and disable when it fails (llama/10906)

5 months agoggml: fix arm build with gcc (llama/10895)
Adrien Gallouët [Thu, 19 Dec 2024 13:20:41 +0000 (14:20 +0100)]
ggml: fix arm build with gcc (llama/10895)

Signed-off-by: Adrien Gallouët <redacted>
5 months agoggml : fix arm build (llama/10890)
Diego Devesa [Wed, 18 Dec 2024 22:21:42 +0000 (23:21 +0100)]
ggml : fix arm build (llama/10890)

* ggml: GGML_NATIVE uses -mcpu=native on ARM

Signed-off-by: Adrien Gallouët <redacted>
* ggml: Show detected features with GGML_NATIVE

Signed-off-by: Adrien Gallouët <redacted>
* remove msvc support, add GGML_CPU_ARM_ARCH option

* disable llamafile in android example

* march -> mcpu, skip adding feature macros

ggml-ci

---------

Signed-off-by: Adrien Gallouët <redacted>
Co-authored-by: Adrien Gallouët <redacted>
5 months agotts : add OuteTTS support (llama/10784)
Georgi Gerganov [Wed, 18 Dec 2024 17:27:21 +0000 (19:27 +0200)]
tts : add OuteTTS support (llama/10784)

* server : add "tokens" output

ggml-ci

* server : output embeddings for all tokens when pooling = none

ggml-ci

* server : be explicit about the pooling type in the tests

ggml-ci

* server : do not normalize embeddings when there is no pooling

ggml-ci

* llama : add OuteTTS support (wip)

* wip

* extract features

* first conv

* group norm

* resnet conv

* resnet

* attn

* pos net

* layer norm

* convnext

* head

* hann window

* fix n_embd + remove llama.cpp hacks

* compute hann window

* fft

* spectrum processing

* clean-up

* tts : receive input text and generate codes

* clip : fix new conv name

* tts : minor fix

* tts : add header + minor fixes

ggml-ci

* tts : add matchematical constant

ggml-ci

* tts : fix sampling + cut initial noise

* tts : fixes

* tts : update default samplers

ggml-ci

* tts : text pre-processing

* tts : outetts-voc -> wavtokenizer-dec

* tts : remove hardcoded constants

ggml-ci

* tts : fix tensor shapes

* llama : refactor wavtokenizer tensors

ggml-ci

* cont

ggml-ci

* cont [no ci]

* llama : update WavTokenizer to non-causal attn

* llama : handle no-vocab detokenization

* tts : add Python example for OuteTTS (wip)

* tts : extend python example to generate spectrogram

ggml-ci

* server : fix rebase artifacts

* tts : enable "return_tokens" in Python example

ggml-ci

* tts : minor fixes

* common : support HF download for vocoder

5 months agotests: add tests for GGUF (llama/10830)
Johannes Gäßler [Tue, 17 Dec 2024 18:09:35 +0000 (19:09 +0100)]
tests: add tests for GGUF (llama/10830)

6 months agoggml : do not install metal source when embed library (#1054)
Georgi Gerganov [Mon, 23 Dec 2024 19:25:52 +0000 (21:25 +0200)]
ggml : do not install metal source when embed library (#1054)

ggml-ci

6 months agoggml : improve inputs log sched_print_assignments (#1053)
Daniel Bevenius [Thu, 19 Dec 2024 02:50:12 +0000 (03:50 +0100)]
ggml : improve inputs log sched_print_assignments (#1053)

This commit attempts to improve the log message for the inputs of the
splits in the sched_print_assignments function.

The motivation for this change is that currently even if there are no
inputs a colon is displayed at the end of the line, which can make it a
little confusing when reading the output as it could be interpreted as
the line below are inputs when they are in fact nodes. With this change
the colon will only be printed if there actually are inputs.

6 months agosync : whisper.cpp
Georgi Gerganov [Wed, 18 Dec 2024 11:03:17 +0000 (13:03 +0200)]
sync : whisper.cpp

6 months agofiles : remove old sources (#1052)
Georgi Gerganov [Wed, 18 Dec 2024 06:38:17 +0000 (08:38 +0200)]
files : remove old sources (#1052)

ggml-ci

6 months agosync : llama.cpp
Georgi Gerganov [Tue, 17 Dec 2024 16:36:44 +0000 (18:36 +0200)]
sync : llama.cpp

ggml-ci

6 months agoggml : update ggml_backend_cpu_device_supports_op (llama/10867)
Georgi Gerganov [Tue, 17 Dec 2024 16:35:42 +0000 (18:35 +0200)]
ggml : update ggml_backend_cpu_device_supports_op (llama/10867)

* ggml : fix cpy op for IQ-quants to use reference impl

ggml-ci

* ggml : disable tests involving i-matrix quantization

* ggml : update ggml_backend_cpu_device_supports_op

ggml-ci

6 months agovulkan: bugfixes for small subgroup size systems + llvmpipe test (llama/10809)
Eve [Tue, 17 Dec 2024 05:52:55 +0000 (05:52 +0000)]
vulkan: bugfixes for small subgroup size systems + llvmpipe test (llama/10809)

* ensure mul mat shaders work on systems with subgroup size less than 32

more fixes

add test

* only s_warptile_mmq needs to be run with 32 threads or more

6 months agorwkv6: add wkv6 support for Vulkan backend (llama/10829)
Zhiyuan Li [Mon, 16 Dec 2024 21:00:46 +0000 (05:00 +0800)]
rwkv6: add wkv6 support for Vulkan backend (llama/10829)

* rwkv_wkv6 vulkan shader

* RWKV_WKV6 Vulkan op tests passed

Signed-off-by: Molly Sophia <redacted>
* Apply code format changes

Signed-off-by: Molly Sophia <redacted>
* add [[unroll]] and remove unnecessary conditions

* add uma support

* fix erros in EditorConfig Checker

---------

Signed-off-by: Molly Sophia <redacted>
Co-authored-by: Molly Sophia <redacted>
6 months agollama : add Qwen2VL support + multimodal RoPE (llama/10361)
HimariO [Sat, 14 Dec 2024 12:43:46 +0000 (20:43 +0800)]
llama : add Qwen2VL support + multimodal RoPE (llama/10361)

* Barebone Qwen2VL LLM convertor

* Add Qwen2VL cli entrypoint

* [WIP] add qwen2vl arch

* Verify m-rope output

* Add vl-rope/2d-rope support for qwen2vl ViT

* update qwen2vl cli tool

* update 5D tensor op workaround

* [WIP] qwen2vl vision model

* make batch and clip utils compatible with qwen2vl

* [WIP] create inference workflow, gguf convert script but fix

* correcting vision-rope behavior, add the missing last layer back to ViT

* add arg parser to qwen2vl_surgery

* replace variable size array with vector

* cuda-gdb cmake preset

* add fp32 mrope, vision rope kernel

* add fp16 support for qwen2vl and m-rope

* add `GGML_ROPE_TYPE_MROPE`, `GGML_ROPE_TYPE_VISION`

* fix rope op mode switching, out dated func args

* update `llama_hparams`

* update to keep up stream changes

* resolve linter, test errors

* add makefile entry, update speical image padding token

* add mrope unit test, fix few compiler warnings

* rename `mrope` related function, params

* minor updates on debug util, bug fixs

* add `m-rope` testcase to `test-backend-ops`

* Apply suggestions from code review

Co-authored-by: Georgi Gerganov <redacted>
* fix traililng whitespce

* store `llama_hparams.rope_sections` with fixed size array

* update position id tensor size check in GGML_OP_ROPE

* minor updates

* update `ggml_backend_*_supports_op` of unsupported backends

* remote old `rope_section` compare operator

---------

Co-authored-by: Georgi Gerganov <redacted>
6 months agoIntroducing experimental OpenCL backend with support for Qualcomm Adreno GPUs (llama...
lhez [Fri, 13 Dec 2024 20:23:52 +0000 (12:23 -0800)]
Introducing experimental OpenCL backend with support for Qualcomm Adreno GPUs (llama/10693)

* [cl][adreno] Add Adreno GPU support

Add new OpenCL backend to support Adreno GPUs

---------

Co-authored-by: Skyler Szot <redacted>
Co-authored-by: Shangqing Gu <redacted>
Co-authored-by: Alexander Angus <redacted>
Co-authored-by: Hongqiang Wang <redacted>
Co-authored-by: Max Krasnyansky <redacted>
* [cl][ci] Add workflow for CL

* [cl][adreno] Fix memory leak for non SMALL_ALLOC path

* opencl: integrate backend dyn.load interface and fix compiler and format warnings

* opencl: remove small-alloc support and fix build errors for non-opencl platforms

* opencl: fixed merge conflict (MUSA added twice in cmake)

* opencl-ci: use RUNNER_TEMP instead of github.workspace

* opencl: fix embed tool invocation with python3

* opencl: CI workflow fixes

* opencl: Clean up small-alloc in CMake files

* opencl: cleanup ggml-opencl2 header file

* opencl: use ulong for offsets and strides in ADD kernel

* opencl: use cl_ulong for all offsets

* opencl: use cl_ulong for sizes and strides

* opencl: use `GGML_LOG_xxx` instead of `fprintf(stderr, ...)`

* opencl: rename backend `opencl2` -> `opencl`

* opencl: rename kernel files `ggml-opencl2` -> `ggml-opencl`

* opencl: make OpenCL required, remove redundant lib and inc directories

* `ggml-base`, `..` and `.` are added by `ggml_add_backend_library`

* opencl: rename backend - funcs, structs, etc `opencl2` -> `opencl`

* opencl: remove copyright marker since main license already covers

* opencl: replace some more OPENCL2 leftovers

* opencl: remove limits on `tensor_extra`

* opencl: use pools for `tensor_extra`

* opencl: fix compiler warnings with GCC and Clang

Still getting the warning about clCreateCmdQueue being obsolete.
Will fix that separately.

* opencl: fail gracefully if opencl devices are not available

Also for unsupported GPUs.

* opencl: fix MSVC builds (string length error)

* opencl: check for various requirements, allow deprecated API

* opencl: update log message for unsupported GPUs

---------

Co-authored-by: Skyler Szot <redacted>
Co-authored-by: Shangqing Gu <redacted>
Co-authored-by: Alexander Angus <redacted>
Co-authored-by: Hongqiang Wang <redacted>
Co-authored-by: Max Krasnyansky <redacted>
6 months agoFix crash caused by ggml_backend_load_all when launching on Android Activity (llama...
谢乃闻 [Fri, 13 Dec 2024 12:56:07 +0000 (12:56 +0000)]
Fix crash caused by ggml_backend_load_all when launching on Android Activity (llama/10812)

* Fix crash caused by ggml_backend_load_all when launching on AndroidActivity.

Details:
Calling ggml_backend_load_all during initialization in the AndroidActivity project leads to a crash with the error:
terminating with uncaught exception of type std::__ndk1::__fs::filesystem::filesystem_error: filesystem error: in directory_iterator::directory_iterator(...): Permission denied [./].
This issue occurs because AndroidActivity restricts file access due to sandboxing.

Reproduction:
In the example folder, the LlamaAndroid project can reproduce the crash by calling ggml_backend_load_all first in Java_android_llama_cpp_LLamaAndroid_backend_1init.

* Update ggml/src/ggml-backend-reg.cpp

---------

Co-authored-by: Diego Devesa <redacted>
6 months agovulkan: small mul_mat_vec optimizations (llama/10665)
Eve [Fri, 13 Dec 2024 08:42:04 +0000 (08:42 +0000)]
vulkan: small mul_mat_vec optimizations (llama/10665)

* double the number of rows per workgroup

* Update ggml-vulkan.cpp

* Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats

* only increase the number of rows for amd and subgroup size 64

* fix missing NUM_ROWS for mul_mat_vec_iq4_nl_f16_f32, untested

* use subgroup min and max to check for gcn (requires https://github.com/ggerganov/llama.cpp/pull/10721)

* manual merge ggml-vulkan.cpp

* set min and max subgroup size in any case

* Also double the number of rows for Intel GPUs

6 months agoSYCL: Reduce most of the compiler warnings (llama/10748)
Akarshan Biswas [Fri, 13 Dec 2024 06:42:15 +0000 (12:12 +0530)]
SYCL: Reduce most of the compiler warnings (llama/10748)

* Try to reduce some unused and typecast warnings

* Reduce compiler warnings step 2

* add a newline at the end of the file

* Initialize nreduce as size_t

* [SYCL] Remove pragma directives from mmq.cpp

* SYCL: mmq add condition to prevent blocks_per_tile_x_row variable from becoming 0

* SYCL softmax: Initialize nreduce as size_t

* ggml-sycl.cpp: fix some trailing whitespaces

* SYCL: remove the unused variables instead of commenting it out

* SYCL poo2d kernel: set NAN for invalid pooling op

* SYCL gemm.hpp: remove pragma directives

* SYCL gemm.hpp: use const cast to properly support dnnl::memory

* SYCL: wkv6 remove a comment

* SYCL: clean comments step 2

* SYCL: clean comments and variables step 3

* SYCL: Use GGML_UNUSED for unused variables

* SYCL: remove extra empty lines and a comment

* Remove TODO

* cleanup spaces

* add a stdout for unsupported op

* use sycl printf over fprintf

* remove prints for CI

* SYCL ggml-sycl: pool2D use sycl::nan and remove if-else block

---------

Co-authored-by: Abhilash Majumder <redacted>
6 months agoggml : Fix compilation issues on ARM platform when building without fp16 (llama/10811)
Karol Kontny [Fri, 13 Dec 2024 00:04:19 +0000 (01:04 +0100)]
ggml : Fix compilation issues on ARM platform when building without fp16 (llama/10811)

6 months agoCUDA: faster non-contiguous concat (llama/10760)
a3sh [Thu, 12 Dec 2024 18:09:50 +0000 (02:09 +0800)]
CUDA: faster non-contiguous concat (llama/10760)

* faster uncontiguous concat

* Use a lambda to avoid code duplication

Co-authored-by: Diego Devesa <redacted>
* Update ggml/src/ggml-cuda/concat.cu

* add constexpr  and static assert

---------

Co-authored-by: Diego Devesa <redacted>
6 months agoremove CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS (llama/10797)
Diego Devesa [Thu, 12 Dec 2024 18:02:49 +0000 (19:02 +0100)]
remove CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS (llama/10797)

other windows build fixes

6 months agoVulkan: Use improved q4_k and q5_k dequant code in dequant shaders (llama/10798)
0cc4m [Thu, 12 Dec 2024 17:36:00 +0000 (18:36 +0100)]
Vulkan: Use improved q4_k and q5_k dequant code in dequant shaders (llama/10798)

6 months agoVulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmat...
0cc4m [Thu, 12 Dec 2024 17:35:37 +0000 (18:35 +0100)]
Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats (llama/10721)

* Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats

* Fix subgroup size control extension support check

Add accf32 and accf16 checks for coopmats

* Also disable coopmats on amdvlk

6 months agoggml: load all backends from a user-provided search path (llama/10699)
Gilad S [Wed, 11 Dec 2024 00:47:21 +0000 (02:47 +0200)]
ggml: load all backends from a user-provided search path (llama/10699)

* feat: load all backends from a user-provided search path

* fix: Windows search path

* refactor: rename `ggml_backend_load_all_in_search_path` to `ggml_backend_load_all_from_path`

* refactor: rename `search_path` to `dir_path`

* fix: change `NULL` to `nullptr`

Co-authored-by: Diego Devesa <redacted>
* fix: change `NULL` to `nullptr`

---------

Co-authored-by: Diego Devesa <redacted>
6 months agovulkan: request round-to-even for fp16 in im2col/rope_head (llama/10767)
Jeff Bolz [Tue, 10 Dec 2024 20:23:17 +0000 (14:23 -0600)]
vulkan: request round-to-even for fp16 in im2col/rope_head (llama/10767)

Vulkan doesn't mandate a specific rounding mode, but the shader_float_controls
feature allows rounding mode to be requested if the implementation supports it.

6 months agovulkan: dynamic subgroup size for the remaining k quants (llama/10745)
Eve [Tue, 10 Dec 2024 19:33:23 +0000 (19:33 +0000)]
vulkan: dynamic subgroup size for the remaining k quants (llama/10745)

* q5_k

q4_k

q3_k

q2_k

q6_k multi row example

* revert as multi row isnt faster for k quants

6 months agoCUDA: rename macros to avoid conflicts with WinAPI (llama/10736)
Andreas Kieslinger [Tue, 10 Dec 2024 17:23:24 +0000 (18:23 +0100)]
CUDA: rename macros to avoid conflicts with WinAPI (llama/10736)

* Renames NVIDIA GPU-architecture flags to avoid name clashes with WinAPI. (e.g. CC_PASCAL, GPU architecture or WinAPI pascal compiler flag?)

* Reverts erroneous rename in SYCL-code.

* Renames GGML_CUDA_MIN_CC_DP4A to GGML_CUDA_CC_DP4A.

* Renames the rest of the compute capability macros for consistency.

6 months agovulkan: disable spirv-opt for coopmat shaders (llama/10763)
Jeff Bolz [Tue, 10 Dec 2024 17:22:20 +0000 (11:22 -0600)]
vulkan: disable spirv-opt for coopmat shaders (llama/10763)

There are some bugs in the 1.3.296 SDK, so disable this. It isn't strictly
necessary anyway.

Add missing dependency on vulkan-shaders-gen, so shaders get recompiled when it
changes.

Fix coopmat support reporting when glslc doesn't support NV_coopmat2.

6 months agosync : whisper.cpp
Georgi Gerganov [Tue, 17 Dec 2024 16:34:54 +0000 (18:34 +0200)]
sync : whisper.cpp

6 months agocmake : fix "amd64" processor string (whisper/2638)
Georgi Gerganov [Tue, 17 Dec 2024 16:34:32 +0000 (18:34 +0200)]
cmake : fix "amd64" processor string (whisper/2638)

6 months agovulkan : fix soft_max.comp division by zero (whisper/2633)
gn64 [Mon, 16 Dec 2024 10:34:38 +0000 (19:34 +0900)]
vulkan : fix soft_max.comp division by zero (whisper/2633)

This change prevents a division by zero error when p.KY is 0.

6 months agoggml : remove return from ggml_gallocr_allocate_node (#1048)
Daniel Bevenius [Sat, 14 Dec 2024 02:23:08 +0000 (03:23 +0100)]
ggml : remove return from ggml_gallocr_allocate_node (#1048)

This commit removes the return statement from ggml_gallocr_allocate_node
function.

The motivation behind this change is to make the code more readable and
consistent.

6 months agoggml : add check for grad_accs (#1046)
Daniel Bevenius [Fri, 13 Dec 2024 07:19:38 +0000 (08:19 +0100)]
ggml : add check for grad_accs (#1046)

* ggml : add check for grad_accs

This commit adds a check for grad_accs in ggml_graph_get_grad and
ggml_graph_get_grad_acc functions. This is necessary to avoid segfaults
when grad_accs is not initialized.

The motivation for this change is that I find it nice to be able to
print out a computation graph using ggml_graph_print but this function
segfaults when grad_accs is not initialized:
```console
(gdb) p g1
$2 = (ggml_cgraph *) 0x7ffff66004b0
(gdb) p *g1
$3 = {size = 2048, n_nodes = 1, n_leafs = 2, nodes = 0x7ffff6600500,
grads = 0x0, grad_accs = 0x0, leafs = 0x7ffff6604500,
visited_hash_set = {size = 4099, used = 0x7ffff6610518,
keys = 0x7ffff6608500}, order = GGML_CGRAPH_EVAL_ORDER_LEFT_TO_RIGHT}
(gdb) p ggml_graph_print(g1)
=== GRAPH ===
n_nodes = 1

Program received signal SIGSEGV, Segmentation fault.
0x0000555555579775 in ggml_graph_get_grad
(cgraph=0x7ffff66004b0,node=0x7ffff6600340)
    at /ggml/ggml/src/ggml.c:5990
5990  return igrad != GGML_HASHSET_FULL &&
          ggml_bitset_get(cgraph->visited_hash_set.used, igrad) ?
          cgraph->grads[igrad] : NULL;
```

* squash! ggml : add check for grad_accs

Fix the check in ggml_graph_get_grad. The check was incorrectly using
cgraph->grad_accs instead of cgraph->grads.

6 months agocommon : remove old types
Georgi Gerganov [Tue, 10 Dec 2024 15:19:09 +0000 (17:19 +0200)]
common : remove old types

ggml-ci

6 months agosync : llama.cpp
Georgi Gerganov [Tue, 10 Dec 2024 15:15:53 +0000 (17:15 +0200)]
sync : llama.cpp

ggml-ci

6 months agoCUDA: fix shared memory access condition for mmv (llama/10740)
Johannes Gäßler [Mon, 9 Dec 2024 19:07:12 +0000 (20:07 +0100)]
CUDA: fix shared memory access condition for mmv (llama/10740)

6 months agovulkan: fix compile warnings (llama/10731)
Jeff Bolz [Mon, 9 Dec 2024 07:24:01 +0000 (01:24 -0600)]
vulkan: fix compile warnings (llama/10731)

6 months agoVulkan: fix NaN in tanh.comp with AMD proprietary driver on Windows (llama/10723)
stduhpf [Sun, 8 Dec 2024 18:19:19 +0000 (19:19 +0100)]
Vulkan: fix NaN in tanh.comp with AMD proprietary driver on Windows (llama/10723)

* Vulkan: fix NaN in tanh.comp

* Faster NaN-free tanh

6 months agovulkan: compile a test shader in cmake to check for coopmat2 support (llama/10713)
Jeff Bolz [Sun, 8 Dec 2024 08:05:55 +0000 (02:05 -0600)]
vulkan: compile a test shader in cmake to check for coopmat2 support (llama/10713)

6 months agoggml : disable iq4_nl interleave size 8 (llama/10709)
Georgi Gerganov [Sat, 7 Dec 2024 16:38:15 +0000 (18:38 +0200)]
ggml : disable iq4_nl interleave size 8 (llama/10709)

ggml-ci

6 months agoggml : refactor online repacking (llama/10446)
Djip007 [Sat, 7 Dec 2024 12:37:50 +0000 (13:37 +0100)]
ggml : refactor online repacking (llama/10446)

* rename ggml-cpu-aarch64.c to .cpp

* reformat extra cpu backend.

- clean Q4_0_N_M and IQ4_0_N_M
  - remove from "file" tensor type
  - allow only with dynamic repack

- extract cpu extra bufts and convert to C++
  - hbm
  - "aarch64"

- more generic use of extra buffer
  - generalise extra_supports_op
  - new API for "cpu-accel":
     - amx
     - aarch64

* clang-format

* Clean Q4_0_N_M ref

Enable restrict on C++

* add op GGML_OP_MUL_MAT_ID for Q4_0_N_M with runtime repack

* added/corrected control on tensor size for Q4 repacking.

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

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

Co-authored-by: Georgi Gerganov <redacted>
* add debug logs on repacks.

---------

Co-authored-by: Georgi Gerganov <redacted>
6 months agoVulkan: VK_KHR_cooperative_matrix support to speed up prompt processing (llama/10597)
0cc4m [Sat, 7 Dec 2024 09:24:15 +0000 (10:24 +0100)]
Vulkan: VK_KHR_cooperative_matrix support to speed up prompt processing (llama/10597)

* Vulkan: Implement VK_KHR_cooperative_matrix support in the matrix matrix multiplication shader

* Improve performance with better q4_k and q5_k dequant and store unrolling

* Add Vulkan MUL_MAT and MUL_MAT_ID accumulator precision selection

* Rework mulmat shader selection and compilation logic, avoid compiling shaders that won't get used by device

* Vulkan: Implement accumulator switch for specific mul mat mat shaders

* Vulkan: Unroll more loops for more mul mat mat performance

* Vulkan: Add VK_AMD_shader_core_properties2 support to read Compute Unit count for split_k logic

* Disable coopmat support on AMD proprietary driver

* Remove redundant checks

* Add environment variable GGML_VK_DISABLE_COOPMAT to disable VK_KHR_cooperative_matrix support

* Fix rebase typo

* Fix coopmat2 MUL_MAT_ID pipeline selection

6 months agometal : Extend how Llama.cpp locates metal resources (llama/10676)
Robert Ormandi [Sat, 7 Dec 2024 07:55:01 +0000 (01:55 -0600)]
metal : Extend how Llama.cpp locates metal resources (llama/10676)

* metal : Extend how Llama.cpp locates metal resources (llama/10675)

  * It searches the resource file in the directory where the current
    binary is located as well.
  * Resolves symbolic links.

Rationale:

When we plug this dependency into a Bazel build and run it in the
context of Bazel (e.g. testing):

  * the execution directory is often very different from where the files
    are located and no direct control over this (Bazel sandboxing),
  * the Bazel sandbox often use symbolic links to make files available.

With this patch, we can have the resource file added to the target,
can build and run tests in the context of Bazel.

* Update ggml/src/ggml-metal/ggml-metal.m

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

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

Co-authored-by: Georgi Gerganov <redacted>
6 months agovulkan: Add VK_NV_cooperative_matrix2 support for mul_mat and flash attention (llama...
Jeff Bolz [Thu, 5 Dec 2024 19:15:05 +0000 (13:15 -0600)]
vulkan: Add VK_NV_cooperative_matrix2 support for mul_mat and flash attention (llama/10206)

6 months agosync : whisper.cpp
Georgi Gerganov [Sun, 8 Dec 2024 20:56:43 +0000 (22:56 +0200)]
sync : whisper.cpp

ggml-ci

6 months agoscripts : update whisper sync
Georgi Gerganov [Sun, 8 Dec 2024 20:55:16 +0000 (22:55 +0200)]
scripts : update whisper sync

6 months agosync : llama.cpp
Georgi Gerganov [Thu, 5 Dec 2024 11:28:22 +0000 (13:28 +0200)]
sync : llama.cpp

ggml-ci

6 months agoggml : add predefined list of CPU backend variants to build (llama/10626)
Diego Devesa [Wed, 4 Dec 2024 13:45:40 +0000 (14:45 +0100)]
ggml : add predefined list of CPU backend variants to build (llama/10626)

* ggml : add predefined list of CPU backend variants to build

* update CPU dockerfiles

6 months agoggml-cpu : fix HWCAP2_I8MM value (llama/10646)
Diego Devesa [Wed, 4 Dec 2024 13:40:44 +0000 (14:40 +0100)]
ggml-cpu : fix HWCAP2_I8MM value (llama/10646)