]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
17 months agometal : show compile log messages
Georgi Gerganov [Thu, 25 Jan 2024 09:26:17 +0000 (11:26 +0200)]
metal : show compile log messages

17 months 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

17 months 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)

17 months 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

17 months 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)

17 months 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

17 months 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>
17 months 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>
17 months 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

17 months 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

17 months 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>
17 months 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>
17 months 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

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

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

17 months 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

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

17 months 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

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

17 months 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

17 months 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

17 months 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

17 months 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

17 months 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

17 months 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>
17 months 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>
17 months 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

17 months 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)

17 months 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>
17 months 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>
17 months agosync : whisper.cpp
Georgi Gerganov [Sun, 14 Jan 2024 09:07:09 +0000 (11:07 +0200)]
sync : whisper.cpp

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

17 months 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

17 months 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>
17 months agosync : whisper.cpp
Georgi Gerganov [Sat, 13 Jan 2024 22:14:03 +0000 (00:14 +0200)]
sync : whisper.cpp

17 months 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)

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

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

17 months 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)

17 months 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

17 months 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)

17 months 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>
17 months 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

17 months 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)

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

17 months 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

17 months 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

17 months 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)

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

ggml-ci

17 months 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

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

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

ggml-ci

17 months 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>
17 months 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)

17 months 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>
17 months agosync : whisper.cpp
Georgi Gerganov [Thu, 11 Jan 2024 19:56:20 +0000 (21:56 +0200)]
sync : whisper.cpp

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

17 months 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)

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

17 months 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>
17 months 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)

17 months 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

17 months 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

17 months 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)

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

17 months 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)

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

17 months 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)

17 months 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

17 months 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.

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

17 months 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>
17 months 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)

17 months 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)

17 months 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

17 months 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)

17 months 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)

17 months 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>
17 months 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

17 months 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>
17 months 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

17 months 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)

17 months 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)

17 months 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)

17 months 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)

17 months 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>
17 months 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

17 months 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

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

ggml-ci

17 months 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

17 months 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

17 months 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

17 months 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)

17 months 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>
17 months 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)

17 months 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>
17 months 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

17 months 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

18 months agoggml : extend ggml_get_rows, ggml_repeat, ggml_concat (#639)
Guillaume Wenzek [Fri, 29 Dec 2023 17:07:03 +0000 (18:07 +0100)]
ggml : extend ggml_get_rows, ggml_repeat, ggml_concat (#639)

* add more int ops

* ggml_compute_forward_dup_bytes

* add tests

* PR comments

* tests : minor indentations

---------

Co-authored-by: Georgi Gerganov <redacted>
18 months agoscripts : do not sync synced commits
Georgi Gerganov [Fri, 29 Dec 2023 13:17:18 +0000 (15:17 +0200)]
scripts : do not sync synced commits