]> git.djapps.eu Git - pkg/ggml/sources/whisper.cpp/log
pkg/ggml/sources/whisper.cpp
7 months agoggml : add asserts for type conversion in fattn kernels (llama/9971)
Georgi Gerganov [Mon, 21 Oct 2024 13:20:46 +0000 (16:20 +0300)]
ggml : add asserts for type conversion in fattn kernels (llama/9971)

ggml-ci

7 months agorpc : pack only RPC structs (llama/9959)
Radoslav Gerganov [Mon, 21 Oct 2024 10:35:40 +0000 (13:35 +0300)]
rpc : pack only RPC structs (llama/9959)

7 months agofix mul_mat_vec_q and *_vec_q error (llama/9939)
Neo Zhang Jianyu [Mon, 21 Oct 2024 06:26:09 +0000 (14:26 +0800)]
fix mul_mat_vec_q and *_vec_q error (llama/9939)

Co-authored-by: arthw <redacted>
7 months agorpc : backend refactoring (llama/9912)
Radoslav Gerganov [Fri, 18 Oct 2024 11:33:58 +0000 (14:33 +0300)]
rpc : backend refactoring (llama/9912)

* rpc : refactor backend

Use structs for RPC request/response messages

* rpc : refactor server

7 months agoAdd SYCL Backend registry, device and Event Interfaces (llama/9705)
Ouadie EL FAROUKI [Fri, 18 Oct 2024 05:46:16 +0000 (06:46 +0100)]
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

7 months agoadd amx kernel for gemm (llama/8998)
Ma Mingfei [Fri, 18 Oct 2024 05:34:36 +0000 (13:34 +0800)]
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

7 months agovulkan : add backend registry / device interfaces (llama/9721)
Diego Devesa [Thu, 17 Oct 2024 00:46:58 +0000 (02:46 +0200)]
vulkan : add backend registry / device interfaces (llama/9721)

* vulkan : add backend registry / device interfaces

* llama : print devices used on model load

7 months agofix: allocating CPU buffer with size `0` (llama/9917)
Gilad S [Wed, 16 Oct 2024 23:34:22 +0000 (02:34 +0300)]
fix: allocating CPU buffer with size `0` (llama/9917)

7 months agofix: use `vm_allocate` to allocate CPU backend buffer on macOS (llama/9875)
Gilad S [Wed, 16 Oct 2024 22:36:51 +0000 (01:36 +0300)]
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

7 months agoCUDA: fix 1D im2col, add tests (ggml/993)
Johannes Gäßler [Fri, 18 Oct 2024 07:24:44 +0000 (09:24 +0200)]
CUDA: fix 1D im2col, add tests (ggml/993)

7 months agoFix cann compilation error (llama/9891)
leo-pony [Wed, 16 Oct 2024 00:51:46 +0000 (08:51 +0800)]
Fix cann compilation error (llama/9891)

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

7 months agoVectorize load instructions in dmmv f16 CUDA kernel (llama/9816)
agray3 [Mon, 14 Oct 2024 00:49:08 +0000 (01:49 +0100)]
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>
7 months agoggml : move more prints to the ggml log system (llama/9839)
Diego Devesa [Fri, 11 Oct 2024 13:34:45 +0000 (15:34 +0200)]
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

7 months agorpc : add backend registry / device interfaces (llama/9812)
Diego Devesa [Thu, 10 Oct 2024 18:14:55 +0000 (20:14 +0200)]
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

7 months agomusa: add docker image support (llama/9685)
R0CKSTAR [Thu, 10 Oct 2024 18:10:37 +0000 (02:10 +0800)]
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>
7 months agoggml : fix BLAS with unsupported types (llama/9775)
Diego Devesa [Tue, 8 Oct 2024 12:21:43 +0000 (14:21 +0200)]
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

7 months agoggml : add backend registry / device interfaces to BLAS backend (llama/9752)
Diego Devesa [Mon, 7 Oct 2024 19:55:08 +0000 (21:55 +0200)]
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

7 months agoUpdate building for Android (llama/9672)
Andrew Minh Nguyen [Mon, 7 Oct 2024 16:37:31 +0000 (09:37 -0700)]
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

7 months agoggml : add metal backend registry / device (llama/9713)
Georgi Gerganov [Mon, 7 Oct 2024 15:27:51 +0000 (18:27 +0300)]
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>
7 months agometal : single allocation of encode_async block (llama/9747)
Paul Tsochantaris [Mon, 7 Oct 2024 12:26:31 +0000 (13:26 +0100)]
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>
7 months agoggml-alloc : remove buffer_id from leaf_alloc (ggml/987)
Daniel Bevenius [Wed, 9 Oct 2024 14:40:35 +0000 (16:40 +0200)]
ggml-alloc : remove buffer_id from leaf_alloc (ggml/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.

7 months agoscripts : sync amx
Georgi Gerganov [Thu, 31 Oct 2024 20:13:24 +0000 (22:13 +0200)]
scripts : sync amx

7 months agoggml : alloc ggml_contexts on the heap (#2525)
Georgi Gerganov [Thu, 31 Oct 2024 20:00:09 +0000 (22:00 +0200)]
ggml : alloc ggml_contexts on the heap (#2525)

* whisper : reduce ggml_context usage

* ggml : allocate contexts on the heap (v2)

* ggml : aligned malloc -> malloc

7 months agoci : fix openblas build (#2511)
Georgi Gerganov [Wed, 30 Oct 2024 10:58:26 +0000 (12:58 +0200)]
ci : fix openblas build (#2511)

* ci : fix openblas build

* cont : would this work?

* ci : I'm sorry, windows

* cont : disabled wrong build

* ci : fix openblas build with pkgconfiglite (#2517)

- choco install pkgconfiglite (vcpkg-pkgconf doesn't contain pkg-config executable?)
- vcpkg install openblas (otherwise it is not detected now)

---------

Co-authored-by: Tamotsu Takahashi <redacted>
7 months agoscripts : add turbo-q8_0 to the benchmark
Georgi Gerganov [Tue, 29 Oct 2024 17:37:24 +0000 (19:37 +0200)]
scripts : add turbo-q8_0 to the benchmark

7 months agowhisper : minor compile warning
Georgi Gerganov [Tue, 29 Oct 2024 17:27:52 +0000 (19:27 +0200)]
whisper : minor compile warning

8 months agowhisper : move new-segment callback after DTW step (#2515)
jettoblack [Tue, 29 Oct 2024 06:47:21 +0000 (02:47 -0400)]
whisper : move new-segment callback after DTW step (#2515)

8 months agoruby : fix installation test (#2519)
KITAITI Makoto [Tue, 29 Oct 2024 06:45:37 +0000 (15:45 +0900)]
ruby : fix installation test (#2519)

8 months agoruby : add more APIs (#2518)
KITAITI Makoto [Mon, 28 Oct 2024 17:23:23 +0000 (02:23 +0900)]
ruby : add more APIs (#2518)

* Add test for built package existence

* Add more tests for Whisper::Params

* Add more Whisper::Params attributes

* Add tests for callbacks

* Add progress and abort callback features

* [skip ci] Add prompt usage in README

* Change prompt text in example

8 months agoruby : support new-segment callback (#2506)
KITAITI Makoto [Mon, 28 Oct 2024 13:43:27 +0000 (22:43 +0900)]
ruby : support new-segment callback (#2506)

* Add Params#new_segment_callback= method

* Add tests for Params#new_segment_callback=

* Group tests for #transcribe

* Don't use static for thread-safety

* Set new_segment_callback only when necessary

* Remove redundant check

* [skip ci] Add Ruby version README

* Revert "Group tests for #transcribe"

This reverts commit 71b65b00ccf1816c9ea8a247fb30f71bc09707d3.

* Revert "Add tests for Params#new_segment_callback="

This reverts commit 81e6df3bab7662da5379db51f28a989db7408c02.

* Add test for Context#full_n_segments

* Add Context#full_n_segments

* Add tests for lang API

* Add lang API

* Add tests for Context#full_lang_id API

* Add Context#full_lang_id

* Add abnormal test cases for lang

* Raise appropriate errors from lang APIs

* Add tests for Context#full_get_segment_t{0,1} API

* Add Context#full_get_segment_t{0,1}

* Add tests for Context#full_get_segment_speaker_turn_next API

* Add Context#full_get_segment_speaker_turn_next

* Add tests for Context#full_get_segment_text

* Add Context#full_get_setgment_text

* Add tests for Params#new_segment_callback=

* Run new segment callback

* Split tests to multiple files

* Use container struct for new segment callback

* Add tests for Params#new_segment_callback_user_data=

* Add Whisper::Params#new_user_callback_user_data=

* Add GC-related test for new segment callback

* Protect new segment callback related structs from GC

* Add meaningful test for build

* Rename: new_segment_callback_user_data -> new_segment_callback_container

* Add tests for Whisper::Segment

* Add Whisper::Segment and Whisper::Context#each_segment

* Extract c_ruby_whisper_callback_container_allocate()

* Add test for Whisper::Params#on_new_segment

* Add Whisper::Params#on_new_egment

* Assign symbol IDs to variables

* Make extsources.yaml simpler

* Update README

* Add document comments

* Add test for calling Whisper::Params#on_new_segment multiple times

* Add file dependencies to GitHub actions config and .gitignore

* Add more files to ext/.gitignore

8 months agoruby : add Metal support (#2516)
KITAITI Makoto [Mon, 28 Oct 2024 11:08:09 +0000 (20:08 +0900)]
ruby : add Metal support (#2516)

8 months agowhisper : fix index overflow in token-level timestamp logic (#2505)
Josscii [Wed, 23 Oct 2024 12:14:03 +0000 (20:14 +0800)]
whisper : fix index overflow in token-level timestamp logic (#2505)

8 months agoreadme : update links and make commands (#2489)
toboil-features [Thu, 17 Oct 2024 10:25:18 +0000 (13:25 +0300)]
readme : update links and make commands (#2489)

* Update links to headers in README.md

* Add link to Vulkan section in README.md

* Add "-j" for parallelism for "make" in README.md

* Update README.md

8 months agoruby : fix bindings (#2484)
KITAITI Makoto [Wed, 16 Oct 2024 15:44:04 +0000 (00:44 +0900)]
ruby : fix bindings (#2484)

* Improve Rakefile

* Remove intermediate files

* Remove unnecessary manipulations from extconf.rb

* Add README and LINCENSE to source files

* Manage ext source files using YAML file

* Use extsources.yaml to include files into gem package file

* Add git-managed source files to build dependency

* Add test task

* Download model for test if not exists

* Add test for build

* Ignore gem package directory

* Enable GitHub action for Ruby binding

* Fix model name

* Build lib file for test

* Use extension for each platform

* Use extension for each platform on testing

* Move built lib file rather than copy

* Add intermediate files to clean targets

8 months agoreadme : add Vulkan notice (#2488)
toboil-features [Wed, 16 Oct 2024 15:43:26 +0000 (18:43 +0300)]
readme : add Vulkan notice (#2488)

* Add Vulkan notice in README.md

* Fix formatting for Vulkan section in README.md

* Fix formatting in README.md

8 months agomake : fix GGML_VULKAN=1 build (#2485)
Georgi Gerganov [Wed, 16 Oct 2024 15:42:47 +0000 (18:42 +0300)]
make : fix GGML_VULKAN=1 build (#2485)

8 months agowhisper : add dtw preset for large-v3-turbo (#2481)
Rotem Dan [Tue, 15 Oct 2024 18:00:21 +0000 (21:00 +0300)]
whisper : add dtw preset for large-v3-turbo (#2481)

8 months agoconvert : handle max_target_positions (#2477)
CrispStrobe [Mon, 14 Oct 2024 07:46:33 +0000 (09:46 +0200)]
convert : handle max_target_positions (#2477)

as needed eg for
https://huggingface.co/primeline/whisper-large-v3-turbo-german/blob/main/config.json

8 months agoreadme : update the Quick Start section (#2475)
Salman Faroz [Mon, 14 Oct 2024 07:44:57 +0000 (13:14 +0530)]
readme : update the Quick Start section (#2475)

navigating into the directory

8 months agowhisper : add OpenVINO init with state (#2464)
Sandro Hanea [Tue, 8 Oct 2024 17:08:00 +0000 (19:08 +0200)]
whisper : add OpenVINO init with state (#2464)

* Fixed OpenVino init on state

* Removed an empty line

* Fixed typo

* Replaced tabs with spaces

---------

Co-authored-by: Sandro Hanea <redacted>
8 months agorelease : v1.7.1
Georgi Gerganov [Mon, 7 Oct 2024 10:06:48 +0000 (13:06 +0300)]
release : v1.7.1

8 months agovulkan : retry allocation with fallback flags (#2451)
SRHMorris [Sun, 6 Oct 2024 07:34:20 +0000 (08:34 +0100)]
vulkan : retry allocation with fallback flags (#2451)

Co-authored-by: Samuel Morris <redacted>
8 months agorelease : v1.7.0
Georgi Gerganov [Sat, 5 Oct 2024 13:43:26 +0000 (16:43 +0300)]
release : v1.7.0

8 months agoscripts : bench v3-turbo
Georgi Gerganov [Sat, 5 Oct 2024 13:22:53 +0000 (16:22 +0300)]
scripts : bench v3-turbo

8 months agowhisper : remove mel leftover constants (396089f)
Georgi Gerganov [Sat, 5 Oct 2024 13:13:03 +0000 (16:13 +0300)]
whisper : remove mel leftover constants (396089f)

8 months agowhisper : zero-out the KV cache upon clear (#2445)
Georgi Gerganov [Sat, 5 Oct 2024 12:22:17 +0000 (15:22 +0300)]
whisper : zero-out the KV cache upon clear (#2445)

8 months agoobjc : fix build
Georgi Gerganov [Sat, 5 Oct 2024 12:18:50 +0000 (15:18 +0300)]
objc : fix build

8 months agometal : zero-init buffer contexts (#0)
Georgi Gerganov [Sat, 5 Oct 2024 11:33:54 +0000 (14:33 +0300)]
metal : zero-init buffer contexts (#0)

8 months agowhisper : revert mel-related changes (#0)
Georgi Gerganov [Sat, 5 Oct 2024 11:29:45 +0000 (14:29 +0300)]
whisper : revert mel-related changes (#0)

too much extra logic and complexity for small benefit

8 months agowhisper : adapt to latest ggml (skip) (#0)
Georgi Gerganov [Sat, 5 Oct 2024 10:14:03 +0000 (13:14 +0300)]
whisper : adapt to latest ggml (skip) (#0)

8 months agoggml : fix typo in example usage ggml_gallocr_new (ggml/984)
Daniel Bevenius [Fri, 4 Oct 2024 13:46:18 +0000 (15:46 +0200)]
ggml : fix typo in example usage ggml_gallocr_new (ggml/984)

8 months agoggml : fixes after sync (ggml/983)
Diego Devesa [Fri, 4 Oct 2024 06:41:40 +0000 (08:41 +0200)]
ggml : fixes after sync (ggml/983)

ggml : remove test-backend-buffer

ggml : fix CUDA build warnings

8 months agoggml-backend : add device and backend reg interfaces (llama/9707)
Diego Devesa [Thu, 3 Oct 2024 18:25:11 +0000 (21:25 +0300)]
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

8 months agoFixed dequant precision issues in Q4_1 and Q5_1 (llama/9711)
Ouadie EL FAROUKI [Thu, 3 Oct 2024 06:50:44 +0000 (07:50 +0100)]
Fixed dequant precision issues in Q4_1 and Q5_1 (llama/9711)

8 months agoggml-backend : add device and backend reg interfaces (llama/9707)
Diego Devesa [Wed, 2 Oct 2024 23:49:47 +0000 (01:49 +0200)]
ggml-backend : add device and backend reg interfaces (llama/9707)

Co-authored-by: Johannes Gäßler <redacted>
8 months agoInitial cmake support of SYCL for AMD GPUs (llama/9658)
Alberto Cabrera Pérez [Wed, 2 Oct 2024 12:57:18 +0000 (13:57 +0100)]
Initial cmake support of SYCL for AMD GPUs (llama/9658)

sycl: initial cmake support of SYCL for AMD GPUs

8 months agovulkan : do not use tensor->extra (llama/9407)
Radoslav Gerganov [Wed, 2 Oct 2024 10:49:16 +0000 (13:49 +0300)]
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>
8 months agoggml/ex: calculate accuracy in graph, adapt MNIST (ggml/980)
Johannes Gäßler [Thu, 3 Oct 2024 15:29:59 +0000 (17:29 +0200)]
ggml/ex: calculate accuracy in graph, adapt MNIST (ggml/980)

8 months agoggml: refactor cross entropy loss CPU impl. (ggml/976)
Johannes Gäßler [Wed, 2 Oct 2024 13:32:39 +0000 (15:32 +0200)]
ggml: refactor cross entropy loss CPU impl. (ggml/976)

8 months agoscripts : sync ggml-backend.cpp
Georgi Gerganov [Sat, 5 Oct 2024 10:09:36 +0000 (13:09 +0300)]
scripts : sync ggml-backend.cpp

8 months agowhisper : fix excessive memory usage (#2443)
Georgi Gerganov [Sat, 5 Oct 2024 09:36:40 +0000 (12:36 +0300)]
whisper : fix excessive memory usage (#2443)

* whisper : fix KV cache allocation

* whisper : reduce memory overhead from unused input tensors

8 months agoexamples : update dr_wav.h to newer version (#2449)
Rahul Vadhyar [Fri, 4 Oct 2024 08:04:51 +0000 (13:34 +0530)]
examples : update dr_wav.h to newer version (#2449)

8 months agotalk-llama : sync llama.cpp
Georgi Gerganov [Wed, 2 Oct 2024 12:14:46 +0000 (15:14 +0300)]
talk-llama : sync llama.cpp

8 months agometal : reduce command encoding overhead (llama/9698)
Georgi Gerganov [Wed, 2 Oct 2024 12:12:16 +0000 (15:12 +0300)]
metal : reduce command encoding overhead (llama/9698)

8 months agosync : ggml
Georgi Gerganov [Wed, 2 Oct 2024 12:11:43 +0000 (15:11 +0300)]
sync : ggml

8 months agotest: fix OPT_STEP_ADAMW for test-backend-ops (ggml/974)
Johannes Gäßler [Mon, 30 Sep 2024 07:55:23 +0000 (09:55 +0200)]
test: fix OPT_STEP_ADAMW for test-backend-ops (ggml/974)

8 months agovulkan : mul_mat: fix UB with small warps (ggml/952)
Salvatore Mesoraca [Mon, 30 Sep 2024 07:14:09 +0000 (09:14 +0200)]
vulkan : mul_mat: fix UB with small warps (ggml/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>
8 months agoggml : fix ggml_cast (ggml/973)
Borislav Stanimirov [Mon, 30 Sep 2024 07:11:41 +0000 (10:11 +0300)]
ggml : fix ggml_cast (ggml/973)

8 months agoggml: fix gradient allocation logic (ggml/966)
Johannes Gäßler [Sun, 29 Sep 2024 21:18:02 +0000 (23:18 +0200)]
ggml: fix gradient allocation logic (ggml/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

8 months agoggml : define missing HWCAP flags (llama/9684)
Georgi Gerganov [Sun, 29 Sep 2024 18:18:23 +0000 (21:18 +0300)]
ggml : define missing HWCAP flags (llama/9684)

ggml-ci

Co-authored-by: Willy Tarreau <redacted>
8 months agoggml : add run-time detection of neon, i8mm and sve (llama/9331)
Dan Johansson [Sat, 28 Sep 2024 12:06:16 +0000 (14:06 +0200)]
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

8 months agoEnable use to the rebar feature to upload buffers to the device. (llama/9251)
Markus Tavenrath [Sat, 28 Sep 2024 10:05:05 +0000 (12:05 +0200)]
Enable use to the rebar feature to upload buffers to the device. (llama/9251)

8 months agomtgpu: enable VMM (llama/9597)
R0CKSTAR [Thu, 26 Sep 2024 01:27:40 +0000 (09:27 +0800)]
mtgpu: enable VMM (llama/9597)

Signed-off-by: Xiaodong Ye <redacted>
8 months agoggml : remove assert for AArch64 GEMV and GEMM Q4 kernels (llama/9217)
Charles Xu [Wed, 25 Sep 2024 13:12:20 +0000 (15:12 +0200)]
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

8 months agocann: fix crash when llama-bench is running on multiple cann devices (llama/9627)
Dou Xinpeng [Wed, 25 Sep 2024 03:30:38 +0000 (11:30 +0800)]
cann: fix crash when llama-bench is running on multiple cann devices (llama/9627)

8 months agoCUDA: remove bad assert (ggml/972)
Johannes Gäßler [Sun, 29 Sep 2024 17:56:17 +0000 (19:56 +0200)]
CUDA: remove bad assert (ggml/972)

8 months agovulkan : multithread pipeline creation (ggml/963)
Jeff Bolz [Sun, 29 Sep 2024 16:50:17 +0000 (11:50 -0500)]
vulkan : multithread pipeline creation (ggml/963)

8 months agovulkan : fix build for GGML_VULKAN_RUN_TESTS, add TFLOPS to log (ggml/961)
Jeff Bolz [Fri, 27 Sep 2024 07:58:01 +0000 (02:58 -0500)]
vulkan : fix build for GGML_VULKAN_RUN_TESTS, add TFLOPS to log (ggml/961)

8 months agovulkan : argsort barriers must be under uniform control flow (ggml/951)
Salvatore Mesoraca [Thu, 26 Sep 2024 06:59:42 +0000 (08:59 +0200)]
vulkan : argsort barriers must be under uniform control flow (ggml/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>
8 months agoggml : fix GGML_MAX_N_THREADS + improve formatting (ggml/969)
Georgi Gerganov [Tue, 24 Sep 2024 10:23:59 +0000 (13:23 +0300)]
ggml : fix GGML_MAX_N_THREADS + improve formatting (ggml/969)

8 months agoserver : ffmpeg overwrite leftover temp file (#2431)
gilbertgong [Wed, 2 Oct 2024 12:06:40 +0000 (05:06 -0700)]
server : ffmpeg overwrite leftover temp file (#2431)

* Remove possible leftover ffmpeg temp file from a previous failed conversion

* Revert "Remove possible leftover ffmpeg temp file from a previous failed conversion"

This reverts commit 00797403bd43ebcb1e0678989a4fc676d417b4af.

* Flag to force ffmpeg to overwrite output file if it exists

8 months agowhisper : add large-v3-turbo (#2440)
Georgi Gerganov [Tue, 1 Oct 2024 12:57:06 +0000 (15:57 +0300)]
whisper : add large-v3-turbo (#2440)

9 months agotests : remove test-backend-ops (#2434)
Georgi Gerganov [Fri, 27 Sep 2024 08:48:33 +0000 (11:48 +0300)]
tests : remove test-backend-ops (#2434)

9 months agoci : disable failing CUDA and Java builds
Georgi Gerganov [Wed, 25 Sep 2024 07:03:34 +0000 (10:03 +0300)]
ci : disable failing CUDA and Java builds

9 months agoreadme : fix references to download-ggml-model.sh (#2427)
Hugo [Tue, 24 Sep 2024 18:07:51 +0000 (20:07 +0200)]
readme : fix references to download-ggml-model.sh (#2427)

The script itself has a hashbang indicating that it is a shell script,
but the README indicates that it must be executed with `bash`.

I checked the script itself, and it seems to be valid POSIX shell. I can
confirm that it works with busybox sh.

Clarify the reference on the README, so it is clear that bash is not
actually a dependency for this script.

9 months agomake : remove "talk" target until updated
Georgi Gerganov [Tue, 24 Sep 2024 11:15:09 +0000 (14:15 +0300)]
make : remove "talk" target until updated

9 months agoggml : add ggml-cpu-impl.h (skip) (#0)
Georgi Gerganov [Tue, 24 Sep 2024 10:27:33 +0000 (13:27 +0300)]
ggml : add ggml-cpu-impl.h (skip) (#0)

9 months agosync : ggml
Georgi Gerganov [Tue, 24 Sep 2024 10:23:04 +0000 (13:23 +0300)]
sync : ggml

9 months agotalk-llama : sync llama.cpp
Georgi Gerganov [Tue, 24 Sep 2024 10:22:55 +0000 (13:22 +0300)]
talk-llama : sync llama.cpp

9 months agoggml : add AVX512DQ requirement for AVX512 builds (llama/9622)
Eric Zhang [Tue, 24 Sep 2024 08:03:21 +0000 (16:03 +0800)]
ggml : add AVX512DQ requirement for AVX512 builds (llama/9622)

9 months agolog : add CONT level for continuing previous log entry (llama/9610)
Georgi Gerganov [Tue, 24 Sep 2024 07:15:35 +0000 (10:15 +0300)]
log : add CONT level for continuing previous log entry (llama/9610)

9 months agothreads: fix msvc build without openmp (llama/9615)
Max Krasnyansky [Tue, 24 Sep 2024 04:18:48 +0000 (21:18 -0700)]
threads: fix msvc build without openmp (llama/9615)

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

9 months agocuda: add q8_0->f32 cpy operation (llama/9571)
Ivan [Tue, 24 Sep 2024 00:14:24 +0000 (03:14 +0300)]
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.

9 months agothreads: improve ggml_barrier scaling with large number of threads (llama/9598)
Max Krasnyansky [Mon, 23 Sep 2024 18:42:43 +0000 (11:42 -0700)]
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.

9 months agoggml : AVX512 gemm for Q4_0_8_8 (llama/9532)
Srihari-mcw [Mon, 23 Sep 2024 14:06:38 +0000 (19:36 +0530)]
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>
9 months agometal : use F32 prec for K*Q in vec FA (llama/9595)
Georgi Gerganov [Mon, 23 Sep 2024 08:27:47 +0000 (11:27 +0300)]
metal : use F32 prec for K*Q in vec FA (llama/9595)

ggml-ci

9 months agoRevert "[SYCL] fallback mmvq (ggml/9088)" (llama/9579)
Akarshan Biswas [Mon, 23 Sep 2024 03:28:06 +0000 (08:58 +0530)]
Revert "[SYCL] fallback mmvq (ggml/9088)" (llama/9579)

This reverts commit 50addec9a532a6518146ab837a85504850627316.

9 months agomusa: enable building fat binaries, enable unified memory, and disable Flash Attentio...
R0CKSTAR [Sun, 22 Sep 2024 14:55:49 +0000 (22:55 +0800)]
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>
9 months agoFix merge error in #9454 (llama/9589)
Molly Sophia [Sun, 22 Sep 2024 13:26:50 +0000 (21:26 +0800)]
Fix merge error in #9454 (llama/9589)

Signed-off-by: Molly Sophia <redacted>
9 months agoCUDA: enable Gemma FA for HIP/Pascal (llama/9581)
Johannes Gäßler [Sun, 22 Sep 2024 07:34:52 +0000 (09:34 +0200)]
CUDA: enable Gemma FA for HIP/Pascal (llama/9581)