]> git.djapps.eu Git - pkg/ggml/sources/whisper.cpp/log
pkg/ggml/sources/whisper.cpp
6 months agoggml : remove redundant copyright notice + update authors
Georgi Gerganov [Thu, 28 Nov 2024 18:46:40 +0000 (20:46 +0200)]
ggml : remove redundant copyright notice + update authors

6 months agoggml : fix row condition for i8mm kernels (llama/10561)
Georgi Gerganov [Thu, 28 Nov 2024 12:56:37 +0000 (14:56 +0200)]
ggml : fix row condition for i8mm kernels (llama/10561)

ggml-ci

6 months agocmake : fix ARM feature detection (llama/10543)
Georgi Gerganov [Thu, 28 Nov 2024 12:56:23 +0000 (14:56 +0200)]
cmake : fix ARM feature detection (llama/10543)

ggml-ci

6 months agoggml-cpu: support IQ4_NL_4_4 by runtime repack (llama/10541)
Shupei Fan [Thu, 28 Nov 2024 12:52:03 +0000 (20:52 +0800)]
ggml-cpu: support IQ4_NL_4_4 by runtime repack (llama/10541)

* ggml-cpu: support IQ4_NL_4_4 by runtime repack

* ggml-cpu: add __ARM_FEATURE_DOTPROD guard

6 months agokompute : improve backend to pass test_backend_ops (llama/10542)
Sergio López [Thu, 28 Nov 2024 11:51:38 +0000 (12:51 +0100)]
kompute : improve backend to pass test_backend_ops (llama/10542)

* kompute: op_unary: reject unsupported parameters

Signed-off-by: Sergio Lopez <redacted>
* kompute: softmax: implement ALiBi support

Signed-off-by: Sergio Lopez <redacted>
* kompute: rope: implement neox and phi3 support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_q4_k permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_[q4_0|q4_1|q8_0] permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_f16 permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_q6_k permutted support

Signed-off-by: Sergio Lopez <redacted>
---------

Signed-off-by: Sergio Lopez <redacted>
6 months agoCANN: Fix SOC_TYPE compile bug (llama/10519)
leo-pony [Thu, 28 Nov 2024 07:25:24 +0000 (15:25 +0800)]
CANN: Fix SOC_TYPE compile bug (llama/10519)

* CANN: Fix the bug build fail on Ascend310P under two cases:
1) Manual specify SOC_TYPE
2) Under some unusual compile environment

* Update the cann backend News content: Support F16 and F32 data type model for Ascend 310P NPU.

* fix CANN  compile fail bug: the assert in ascend kernel function doesn't supportted on some CANN version

6 months agoCANN: ROPE operator optimization (llama/10540)
Chenguang Li [Thu, 28 Nov 2024 06:24:46 +0000 (14:24 +0800)]
CANN: ROPE operator optimization (llama/10540)

* [cann] ROPE operator optimization

Co-authored-by: noemotiovon <redacted>
6 months agoAdd some minimal optimizations for CDNA (llama/10498)
uvos [Wed, 27 Nov 2024 16:10:08 +0000 (17:10 +0100)]
Add some minimal optimizations for CDNA (llama/10498)

* Add some minimal optimizations for CDNA

* ggml_cuda: set launch bounds also for GCN as it helps there too

6 months agometal : fix group_norm support condition (llama/0)
Georgi Gerganov [Wed, 27 Nov 2024 09:22:14 +0000 (11:22 +0200)]
metal : fix group_norm support condition (llama/0)

6 months agovulkan: define all quant data structures in types.comp (llama/10440)
Jeff Bolz [Wed, 27 Nov 2024 07:32:54 +0000 (01:32 -0600)]
vulkan: define all quant data structures in types.comp (llama/10440)

6 months agovulkan: Handle GPUs with less shared memory (llama/10468)
Jeff Bolz [Wed, 27 Nov 2024 07:30:27 +0000 (01:30 -0600)]
vulkan: Handle GPUs with less shared memory (llama/10468)

There have been reports of failure to compile on systems with <= 32KB
of shared memory (e.g. #10037). This change makes the large tile size
fall back to a smaller size if necessary, and makes mul_mat_id fall
back to CPU if there's only 16KB of shared memory.

6 months agovulkan: further optimize q5_k mul_mat_vec (llama/10479)
Jeff Bolz [Wed, 27 Nov 2024 07:21:59 +0000 (01:21 -0600)]
vulkan: further optimize q5_k mul_mat_vec (llama/10479)

6 months agovulkan: skip integer div/mod in get_offsets for batch_idx==0 (llama/10506)
Jeff Bolz [Wed, 27 Nov 2024 07:08:54 +0000 (01:08 -0600)]
vulkan: skip integer div/mod in get_offsets for batch_idx==0 (llama/10506)

6 months agovulkan: optimize Q2_K and Q3_K mul_mat_vec (llama/10459)
Jeff Bolz [Wed, 27 Nov 2024 07:00:50 +0000 (01:00 -0600)]
vulkan: optimize Q2_K and Q3_K mul_mat_vec (llama/10459)

6 months agomtgpu: Add MUSA_DOCKER_ARCH in Dockerfiles && update cmake and make (llama/10516)
R0CKSTAR [Tue, 26 Nov 2024 16:00:41 +0000 (00:00 +0800)]
mtgpu: Add MUSA_DOCKER_ARCH in Dockerfiles && update cmake and make (llama/10516)

Signed-off-by: Xiaodong Ye <redacted>
6 months agovulkan: fix group_norm (llama/10496)
Jeff Bolz [Tue, 26 Nov 2024 15:45:05 +0000 (09:45 -0600)]
vulkan: fix group_norm (llama/10496)

Fix bad calculation of the end of the range. Add a backend test that
covers the bad case (taken from stable diffusion).

Fixes https://github.com/leejet/stable-diffusion.cpp/issues/439.

6 months agocmake : enable warnings in llama (llama/10474)
Georgi Gerganov [Tue, 26 Nov 2024 12:18:08 +0000 (14:18 +0200)]
cmake : enable warnings in llama (llama/10474)

* cmake : enable warnings in llama

ggml-ci

* cmake : add llama_get_flags and respect LLAMA_FATAL_WARNINGS

* cmake : get_flags -> ggml_get_flags

* speculative-simple : fix warnings

* cmake : reuse ggml_get_flags

ggml-ci

* speculative-simple : fix compile warning

ggml-ci

6 months agoggml-cpu: cmake add arm64 cpu feature check for macos (llama/10487)
Charles Xu [Tue, 26 Nov 2024 11:37:05 +0000 (12:37 +0100)]
ggml-cpu: cmake add arm64 cpu feature check for macos (llama/10487)

* ggml-cpu: cmake add arm64 cpu feature check for macos

* use vmmlaq_s32 for compile option i8mm check

6 months agoCANN: Improve the Inferencing Performance for Ascend NPU Device (llama/10454)
Shanshan Shen [Tue, 26 Nov 2024 10:08:37 +0000 (18:08 +0800)]
CANN: Improve the Inferencing Performance for Ascend NPU Device (llama/10454)

* improve inferencing performance for ascend npu.

Co-authored-by: Frank Mai <redacted>
* some modification after review

* some modifications after review

* restore some modifications

* restore some modifications

---------

Co-authored-by: shanshan shen <redacted>
Co-authored-by: Frank Mai <redacted>
6 months agoCANN: RoPE and CANCAT operator optimization (llama/10488)
Chenguang Li [Tue, 26 Nov 2024 09:31:05 +0000 (17:31 +0800)]
CANN: RoPE and CANCAT operator optimization (llama/10488)

Co-authored-by: noemotiovon <redacted>
6 months agovulkan: Fix a vulkan-shaders-gen arugment parsing error (llama/10484)
Junil Kim [Tue, 26 Nov 2024 01:47:20 +0000 (10:47 +0900)]
vulkan: Fix a vulkan-shaders-gen arugment parsing error (llama/10484)

The vulkan-shaders-gen was not parsing the --no-clean argument correctly.
Because the previous code was parsing the arguments which have a value only
and the --no-clean argument does not have a value, it was not being parsed
correctly. This commit can now correctly parse arguments that don't have values.

6 months agometal : enable mat-vec kernels for bs <= 4 (llama/10491)
Georgi Gerganov [Mon, 25 Nov 2024 19:49:31 +0000 (21:49 +0200)]
metal : enable mat-vec kernels for bs <= 4 (llama/10491)

6 months agollama : accept a list of devices to use to offload a model (llama/10497)
Diego Devesa [Mon, 25 Nov 2024 18:30:06 +0000 (19:30 +0100)]
llama : accept a list of devices to use to offload a model (llama/10497)

* llama : accept a list of devices to use to offload a model

* accept `--dev none` to completely disable offloading

* fix dev list with dl backends

* rename env parameter to LLAMA_ARG_DEVICE for consistency

6 months agoggml : add support for dynamic loading of backends (llama/10469)
Diego Devesa [Mon, 25 Nov 2024 14:13:39 +0000 (15:13 +0100)]
ggml : add support for dynamic loading of backends (llama/10469)

* ggml : add support for dynamic loading of backends

---------

Co-authored-by: Georgi Gerganov <redacted>
6 months agometal : minor code formatting
Georgi Gerganov [Mon, 25 Nov 2024 13:08:04 +0000 (15:08 +0200)]
metal : minor code formatting

6 months agoggml : do not use ARM features not included in the build (llama/10457)
Diego Devesa [Sat, 23 Nov 2024 13:41:12 +0000 (14:41 +0100)]
ggml : do not use ARM features not included in the build (llama/10457)

6 months agoCANN: Support Ascend310P to accelerate F32 and F16 Model (llama/10216)
leo-pony [Fri, 22 Nov 2024 06:07:20 +0000 (14:07 +0800)]
CANN: Support Ascend310P to accelerate F32 and F16 Model (llama/10216)

* CANN Support Ascend310P to accelerate F32 and F16 Model

* Add compile option soc type macro ASCEND_310P to ggml-cann lib

* Remove unused code

* Remove the ascend soc_type hard code compile option in CMakelist.txt

6 months agocuda : optimize argmax (llama/10441)
Diego Devesa [Thu, 21 Nov 2024 17:18:50 +0000 (18:18 +0100)]
cuda : optimize argmax (llama/10441)

* cuda : optimize argmax

* remove unused parameter

ggml-ci

* fixup : use full warps

ggml-ci

* Apply suggestions from code review

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

* ggml : check ne00 <= INT32_MAX in argmax and argsort

---------

Co-authored-by: Johannes Gäßler <redacted>
6 months agovulkan: predicate max operation in soft_max shaders/soft_max (llama/10437)
Jeff Bolz [Wed, 20 Nov 2024 19:47:36 +0000 (13:47 -0600)]
vulkan: predicate max operation in soft_max shaders/soft_max (llama/10437)

Fixes #10434

6 months agovulkan: copy iq4_nl LUT into shared memory (llama/10409)
Jeff Bolz [Wed, 20 Nov 2024 07:40:18 +0000 (01:40 -0600)]
vulkan: copy iq4_nl LUT into shared memory (llama/10409)

6 months agovulkan: further optimize mul_mat_vec using larger loads (llama/10387)
Jeff Bolz [Wed, 20 Nov 2024 07:11:00 +0000 (01:11 -0600)]
vulkan: further optimize mul_mat_vec using larger loads (llama/10387)

* vulkan: Use pipeline_robustness to disable robustness in mul_mat_vec.

Add some early returns for nonexistent rows in mul_mat_vec shaders. These
can only be hit when dispatching a 2D grid of workgroups. Fix the logic
for the 2D grid of workgroups to round up.

Enable the pipeline robustness extension if it's available, and use it to
disable robustness for these pipelines. The instructions to do the bounds
checking contend for the same ALU resources as the bit twiddling dequant
instructions.

* vulkan: Add GLSL structure aliases for quant types to allow larger loads

In Vulkan it's not possible to cast pointer types, so instead you have to
declare an aliased binding for the memory with a different type. This
commit adds aliases for the quant formats using 16b ints, and in a few
places where the struct size is a multiple of 4 also using 32b ints.
Currently only q4_k's aliases are used, but others will be used in
subsequent commits.

* vulkan: use larger loads in q5_k and q6_k shaders.

Similar to the optimization I did in q4_k recently, this vectorizes some loads
and reduces the number of bit twiddling instructions.

* vulkan: use larger K step per iteration in mul_mat_vec.

Add vec4 dequantization functions, and use them to do K=8 per iteration in
mul_mat_vec. This uses 16b loads for the quant values and 128b loads for B
which helps reduce the load on the memory system.

The K_PER_ITER==2 logic is still there, just for F16/F32, and really only
because they support unaligned sizes.

Tweak the num_iters/unrolling logic to be simpler and catch a couple missed
unrolling opportunities.

6 months agoadd cmake rvv support (llama/10411)
haopeng [Tue, 19 Nov 2024 20:10:31 +0000 (04:10 +0800)]
add cmake rvv support (llama/10411)

6 months agoCUDA: remove unnecessary warp reduce in FA (ggml/1032)
mahorozte [Tue, 3 Dec 2024 13:11:43 +0000 (21:11 +0800)]
CUDA: remove unnecessary warp reduce in FA (ggml/1032)

* kqmax_new_j in every thread within warp is same after operate at line 199,this reduce can be omit

* same problem in vec32

---------

Co-authored-by: ZhaoXiaoYu <redacted>
6 months agofeat: add `GGML_UNARY_OP_ARGMAX` Metal kernel (ggml/1019)
PAB [Mon, 2 Dec 2024 18:27:24 +0000 (19:27 +0100)]
feat: add `GGML_UNARY_OP_ARGMAX` Metal kernel (ggml/1019)

* implemented argmax kernel

* tpig -> tgpig

* change to strides

* contiguous assertions

* kernel working and tested

* argmax simd parallel implementation

* added 2 new tests for argmax in test-backend-ops

* cosmit

* added 3 tests cases for perf eval

* add test_argmax in make_test_cases_perf

* Update test-backend-ops.cpp

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

Co-authored-by: Diego Devesa <redacted>
6 months agometal : add `GGML_OP_CONV_TRANSPOSE_1D` kernels (ggml/1026)
PAB [Thu, 28 Nov 2024 08:25:06 +0000 (09:25 +0100)]
metal : add `GGML_OP_CONV_TRANSPOSE_1D` kernels (ggml/1026)

* wip

* wip implementation f32

* kernel conv transpose 1d f32 working

* initial commit

6 months agoDo not include arm_neon.h when compiling CUDA code (ggml/1028)
Frankie Robertson [Tue, 26 Nov 2024 13:50:26 +0000 (15:50 +0200)]
Do not include arm_neon.h when compiling CUDA code (ggml/1028)

6 months agoggml-opt: fix data corruption (ggml/1022)
Johannes Gäßler [Wed, 20 Nov 2024 13:56:04 +0000 (14:56 +0100)]
ggml-opt: fix data corruption (ggml/1022)

7 months agoruby : Add low-level methods to transcribe (#2585)
KITAITI Makoto [Thu, 28 Nov 2024 08:33:07 +0000 (17:33 +0900)]
ruby : Add low-level methods to transcribe (#2585)

* Add tests for Whisper::Context#full

* Add Whisper::Context#full

* Add tests for Whisper::Error

* Add document of Whisper::Context#full [skip ci]

* Add additional signature for Whisper::Context#full

* Add description to Whisper::Context#full

* Add test for Whisper::Context#full_parallel

* Add Whisper::Context#full_parallel

* Hide Whisper's instance methods from Ruby code

* Add class to test MemoryView

* Build test class before running test

* Add test for MemoryView

* Make Whisper::Context#full and #full_parallel accept MemoryView

* Use Ruby 3.1 on CI

* Add comment on samples data type

* Update README

* Update README

* Remove unused code

7 months agomodels : add `q8_0` models to `download-ggml-model.sh` (#2589)
Michael Rienstra [Thu, 28 Nov 2024 08:31:54 +0000 (00:31 -0800)]
models : add `q8_0` models to `download-ggml-model.sh` (#2589)

7 months agoruby : Follow source tree change (#2580)
KITAITI Makoto [Thu, 21 Nov 2024 15:04:29 +0000 (00:04 +0900)]
ruby : Follow source tree change (#2580)

* Follow whisper.cpp source tree change

* Update whispercpp.gemspec

* Follow whisper.cpp log level change

* Fix paths in GitHub workflow for Ruby bindings

* Use GitHub workflow setting for dependency definition

* Use ternary operator

7 months agowhisper : use backend registry (#0)
Georgi Gerganov [Wed, 20 Nov 2024 13:32:34 +0000 (15:32 +0200)]
whisper : use backend registry (#0)

7 months agoggml/sched : do not skip views in pre-assignments
slaren [Wed, 20 Nov 2024 12:25:08 +0000 (13:25 +0100)]
ggml/sched : do not skip views in pre-assignments

7 months agowhisper : adapt to new ggml (wip)
Georgi Gerganov [Tue, 19 Nov 2024 17:09:07 +0000 (19:09 +0200)]
whisper : adapt to new ggml (wip)

7 months agotalk-llama : sync llama.cpp
Georgi Gerganov [Tue, 19 Nov 2024 17:08:57 +0000 (19:08 +0200)]
talk-llama : sync llama.cpp

7 months agosync : ggml
Georgi Gerganov [Tue, 19 Nov 2024 17:04:21 +0000 (19:04 +0200)]
sync : ggml

7 months agoggml : sync resolve (skip) (#0)
Georgi Gerganov [Tue, 19 Nov 2024 17:03:47 +0000 (19:03 +0200)]
ggml : sync resolve (skip) (#0)

7 months agoAdd required ggml-base and backend libs to cmake pkg (llama/10407)
bandoti [Tue, 19 Nov 2024 16:10:30 +0000 (12:10 -0400)]
Add required ggml-base and backend libs to cmake pkg (llama/10407)

7 months agocuda : fix CUDA_FLAGS not being applied (llama/10403)
Diego Devesa [Tue, 19 Nov 2024 13:29:38 +0000 (14:29 +0100)]
cuda : fix CUDA_FLAGS not being applied (llama/10403)

7 months agosycl : Add option to set the SYCL architecture for all targets (llama/10266)
Romain Biessy [Tue, 19 Nov 2024 08:02:23 +0000 (09:02 +0100)]
sycl : Add option to set the SYCL architecture for all targets (llama/10266)

* Add option to set the SYCL architecture for all targets
* Convert GGML_SYCL_HIP_TARGET to the more generic GGML_SYCL_ARCH option
* Document that setting GGML_SYCL_ARCH can improve the performance

7 months agovulkan: Optimize soft_max (llama/10301)
Jeff Bolz [Tue, 19 Nov 2024 07:25:17 +0000 (01:25 -0600)]
vulkan: Optimize soft_max (llama/10301)

* vulkan: Optimize soft_max

Large soft_max could already saturate memory, but small/medium sizes were
pretty slow. The bulk of the gains for them comes from using a smaller
workgroup size, and making the workgroup size match the subgroup size also
makes the barriers much cheaper.

Cache some values in locals to avoid refetching/recomputing. And stamp
out a few "template instantiations" so smaller cases will fully unroll.

Add a missing early return for OOB rows. This happens when there are more
than 512 rows and the dispatch is 512 x H.

* vulkan: Further soft_max optimizations

Restore the workgroup size of 512 case, use it for >1024.

Use unrollable loops for more iteration counts.

7 months agosycl: Revert MUL_MAT_OP support changes (llama/10385)
Alberto Cabrera Pérez [Tue, 19 Nov 2024 00:50:04 +0000 (00:50 +0000)]
sycl: Revert MUL_MAT_OP support changes (llama/10385)

7 months agocuda : only use native when supported by cmake (llama/10389)
Diego Devesa [Mon, 18 Nov 2024 17:43:40 +0000 (18:43 +0100)]
cuda : only use native when supported by cmake (llama/10389)

7 months agovulkan: remove use of null initializer (llama/10372)
Jeff Bolz [Mon, 18 Nov 2024 14:28:42 +0000 (08:28 -0600)]
vulkan: remove use of null initializer (llama/10372)

Seems like this isn't working for vulkan-over-metal when the array is sized
by a spec constant. Maybe a spirv-cross limitation?

7 months agometal : fox offset integer overflows in im2col (ggml/1015)
Plamen Minev [Mon, 18 Nov 2024 13:02:27 +0000 (15:02 +0200)]
metal : fox offset integer overflows in im2col (ggml/1015)

-- While running StableDiffusion.cpp locally with Metal some offsets overflow and results in incorrect calculations

7 months agoVulkan: Fix device info output format specifiers (llama/10366)
0cc4m [Mon, 18 Nov 2024 10:02:43 +0000 (11:02 +0100)]
Vulkan: Fix device info output format specifiers (llama/10366)

* Vulkan: Fix device info output format specifiers

* Vulkan: Use zu printf specifier for size_t instead of ld

7 months agometal : add `GGML_UNARY_OP_ELU` kernel (ggml/1018)
PAB [Mon, 18 Nov 2024 09:02:49 +0000 (10:02 +0100)]
metal : add `GGML_UNARY_OP_ELU` kernel (ggml/1018)

7 months agoCUDA: fix MMV kernel being used for FP16 src1 (llama/10357)
Johannes Gäßler [Sun, 17 Nov 2024 22:20:42 +0000 (23:20 +0100)]
CUDA: fix MMV kernel being used for FP16 src1 (llama/10357)

7 months agoCMake: fix typo in comment [no ci] (llama/10360)
Johannes Gäßler [Sun, 17 Nov 2024 11:59:38 +0000 (12:59 +0100)]
CMake: fix typo in comment [no ci] (llama/10360)

7 months agollama : only use default buffer types for the KV cache (llama/10358)
Diego Devesa [Sun, 17 Nov 2024 11:25:45 +0000 (12:25 +0100)]
llama : only use default buffer types for the KV cache (llama/10358)

7 months agometal : refactor kernel args into structs (llama/10238)
Georgi Gerganov [Sun, 17 Nov 2024 09:23:01 +0000 (11:23 +0200)]
metal : refactor kernel args into structs (llama/10238)

* metal : add kernel arg structs (wip)

* metal : fattn args

ggml-ci

* metal : cont + avoid potential int overflow [no ci]

* metal : mul mat struct (wip)

* cont : mul mat vec

* cont : pass by reference

* cont : args is first argument

* cont : use char ptr

* cont : shmem style

* cont : thread counters style

* cont : mul mm id

ggml-ci

* cont : int safety + register optimizations

ggml-ci

* metal : GGML_OP_CONCAT

ggml-ci

* metal : GGML_OP_ADD, GGML_OP_SUB, GGML_OP_MUL, GGML_OP_DIV

* metal : GGML_OP_REPEAT

* metal : GGML_OP_CPY

* metal : GGML_OP_RMS_NORM

* metal : GGML_OP_NORM

* metal : add TODOs for rest of ops

* ggml : add ggml-metal-impl.h

ggml-ci

7 months agoggml : fix undefined reference to 'getcpu' (llama/10354)
FirstTimeEZ [Sun, 17 Nov 2024 08:39:22 +0000 (21:39 +1300)]
ggml : fix undefined reference to 'getcpu' (llama/10354)

https://github.com/ggerganov/llama.cpp/issues/10352

7 months agoCUDA: remove DMMV, consolidate F16 mult mat vec (llama/10318)
Johannes Gäßler [Sun, 17 Nov 2024 08:09:55 +0000 (09:09 +0100)]
CUDA: remove DMMV, consolidate F16 mult mat vec (llama/10318)

7 months agoCMake: default to -arch=native for CUDA build (llama/10320)
Johannes Gäßler [Sun, 17 Nov 2024 08:06:34 +0000 (09:06 +0100)]
CMake: default to -arch=native for CUDA build (llama/10320)

7 months agoggml : fix possible buffer use after free in sched reserve (llama/9930)
Diego Devesa [Sun, 17 Nov 2024 06:31:17 +0000 (07:31 +0100)]
ggml : fix possible buffer use after free in sched reserve (llama/9930)

7 months agoggml : inttypes.h -> cinttypes (llama/0)
Georgi Gerganov [Sat, 16 Nov 2024 21:40:39 +0000 (23:40 +0200)]
ggml : inttypes.h -> cinttypes (llama/0)

ggml-ci

7 months agoggml : adapt AMX to tensor->grad removal (llama/0)
Georgi Gerganov [Sat, 16 Nov 2024 19:38:01 +0000 (21:38 +0200)]
ggml : adapt AMX to tensor->grad removal (llama/0)

ggml-ci

7 months agoggml : fix compile warnings (llama/0)
Georgi Gerganov [Sat, 16 Nov 2024 19:32:41 +0000 (21:32 +0200)]
ggml : fix compile warnings (llama/0)

ggml-ci

7 months agollamafile : fix include path (llama/0)
Georgi Gerganov [Sat, 16 Nov 2024 15:58:56 +0000 (17:58 +0200)]
llamafile : fix include path (llama/0)

ggml-ci

7 months agovulkan: Optimize some mat-vec mul quant shaders (llama/10296)
Jeff Bolz [Sat, 16 Nov 2024 06:26:57 +0000 (00:26 -0600)]
vulkan: Optimize some mat-vec mul quant shaders (llama/10296)

Compute two result elements per workgroup (for Q{4,5}_{0,1}). This reuses
the B loads across the rows and also reuses some addressing calculations.
This required manually partially unrolling the loop, since the compiler
is less willing to unroll outer loops.

Add bounds-checking on the last iteration of the loop. I think this was at
least partly broken before.

Optimize the Q4_K shader to vectorize most loads and reduce the number of
bit twiddling instructions.

7 months agoggml : optimize Q4_0 into Q4_0_X_Y repack (llama/10324)
Dan Johansson [Sat, 16 Nov 2024 00:53:37 +0000 (01:53 +0100)]
ggml : optimize Q4_0 into Q4_0_X_Y repack (llama/10324)

7 months agoMake updates to fix issues with clang-cl builds while using AVX512 flags (llama/10314)
Srihari-mcw [Fri, 15 Nov 2024 21:27:00 +0000 (02:57 +0530)]
Make updates to fix issues with clang-cl builds while using AVX512 flags (llama/10314)

7 months agoggml: new optimization interface (ggml/988)
Johannes Gäßler [Sat, 16 Nov 2024 12:49:35 +0000 (13:49 +0100)]
ggml: new optimization interface (ggml/988)

* ggml: new optimization interface

remove test2.c, test3.c

store adamw params in tensor

move grads from tensor to graph

* avoid segfault upon API misuse

* add ggml-opt.h to public headers

* remove dependence of ggml-opt.cpp on ggml-cpu.h

7 months agoggml : remove duplicated sources from the last sync (ggml/1017)
Georgi Gerganov [Fri, 15 Nov 2024 21:52:31 +0000 (23:52 +0200)]
ggml : remove duplicated sources from the last sync (ggml/1017)

* ggml : remove duplicated sources from the last sync

ggml-ci

* cont : remove FindSIMD.cmake [no ci]

7 months agoggml : fix some build issues
slaren [Fri, 15 Nov 2024 19:20:54 +0000 (20:20 +0100)]
ggml : fix some build issues

7 months agosync : leftovers (ggml/0)
Georgi Gerganov [Fri, 15 Nov 2024 19:43:41 +0000 (21:43 +0200)]
sync : leftovers (ggml/0)

ggml-ci

7 months agocmake : restore CMakeLists.txt (llama/10256)
Georgi Gerganov [Fri, 15 Nov 2024 19:35:51 +0000 (21:35 +0200)]
cmake : restore CMakeLists.txt (llama/10256)

ggml-ci

7 months agoAVX BF16 and single scale quant optimizations (llama/10212)
Eve [Fri, 15 Nov 2024 11:47:58 +0000 (11:47 +0000)]
AVX BF16 and single scale quant optimizations (llama/10212)

* use 128 bit loads (i've tried 256->128 to death and its slower)

* double accumulator

* avx bf16 vec dot

* +3% q4_0 inference

* +7% tg +5% pp compared to master

* slower f16c version, kep for reference

* 256b version, also slow. i tried :)

* revert f16

* faster with madd

* split to functions

* Q8_0 and IQ4_NL, 5-7% faster

* fix potential overflow (performance reduced)

* 16 bit add for q4_0 only

* merge

7 months agosycl: Use syclcompat::dp4a (llama/10267)
Romain Biessy [Fri, 15 Nov 2024 03:09:12 +0000 (04:09 +0100)]
sycl: Use syclcompat::dp4a (llama/10267)

* sycl: Use syclcompat::dp4a

* Using the syclcompat version allow the compiler to optimize the
  operation with native function

* Update news section

* Update CI Windows oneAPI version to 2025.0

* Reword doc

* Call syclcompat::dp4a inside dpct::dp4a

This reverts commit 90cb61d692d61360b46954a1c7f780bd2e569b73.

7 months agobackend cpu: add online flow for aarch64 Q4_0 GEMV/GEMM kernels (llama/9921)
Charles Xu [Fri, 15 Nov 2024 00:28:50 +0000 (01:28 +0100)]
backend cpu: add online flow for aarch64 Q4_0 GEMV/GEMM kernels (llama/9921)

* backend-cpu: add online flow for aarch64 Q4_0 GEMV/GEMM kernels

---------

Co-authored-by: Diego Devesa <redacted>
7 months agoggml : build backends as libraries (llama/10256)
Diego Devesa [Thu, 14 Nov 2024 17:04:35 +0000 (18:04 +0100)]
ggml : build backends as libraries (llama/10256)

* ggml : build backends as libraries

---------

Signed-off-by: Xiaodong Ye <redacted>
Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: R0CKSTAR <redacted>
7 months agoscripts : update sync
Georgi Gerganov [Tue, 19 Nov 2024 16:59:18 +0000 (18:59 +0200)]
scripts : update sync

7 months agorelease : v1.7.2
Georgi Gerganov [Tue, 19 Nov 2024 16:54:22 +0000 (18:54 +0200)]
release : v1.7.2

7 months agosycl: fix example build (#2570)
Stefan Sydow [Mon, 18 Nov 2024 12:57:23 +0000 (13:57 +0100)]
sycl: fix example build (#2570)

7 months agoci : use local ggml in Android build (#2567)
Georgi Gerganov [Sat, 16 Nov 2024 18:45:41 +0000 (20:45 +0200)]
ci : use local ggml in Android build (#2567)

7 months agoggml : tmp workaround for whisper.cpp (skip) (#2565)
Georgi Gerganov [Sat, 16 Nov 2024 18:19:02 +0000 (20:19 +0200)]
ggml : tmp workaround for whisper.cpp (skip) (#2565)

7 months agoupdate : readme
Georgi Gerganov [Fri, 15 Nov 2024 14:00:10 +0000 (16:00 +0200)]
update : readme

7 months agoscripts : fix sync path
Georgi Gerganov [Fri, 15 Nov 2024 13:24:09 +0000 (15:24 +0200)]
scripts : fix sync path

7 months agowhisper.swiftui : switch Mac dest to Mac (Designed for iPad) (#2562)
Jhen-Jie Hong [Fri, 15 Nov 2024 13:21:53 +0000 (21:21 +0800)]
whisper.swiftui : switch Mac dest to Mac (Designed for iPad) (#2562)

7 months agocmake : fix ppc64 check (#0)
Georgi Gerganov [Fri, 15 Nov 2024 07:04:34 +0000 (09:04 +0200)]
cmake : fix ppc64 check (#0)

7 months agowhisper : include ggml-cpu.h (#0)
Georgi Gerganov [Fri, 15 Nov 2024 09:01:47 +0000 (11:01 +0200)]
whisper : include ggml-cpu.h (#0)

7 months agobuild : fixes
Georgi Gerganov [Fri, 15 Nov 2024 07:07:53 +0000 (09:07 +0200)]
build : fixes

7 months agotalk-llama : sync llama.cpp
Georgi Gerganov [Fri, 15 Nov 2024 06:41:06 +0000 (08:41 +0200)]
talk-llama : sync llama.cpp

7 months agowhisper : fix build (#0)
Georgi Gerganov [Fri, 15 Nov 2024 06:40:47 +0000 (08:40 +0200)]
whisper : fix build (#0)

7 months agosync : ggml
Georgi Gerganov [Fri, 15 Nov 2024 06:40:34 +0000 (08:40 +0200)]
sync : ggml

7 months agosycl : Fixes to broken builds and test-backend-ops (llama/10257)
Alberto Cabrera Pérez [Wed, 13 Nov 2024 09:40:57 +0000 (09:40 +0000)]
sycl : Fixes to broken builds and test-backend-ops (llama/10257)

* Fixes broken build for the SYCL CUDA backend caused by non-explicit gemm call in outprod (merged in with RWKV6 in
Optimize RWKV6 Operator Naming and Implement Multi-core CPU/ SYCL Acceleration #10133)

* Marks permuted MUL_MAT as unsupported to be able to run test-backend-ops

* Fixes asserts in norm to fix debug builds.

7 months agovulkan: Optimize contiguous copies (llama/10254)
Jeff Bolz [Wed, 13 Nov 2024 06:58:57 +0000 (00:58 -0600)]
vulkan: Optimize contiguous copies (llama/10254)

* tests: Fix memory bandwidth calculation for perf tests

Add a flops calculation for flash attention.

Add one GGML_OP_CPY perf test.

* vulkan: Optimize contiguous copies

Add a variant of the copy shader for when the tensors are contiguous. Avoid
the complex addressing calculations, and do four elements per invocation
to hide some other overhead.

Apply similar changes to the scale shader, since scale is always contiguous.

Add a "progress bar" for shader compiles.

7 months agovulkan: Throttle the number of shader compiles during the build step. (llama/10222)
Jeff Bolz [Mon, 11 Nov 2024 17:13:51 +0000 (11:13 -0600)]
vulkan: Throttle the number of shader compiles during the build step. (llama/10222)

Fixes #9582

Spawning too many concurrent copies of glslc leads to "Failed to create pipes"
errors on Linux. This change applies the same throttling we use for
multithreaded pipeline creation.

7 months agometal : more precise Q*K in FA vec kernel (llama/10247)
Georgi Gerganov [Mon, 11 Nov 2024 06:39:13 +0000 (08:39 +0200)]
metal : more precise Q*K in FA vec kernel (llama/10247)

7 months agovulkan: Fix newly added tests for permuted mul_mat and 1D im2col (llama/10226)
Jeff Bolz [Sun, 10 Nov 2024 11:37:56 +0000 (05:37 -0600)]
vulkan: Fix newly added tests for permuted mul_mat and 1D im2col (llama/10226)

7 months agometal : reorder write loop in mul mat kernel + style (llama/10231)
Georgi Gerganov [Sat, 9 Nov 2024 09:53:13 +0000 (11:53 +0200)]
metal : reorder write loop in mul mat kernel + style (llama/10231)

* metal : reorder write loop

* metal : int -> short, style

ggml-ci