]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
2 years agoggml : update softmax n_task calculation (llama/5126)
snadampal [Fri, 26 Jan 2024 17:17:59 +0000 (11:17 -0600)]
ggml : update softmax n_task calculation (llama/5126)

updated the n_task calculation to use max number of
threads possible. This has improved the prompt eval
performance by around 5% for DOT kernels and by
around 10% for MMLA kernels on AWS Graviton3.

2 years agometal : remove unused `n_buffers` and `buffers` (llama/5129)
Paul Tsochantaris [Fri, 26 Jan 2024 12:16:07 +0000 (12:16 +0000)]
metal : remove unused `n_buffers` and `buffers` (llama/5129)

2 years agometal : show compile log messages
Georgi Gerganov [Thu, 25 Jan 2024 09:26:17 +0000 (11:26 +0200)]
metal : show compile log messages

2 years agocuda : fix 2-bit quants on amd hip (llama/5105)
Engininja2 [Wed, 24 Jan 2024 22:18:15 +0000 (16:18 -0600)]
cuda : fix 2-bit quants on amd hip (llama/5105)

* cuda : fix 2-bit quants on amd hip

* use __low2float intrinsic function for new quants

2 years agollama : pre-allocate input tensors in a separate buffer (llama/5100)
slaren [Wed, 24 Jan 2024 11:48:14 +0000 (12:48 +0100)]
llama : pre-allocate input tensors in a separate buffer (llama/5100)

2 years agometal : disable support for MUL_MAT F32 x F16
Georgi Gerganov [Tue, 23 Jan 2024 13:50:56 +0000 (15:50 +0200)]
metal : disable support for MUL_MAT F32 x F16

2 years agoCUDA: more info when no device code (llama/5088)
Johannes Gäßler [Tue, 23 Jan 2024 12:31:56 +0000 (13:31 +0100)]
CUDA: more info when no device code (llama/5088)

2 years agominor : clean-up some warnings and style (llama/5094)
Georgi Gerganov [Tue, 23 Jan 2024 12:12:57 +0000 (14:12 +0200)]
minor : clean-up some warnings and style (llama/5094)

* minor : clean-up some warnings and style

ggml-ci

* ggml : add comment

2 years agoggml : parallelize FP32 conversion when using BLAS (llama/5045)
Reinforce-II [Mon, 22 Jan 2024 13:15:08 +0000 (21:15 +0800)]
ggml : parallelize FP32 conversion when using BLAS (llama/5045)

* make GGML_TASK_INIT phase can be run in multithread

* multithreaded dequantize in mul_mat when using blas library

* minor fixes

* update outdated comment
* fix coding style

* simplify code

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

Co-authored-by: Georgi Gerganov <redacted>
2 years agollava : MobileVLM support (llama/4954)
XiaotaoChen [Mon, 22 Jan 2024 13:09:35 +0000 (21:09 +0800)]
llava : MobileVLM support (llama/4954)

* MobileVLM native implementation

* delete depthwise_conv_2d and permute_cpy relative code, replace the two by the existed functions, and opt ldp definition, support LLAMA_PERF option for CMake

* move android script to example/llava directory

* Fix the editor config checks

---------

Co-authored-by: Chenxiaotao03 <redacted>
2 years agollama : run all KQV ops on the CPU with no KV offload (llama/5049)
slaren [Sat, 20 Jan 2024 15:05:49 +0000 (16:05 +0100)]
llama : run all KQV ops on the CPU with no KV offload (llama/5049)

ggml-ci

2 years agocuda : fix compile error in jetson platform (llama/4975)
Kylin [Sat, 20 Jan 2024 07:01:46 +0000 (15:01 +0800)]
cuda : fix compile error in jetson platform (llama/4975)

* cuda: fix compile error in jetson platform

* cuda: update comment in ggml-cuda.cu

* cuda: update ggml-cuda.cu comment

2 years agogpt-2 : clarify instructions for CLBlast on Android (#706)
Neuman Vong [Fri, 26 Jan 2024 13:14:58 +0000 (00:14 +1100)]
gpt-2 : clarify instructions for CLBlast on Android (#706)

* Tabs to spaces

* CLBlast build

* Update README

* Clarify pwd

* Omit output

* How to get libOpenCL.so

* Clarify OpenCL limitations

* Prefer parameters over envvar

* @slaren

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

Co-authored-by: slaren <redacted>
2 years agoggml : check ggml_add src1 type (#708)
Judd [Fri, 26 Jan 2024 13:04:01 +0000 (21:04 +0800)]
ggml : check ggml_add src1 type (#708)

Co-authored-by: Judd <redacted>
2 years agomnist : add tensorflow and keras to requirements.txt (#702)
Jack Vial [Mon, 22 Jan 2024 12:03:48 +0000 (07:03 -0500)]
mnist : add tensorflow and keras to requirements.txt (#702)

* Added tensorflow and keras to requirements.txt to streamline running MNIST example. Updated instructions to include installing python dependenices for the training script.

* add output example whitespace back

* edit readme and fixed typo

2 years agosync : whisper.cpp
Georgi Gerganov [Thu, 18 Jan 2024 09:04:35 +0000 (11:04 +0200)]
sync  : whisper.cpp

2 years agosync : llama.cpp
Georgi Gerganov [Thu, 18 Jan 2024 08:47:52 +0000 (10:47 +0200)]
sync : llama.cpp

2 years agometal : fix memory leak, dangling pointer and unused autorel (llama/5007)
Paul Tsochantaris [Thu, 18 Jan 2024 08:47:24 +0000 (08:47 +0000)]
metal : fix memory leak, dangling pointer and unused autorel (llama/5007)

* Metal memory: Small memory leak on init, dangling pointer, and unused autorelease pool in graph compute

* SPM header potential fix

* Reverting symlinks

2 years agoggml : fix SPM package headers
Georgi Gerganov [Thu, 18 Jan 2024 08:35:42 +0000 (10:35 +0200)]
ggml : fix SPM package headers

2 years agoreadme : add link (#699)
Judd [Wed, 17 Jan 2024 19:43:53 +0000 (03:43 +0800)]
readme : add link (#699)

add a link to ChatLLM.cpp

2 years agosync : llama.cpp
Georgi Gerganov [Wed, 17 Jan 2024 18:53:31 +0000 (20:53 +0200)]
sync : llama.cpp

2 years agometal : update ggml-metal.m from llama.cpp
Georgi Gerganov [Wed, 17 Jan 2024 18:53:14 +0000 (20:53 +0200)]
metal : update ggml-metal.m from llama.cpp

2 years agoggml : add IQ2 to test-backend-ops + refactoring (llama/4990)
Georgi Gerganov [Wed, 17 Jan 2024 16:54:56 +0000 (18:54 +0200)]
ggml : add IQ2 to test-backend-ops + refactoring (llama/4990)

* ggml : add IQ2 to test-backend-ops + refactoring

ggml-ci

* cuda : update supports_op for IQ2

ggml-ci

* ci : enable LLAMA_CUBLAS=1 for CUDA nodes

ggml-ci

* cuda : fix out-of-bounds-access in `mul_mat_vec_q`

ggml-ci

* tests : avoid creating RNGs for each Q tensor

ggml-ci

* tests : avoid creating RNGs for each tensor

ggml-ci

2 years agoimatrix : offload to GPU support (llama/4957)
Georgi Gerganov [Wed, 17 Jan 2024 16:46:30 +0000 (18:46 +0200)]
imatrix : offload to GPU support (llama/4957)

* backend : add eval callback

ggml-ci

* backend : group nodes in a single compute when user don't need them

* backend : clean-up the implementation

ggml-ci

* simple : do not perform tensor data copy if not needed

* simple : fix

* imatrix : offload to GPU support

* imatrix : fix ggml_mul_mat_id hanlding

ggml-ci

* ci : add imatrix test

ggml-ci

* ci : rearrange output

ggml-ci

2 years agobackend : add eval callback (llama/4935)
Georgi Gerganov [Wed, 17 Jan 2024 16:39:41 +0000 (18:39 +0200)]
backend : add eval callback (llama/4935)

* backend : add eval callback

ggml-ci

* backend : group nodes in a single compute when user don't need them

* backend : clean-up the implementation

ggml-ci

* simple : do not perform tensor data copy if not needed

* simple : fix

* simple : no need for ggml_is_contiguous + fix bool parse

* llama : fix callback placement in llama_context_params

* backend : avoid double-ask callback calls

* simple : restore examples, imatrix will serve as a demo

2 years agometal : create autorelease pool during library build (llama/4970)
Georgi Gerganov [Wed, 17 Jan 2024 16:38:39 +0000 (18:38 +0200)]
metal : create autorelease pool during library build (llama/4970)

* metal : create autorelease pool during library build

ggml-ci

* test : simplify

ggml-ci

2 years agoggml : importance matrix support for legacy quants (llama/4969)
Kawrakow [Tue, 16 Jan 2024 17:51:26 +0000 (19:51 +0200)]
ggml : importance matrix support for legacy quants (llama/4969)

* imatrix: adding support for legacy quants

* imatrix: guard Q4_0/Q5_0 against ffn_down craziness

---------

Co-authored-by: Iwan Kawrakow <redacted>
2 years agometal : log `recommendedMaxWorkingSetSize` on iOS 16+ (llama/4936)
Alex Azarov [Tue, 16 Jan 2024 13:33:02 +0000 (14:33 +0100)]
metal : log `recommendedMaxWorkingSetSize` on iOS 16+ (llama/4936)

* metal: Log `recommendedMaxWorkingSetSize` on iOS 16+

* Only log on iOS and macOS, ignoring tvOS and other platforms

* Check for Xcode version before using recommendedMaxWorkingSetSize

---------

Co-authored-by: Georgi Gerganov <redacted>
2 years agoggml : introduce GGML_CALL function annotation (llama/4850)
Justine Tunney [Tue, 16 Jan 2024 11:16:33 +0000 (03:16 -0800)]
ggml : introduce GGML_CALL function annotation (llama/4850)

This change makes it possible to build ggml-cuda.cu and ggml-metal.m as
independent dynamic shared objects, that may be conditionally linked at
runtime in a multiplatform binary. It introduces a GGML_CALL annotation
that documents which functions have a cyclic call relationship, between
the application code and GPU modules.

This change does nothing, unless the build defines -DGGML_MULTIPLATFORM
which causes back-references and function pointers to conform to MS ABI
which is supported by NVCC, ROCm, XCode, GCC and Clang across platforms

2 years agocuda : fix dequantize kernel names (llama/4938)
Georgi Gerganov [Mon, 15 Jan 2024 11:27:00 +0000 (13:27 +0200)]
cuda : fix dequantize kernel names (llama/4938)

2 years agoCUDA: faster dequantize kernels for Q4_0 and Q4_1 (llama/4938)
Kawrakow [Mon, 15 Jan 2024 05:48:06 +0000 (07:48 +0200)]
CUDA: faster dequantize kernels for Q4_0 and Q4_1 (llama/4938)

Co-authored-by: Iwan Kawrakow <redacted>
2 years agoAdd ability to use importance matrix for all k-quants (llama/4930)
Kawrakow [Sun, 14 Jan 2024 14:21:12 +0000 (16:21 +0200)]
Add ability to use importance matrix for all k-quants (llama/4930)

Co-authored-by: Iwan Kawrakow <redacted>
2 years agosync : whisper.cpp
Georgi Gerganov [Sun, 14 Jan 2024 09:07:09 +0000 (11:07 +0200)]
sync : whisper.cpp

2 years agosync : llama.cpp
Georgi Gerganov [Sun, 14 Jan 2024 08:49:50 +0000 (10:49 +0200)]
sync : llama.cpp

2 years agometal : correctly set SIMD support flags on iOS (llama/4923)
Alex Azarov [Sun, 14 Jan 2024 08:44:39 +0000 (09:44 +0100)]
metal : correctly set SIMD support flags on iOS (llama/4923)

* Correctly set support_simdgroup_reduction and support_simdgroup_mm on iPhone/iPad

* log a little bit more info on iOS

2 years ago2-bit quantizations (llama/4897)
Kawrakow [Sun, 14 Jan 2024 07:45:56 +0000 (09:45 +0200)]
2-bit quantizations (llama/4897)

* imatrix: load

* imatrix: WIP

* imatrix: Add Q2_K quantization

* imatrix: also guard against Q2_K_S quantization without importance matrix

* imatrix: guard even more against low-bit quantization misuse

---------

Co-authored-by: Iwan Kawrakow <redacted>
2 years agosync : whisper.cpp
Georgi Gerganov [Sat, 13 Jan 2024 22:14:03 +0000 (00:14 +0200)]
sync : whisper.cpp

2 years agowhisper : load the model into multiple buffers of max size 1GB (whisper/1763)
Georgi Gerganov [Sat, 13 Jan 2024 15:47:40 +0000 (17:47 +0200)]
whisper : load the model into multiple buffers of max size 1GB (whisper/1763)

2 years agosync : llama.cpp
Georgi Gerganov [Sat, 13 Jan 2024 22:09:33 +0000 (00:09 +0200)]
sync : llama.cpp

2 years agoexamples : adapt to metal API
Georgi Gerganov [Sat, 13 Jan 2024 22:09:26 +0000 (00:09 +0200)]
examples : adapt to metal API

2 years agoggml: cache sin/cos for RoPE (llama/4908)
Johannes Gäßler [Sat, 13 Jan 2024 20:41:37 +0000 (21:41 +0100)]
ggml: cache sin/cos for RoPE (llama/4908)

2 years agometal : remove old API (llama/4919)
Georgi Gerganov [Sat, 13 Jan 2024 18:45:45 +0000 (20:45 +0200)]
metal : remove old API (llama/4919)

ggml-ci

2 years agometal : disable log for loaded kernels (llama/4794)
Georgi Gerganov [Sat, 13 Jan 2024 16:46:37 +0000 (18:46 +0200)]
metal : disable log for loaded kernels (llama/4794)

2 years agogguf : fix potential infinite for-loop (llama/4600)
texmex76 [Sat, 13 Jan 2024 16:06:20 +0000 (17:06 +0100)]
gguf : fix potential infinite for-loop (llama/4600)

Co-authored-by: Bernhard Gstrein <redacted>
2 years agometal : refactor kernel loading code (llama/4794)
Georgi Gerganov [Sat, 13 Jan 2024 16:03:45 +0000 (18:03 +0200)]
metal : refactor kernel loading code (llama/4794)

* metal : detect more GPU families

* metal : refactor kernel loading

* metal : set kernel family requirements

* metal : fix kernel init + fix compile options

* metal : take into account simdgroup reduction support

* metal : print only skipped kernels

* metal : fix check for simdgroup reduction support

* metal : check for Metal 3

* metal : free allocations

* metal : normalize encoder:setComputePipelineStatus calls

ggml-ci

* metal : fix Metal3 family check

ggml-ci

* metal : check for simdgroup matrix mul. feature

ggml-ci

2 years agoCUDA: faster q8_0 -> f16 dequantization (llama/4895)
Johannes Gäßler [Fri, 12 Jan 2024 19:38:54 +0000 (20:38 +0100)]
CUDA: faster q8_0 -> f16 dequantization (llama/4895)

2 years agosync : whisper.cpp
Georgi Gerganov [Fri, 12 Jan 2024 20:01:22 +0000 (22:01 +0200)]
sync : whisper.cpp

2 years agoggml : fix 32-bit ARM compat for IQ2_XS (whisper/1758)
Georgi Gerganov [Fri, 12 Jan 2024 12:02:30 +0000 (14:02 +0200)]
ggml : fix 32-bit ARM compat for IQ2_XS (whisper/1758)

* ggml : fix 32-bit ARM compat

* ggml : fix fix

* ggml : fix fix fix

2 years agowhisper : fix segment length with params.no_timestamps == true
Georgi Gerganov [Fri, 12 Jan 2024 11:37:38 +0000 (13:37 +0200)]
whisper : fix segment length with params.no_timestamps == true

2 years agoparams : don't compute timestamps when not printing them (whisper/1755)
George Hindle [Fri, 12 Jan 2024 11:24:38 +0000 (11:24 +0000)]
params : don't compute timestamps when not printing them (whisper/1755)

2 years agobackend_sched : fix assignments
slaren [Fri, 12 Jan 2024 19:38:34 +0000 (20:38 +0100)]
backend_sched : fix assignments

ggml-ci

2 years agoexamples : remove obsolete starcoder mmap example
Georgi Gerganov [Fri, 12 Jan 2024 19:28:02 +0000 (21:28 +0200)]
examples : remove obsolete starcoder mmap example

ggml-ci

2 years agoupdate gpt-2 example
slaren [Fri, 12 Jan 2024 19:23:47 +0000 (20:23 +0100)]
update gpt-2 example

2 years agosync : llama.cpp
Georgi Gerganov [Fri, 12 Jan 2024 19:15:10 +0000 (21:15 +0200)]
sync : llama.cpp

ggml-ci

2 years agollama : ggml-backend integration (llama/4766)
slaren [Fri, 12 Jan 2024 19:07:38 +0000 (20:07 +0100)]
llama : ggml-backend integration (llama/4766)

* llama : ggml-backend integration

* ggml-backend : add names to buffers

* fix unmap after loading

* batched-bench : add tensor_split param

* llama : check for null tensor_split

* ggml-backend : increase GGML_MAX_BACKENDS

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

* cuda : add ggml-backend split buffer support

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

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

* ggml : fix null backend dereference

* ggml : also check ggml_backend_is_cpu

* test-backend-ops : check buffer allocation failures

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

* ggml : fix mul_mat_id work size

* llama : rewrite session kv load/set without graphs

* minor

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

* llama : abort ctx if cuda backend init fails

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

ggml-ci

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

* opencl : add ggml-backend buffer type

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

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

ggml-ci

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

* Apply suggestions from code review

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

* address review comments

* llama-bench : add split-mode parameter

* fix whitespace

* opencl : fix double initialization

* server : add --split-mode parameter

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

ggml-ci

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

* fix opencl

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

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Johannes Gäßler <redacted>
2 years agoCUDA: fix softmax compile for old CUDA versions (llama/4862)
Johannes Gäßler [Fri, 12 Jan 2024 11:30:41 +0000 (12:30 +0100)]
CUDA: fix softmax compile for old CUDA versions (llama/4862)

2 years agoImportance Matrix calculation (llama/4861)
Kawrakow [Fri, 12 Jan 2024 05:59:57 +0000 (06:59 +0100)]
Importance Matrix calculation (llama/4861)

* imatrix: 1st version

* imatrix: WIP

* Cleanup

* Update examples/imatrix/imatrix.cpp

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

Co-authored-by: Iwan Kawrakow <redacted>
Co-authored-by: Georgi Gerganov <redacted>
2 years agosync : whisper.cpp
Georgi Gerganov [Thu, 11 Jan 2024 19:56:20 +0000 (21:56 +0200)]
sync : whisper.cpp

2 years agosync : ggml
Georgi Gerganov [Thu, 11 Jan 2024 19:54:17 +0000 (21:54 +0200)]
sync : ggml

2 years agomain : add cli option to disable system prints (whisper/1740)
Georgi Gerganov [Mon, 8 Jan 2024 14:41:28 +0000 (16:41 +0200)]
main : add cli option to disable system prints (whisper/1740)

2 years agosync : llama.cpp
Georgi Gerganov [Thu, 11 Jan 2024 19:49:13 +0000 (21:49 +0200)]
sync : llama.cpp

2 years agoggml : SOTA 2-bit quants (add IQ2_XS) (llama/4856)
Kawrakow [Thu, 11 Jan 2024 19:39:39 +0000 (20:39 +0100)]
ggml : SOTA 2-bit quants (add IQ2_XS) (llama/4856)

* iq2_xs: basics

* iq2_xs: this should have been in the basics

* iq2_xs: CUDA and scalar CPU works

* iq2_xs: WIP Metal

* iq2_xs: Metal now works

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

* iq2_xs: better ARM_NEON dot product

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

* iq2_xs: AVX2 dot product - 19.5 t/s

* iq2_xs: faster AVX2 dit product

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

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

* Add llama enum for IQ2_XS

---------

Co-authored-by: Iwan Kawrakow <redacted>
2 years agometal : put encoder debug group behind a define (llama/4873)
Paul Tsochantaris [Thu, 11 Jan 2024 14:31:52 +0000 (14:31 +0000)]
metal : put encoder debug group behind a define (llama/4873)

2 years agometal : improve dequantize precision to match CPU (llama/4836)
Georgi Gerganov [Tue, 9 Jan 2024 17:37:08 +0000 (19:37 +0200)]
metal : improve dequantize precision to match CPU (llama/4836)

ggml-ci

2 years agoggml : fix vld1q_s8_x4 32-bit compat (llama/4828)
Georgi Gerganov [Tue, 9 Jan 2024 08:42:06 +0000 (10:42 +0200)]
ggml : fix vld1q_s8_x4 32-bit compat (llama/4828)

* ggml : fix vld1q_s8_x4 32-bit compat

ggml-ci

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

ggml-ci

2 years agoCUDA: faster softmax via shared memory + fp16 math (llama/4742)
Johannes Gäßler [Tue, 9 Jan 2024 07:58:55 +0000 (08:58 +0100)]
CUDA: faster softmax via shared memory + fp16 math (llama/4742)

2 years agometal : fix deprecation warning (#690)
Georgi Gerganov [Thu, 11 Jan 2024 07:34:59 +0000 (09:34 +0200)]
metal : fix deprecation warning (#690)

2 years agoggml : remove ggml_cpy_inplace and ggml_cont_inplace (#693)
Timothy Cronin [Thu, 11 Jan 2024 07:27:48 +0000 (02:27 -0500)]
ggml : remove ggml_cpy_inplace and ggml_cont_inplace (#693)

2 years agoUpdate README.md (#692)
otaGran [Wed, 10 Jan 2024 16:12:18 +0000 (11:12 -0500)]
Update README.md (#692)

2 years agometal : wrap each operation in debug group (#690)
Jack Mousseau [Wed, 10 Jan 2024 14:19:19 +0000 (06:19 -0800)]
metal : wrap each operation in debug group (#690)

2 years agoggml : change GGML_MAX_NAME at compile time (#682)
leejet [Wed, 10 Jan 2024 13:13:42 +0000 (21:13 +0800)]
ggml : change GGML_MAX_NAME at compile time (#682)

* change GGML_MAX_NAME to 128

* allow controlling the value of GGML_MAX_NAME through external macro definitions

2 years agoFix execlp call (#689)
Halalaluyafail3 [Tue, 9 Jan 2024 16:16:37 +0000 (11:16 -0500)]
Fix execlp call (#689)

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

2 years agosync : llama.cpp
Georgi Gerganov [Mon, 8 Jan 2024 21:33:18 +0000 (23:33 +0200)]
sync : llama.cpp

2 years agoSOTA 2-bit quants (llama/4773)
Kawrakow [Mon, 8 Jan 2024 15:02:32 +0000 (16:02 +0100)]
SOTA 2-bit quants (llama/4773)

* iq2_xxs: basics

* iq2_xxs: scalar and AVX2 dot products

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

* iq2_xxs: ARM_NEON dot product

Somehow strangely slow (112 ms/token).

* iq2_xxs: WIP Metal

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

* iq2_xxs: Metal dot product now works

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

Not the greatest performance, but not complete garbage either.

* iq2_xxs: slighty faster dot product

TG-128 is now 48.4 t/s

* iq2_xxs: slighty faster dot product

TG-128 is now 50.9 t/s

* iq2_xxs: even faster Metal dot product

TG-128 is now 54.1 t/s.

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

* iq2_xxs: dequantize CUDA kernel - fix conflict with master

* iq2_xxs: quantized CUDA dot product (MMVQ)

We get TG-128 = 153.1 t/s

* iq2_xxs: slightly faster CUDA dot product

TG-128 is now at 155.1 t/s.

* iq2_xxs: add to llama ftype enum

* iq2_xxs: fix MoE on Metal

* Fix missing MMQ ops when on hipBLAS

I had put the ggml_supports_mmq call at the wrong place.

* Fix bug in qequantize_row_iq2_xxs

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

* Fixing tests

* PR suggestion

---------

Co-authored-by: Iwan Kawrakow <redacted>
2 years agoCUDA: fixed redundant value dequantization (llama/4809)
Johannes Gäßler [Sun, 7 Jan 2024 16:24:08 +0000 (17:24 +0100)]
CUDA: fixed redundant value dequantization (llama/4809)

2 years agoggml : use __builtin_amdgcn_sudot4 in __dp4a for gfx11 (llama/4787)
Konstantin Zhuravlyov [Sun, 7 Jan 2024 06:52:42 +0000 (01:52 -0500)]
ggml : use __builtin_amdgcn_sudot4 in __dp4a for gfx11 (llama/4787)

2 years agoggml : do not sched_yield when calling BLAS (llama/4761)
Georgi Gerganov [Fri, 5 Jan 2024 13:18:21 +0000 (15:18 +0200)]
ggml : do not sched_yield when calling BLAS (llama/4761)

* ggml : do not sched_yield when calling BLAS

ggml-ci

* ggml : fix do_yield logic

ggml-ci

* ggml : simplify do_yield logic

ggml-ci

2 years agoPrint backend name on test-backend-ops failure (llama/4751)
Johannes Gäßler [Thu, 4 Jan 2024 08:43:23 +0000 (09:43 +0100)]
Print backend name on test-backend-ops failure (llama/4751)

2 years agoggml : include stdlib.h before intrin.h (llama/4736)
Georgi Gerganov [Thu, 4 Jan 2024 08:12:26 +0000 (10:12 +0200)]
ggml : include stdlib.h before intrin.h (llama/4736)

2 years agofix ROCm on Windows (#683)
旺旺碎冰冰 [Sat, 6 Jan 2024 03:17:47 +0000 (11:17 +0800)]
fix ROCm on Windows (#683)

* fix ROCm on windows

* keep windows does not link m.lib

* add hip missed compile definitions

* never link m.lib

* use find_library for m

---------

Co-authored-by: slaren <redacted>
2 years agogguf : add keys for kv sizes to spec (#676)
postmasters [Fri, 5 Jan 2024 15:25:38 +0000 (07:25 -0800)]
gguf : add keys for kv sizes to spec (#676)

* Add keys for kv sizes to GGUF spec

* Fix types of key_length and value_length

2 years agofix : cuda order of synchronization when setting a buffer (#679)
Erik Scholz [Fri, 5 Jan 2024 15:00:00 +0000 (16:00 +0100)]
fix : cuda order of synchronization when setting a buffer (#679)

* fix : cuda order of synchronization when setting a buffer

* also sync before memcpy

---------

Co-authored-by: slaren <redacted>
2 years agometal : switch back to default.metallib (#681)
Georgi Gerganov [Fri, 5 Jan 2024 14:30:52 +0000 (16:30 +0200)]
metal : switch back to default.metallib (#681)

ggml-ci

2 years agoggml : fix q2_k bpw in comments (#680)
Georgi Gerganov [Fri, 5 Jan 2024 13:36:04 +0000 (15:36 +0200)]
ggml : fix q2_k bpw in comments (#680)

2 years agowhisper : reset the "batched" timings (whisper/1721)
Georgi Gerganov [Thu, 4 Jan 2024 11:37:25 +0000 (13:37 +0200)]
whisper : reset the "batched" timings (whisper/1721)

2 years agoggml : add error handling to graph_compute (whisper/1714)
Finn Voorhees [Wed, 3 Jan 2024 13:39:43 +0000 (08:39 -0500)]
ggml : add error handling to graph_compute (whisper/1714)

2 years agoscripts : allow to skip commits during sync (#678)
Georgi Gerganov [Wed, 3 Jan 2024 12:50:34 +0000 (14:50 +0200)]
scripts : allow to skip commits during sync (#678)

2 years agoUpdate src/ggml-cuda.cu
Georgi Gerganov [Wed, 3 Jan 2024 12:18:46 +0000 (14:18 +0200)]
Update src/ggml-cuda.cu

Co-authored-by: slaren <redacted>
2 years agocuda : mark I16 and I32 ops as unsupported
Georgi Gerganov [Wed, 3 Jan 2024 11:01:44 +0000 (13:01 +0200)]
cuda : mark I16 and I32 ops as unsupported

ggml-ci

2 years agometal : add kernel_get_rows_i32
Georgi Gerganov [Wed, 3 Jan 2024 09:35:46 +0000 (11:35 +0200)]
metal : add kernel_get_rows_i32

ggml-ci

2 years agosync : llama.cpp
Georgi Gerganov [Wed, 3 Jan 2024 09:25:18 +0000 (11:25 +0200)]
sync : llama.cpp

ggml-ci

2 years agometal : optimize ggml_mul_mat_id (faster Mixtral PP) (llama/4725)
Georgi Gerganov [Tue, 2 Jan 2024 19:07:47 +0000 (21:07 +0200)]
metal : optimize ggml_mul_mat_id (faster Mixtral PP) (llama/4725)

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

ggml-ci

* metal : fix Metal API debug warnings

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

* metal : fix API debug warnings

* metal : fix compile warnings

* metal : use uint64_t for strides

* cmake : rename option to LLAMA_METAL_SHADER_DEBUG

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

* metal : normalize mat-vec kernel signatures

* cmake : respect LLAMA_QKK_64 option

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

* metal : optimizing ggml_mul_mat_id (wip)

* metal : minor fix

* metal : opt mul_mm_id

2 years agometal : enable shader debugging (cmake option) (llama/4705)
Georgi Gerganov [Tue, 2 Jan 2024 08:57:44 +0000 (10:57 +0200)]
metal : enable shader debugging (cmake option) (llama/4705)

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

ggml-ci

* metal : fix Metal API debug warnings

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

* metal : fix API debug warnings

* metal : fix compile warnings

* metal : use uint64_t for strides

* cmake : rename option to LLAMA_METAL_SHADER_DEBUG

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

* metal : normalize mat-vec kernel signatures

* cmake : respect LLAMA_QKK_64 option

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

ggml-ci

2 years agoggml : add ggml_vdotq_s32 alias (llama/4715)
Georgi Gerganov [Sun, 31 Dec 2023 09:43:31 +0000 (11:43 +0200)]
ggml : add ggml_vdotq_s32 alias (llama/4715)

ggml-ci

2 years agoCUDA: fixed tensor cores not being used on RDNA3 (llama/4697)
Johannes Gäßler [Sat, 30 Dec 2023 12:52:01 +0000 (13:52 +0100)]
CUDA: fixed tensor cores not being used on RDNA3 (llama/4697)

2 years agoggml : add ggml_cpu_has_avx_vnni() (llama/4589)
automaticcat [Sat, 30 Dec 2023 08:07:48 +0000 (15:07 +0700)]
ggml : add ggml_cpu_has_avx_vnni() (llama/4589)

* feat: add avx_vnni based on intel documents

* ggml: add avx vnni based on intel document

* llama: add avx vnni information display

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

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

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

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

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

* Update ggml.c

Fix indentation upgate

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

Co-authored-by: Georgi Gerganov <redacted>
2 years agoCUDA: fix tensor core logic for Pascal and HIP (llama/4682)
Johannes Gäßler [Fri, 29 Dec 2023 22:12:53 +0000 (23:12 +0100)]
CUDA: fix tensor core logic for Pascal and HIP (llama/4682)

2 years agocuda: fix vmm oom issue on NVIDIA AGX Orin (llama/4687)
hydai [Fri, 29 Dec 2023 16:31:19 +0000 (00:31 +0800)]
cuda: fix vmm oom issue on NVIDIA AGX Orin (llama/4687)

Signed-off-by: hydai <redacted>
2 years agoscripts : fix sync order + metal sed
Georgi Gerganov [Wed, 3 Jan 2024 09:24:48 +0000 (11:24 +0200)]
scripts : fix sync order + metal sed

2 years agoswift : add Swift Package declaration (#674)
Ashraful Islam [Wed, 3 Jan 2024 09:07:30 +0000 (03:07 -0600)]
swift : add Swift Package declaration (#674)

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

* fixes the src paths in package.swift