]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
14 months agobuild: fix and ignore msvc warnings (#805)
Borislav Stanimirov [Thu, 25 Apr 2024 14:24:07 +0000 (17:24 +0300)]
build: fix and ignore msvc warnings (#805)

14 months agospec : fix typo in gguf.md (#798)
Daniel Bevenius [Thu, 18 Apr 2024 16:47:17 +0000 (18:47 +0200)]
spec : fix typo in gguf.md (#798)

Signed-off-by: Daniel Bevenius <redacted>
14 months agosync : whisper.cpp
Georgi Gerganov [Tue, 9 Apr 2024 17:28:35 +0000 (20:28 +0300)]
sync : whisper.cpp

14 months agowhisper : fix DTW memory access (whisper/2012)
Brad Murray [Tue, 9 Apr 2024 15:38:19 +0000 (11:38 -0400)]
whisper : fix DTW memory access (whisper/2012)

* Fix DTW memory access

* Memory fix - Apply changes from denersc

14 months agocommon : fix file-handle leak in read_wav() (whisper/2026)
ulatekh [Tue, 9 Apr 2024 15:34:34 +0000 (08:34 -0700)]
common : fix file-handle leak in read_wav() (whisper/2026)

Now it cleans up in case of error.

14 months agomain : set stdin to binary mode on Windows (whisper/2025)
Rotem Dan [Tue, 9 Apr 2024 15:33:32 +0000 (18:33 +0300)]
main : set stdin to binary mode on Windows (whisper/2025)

14 months agomain : allow a response-file as the sole parameter (whisper/2019)
ulatekh [Tue, 9 Apr 2024 15:31:16 +0000 (08:31 -0700)]
main : allow a response-file as the sole parameter (whisper/2019)

* The "main" example now allows a response-file as the sole parameter.

A response-file is a text file with command-line parameters, one per line.
Prefix the name of the response-file with "@" to identify it as such.
It's used under MS Windows to work around command-line length limits.
It may be useful under other platforms to simplify character-escaping.

* minor : style

---------

Co-authored-by: Georgi Gerganov <redacted>
14 months agowhisper : suppress tokens with a regex (whisper/1997)
ulatekh [Tue, 9 Apr 2024 15:27:28 +0000 (08:27 -0700)]
whisper : suppress tokens with a regex (whisper/1997)

* Allow a regular expression to describe tokens to suppress.

Example: --suppress-tokens-re "[,\.]|[ ]?[0-9]+" will suppress commas, periods, and numeric tokens.

Technique inspired by https://github.com/openai/whisper/discussions/1041

Co-authored-by: Georgi Gerganov <redacted>
* Blind change to fix Java test.

---------

Co-authored-by: Georgi Gerganov <redacted>
14 months agosync : llama.cpp
Georgi Gerganov [Tue, 9 Apr 2024 17:18:09 +0000 (20:18 +0300)]
sync : llama.cpp

14 months agolicense : update copyright notice + add AUTHORS
Georgi Gerganov [Tue, 9 Apr 2024 17:17:51 +0000 (20:17 +0300)]
license : update copyright notice + add AUTHORS

14 months agollama : add Command R Plus support (llama/6491)
Carolinabanana [Tue, 9 Apr 2024 08:16:13 +0000 (09:16 +0100)]
llama : add Command R Plus support (llama/6491)

* Add Command R Plus GGUF

* Add Command R Plus GGUF

* Loading works up to LayerNorm2D

* Export new tensors in 1D so they are not quantized.

* Fix embedding layer based on Noeda's example

* Whitespace

* Add line

* Fix unexpected tokens on MPS. Re-add F16 fix. ((Noeda)

* dranger003: Fix block index overflow in CUDA dequantizing.

* Reverted blocked multiplication code as it still has issues and could affect other Llama arches

* export norms as f32

* fix overflow issues during quant and other cleanup

* Type convention

Co-authored-by: Georgi Gerganov <redacted>
* dranger003: Fix more int overflow during quant.

---------

Co-authored-by: S <redacted>
Co-authored-by: S <redacted>
Co-authored-by: slaren <redacted>
Co-authored-by: Georgi Gerganov <redacted>
14 months agoremove row=1 cond (llama/6532)
Abhilash Majumder [Mon, 8 Apr 2024 08:26:01 +0000 (13:56 +0530)]
remove row=1 cond (llama/6532)

14 months agosupport/fix OPs GGML_TYPE_IQ4_NL, GGML_TYPE_IQ4_XS, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ3_...
Neo Zhang Jianyu [Sun, 7 Apr 2024 02:55:59 +0000 (10:55 +0800)]
support/fix OPs GGML_TYPE_IQ4_NL, GGML_TYPE_IQ4_XS, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ3_S, GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M (llama/6521)

14 months agoscripts : sync LICENSE and gen-authors.sh
Georgi Gerganov [Tue, 9 Apr 2024 17:15:04 +0000 (20:15 +0300)]
scripts : sync LICENSE and gen-authors.sh

14 months agosync : whisper.cpp (grammar-parser, skip)
Georgi Gerganov [Sun, 7 Apr 2024 14:02:17 +0000 (17:02 +0300)]
sync : whisper.cpp (grammar-parser, skip)

14 months agosync : whisper.cpp
Georgi Gerganov [Sun, 7 Apr 2024 13:22:06 +0000 (16:22 +0300)]
sync : whisper.cpp

14 months agoggml: bypass code incompatible with CUDA < 11.1 (whisper/2020)
Slava Primenko [Thu, 4 Apr 2024 12:49:24 +0000 (14:49 +0200)]
ggml: bypass code incompatible with CUDA < 11.1 (whisper/2020)

`cudaHostRegisterReadOnly` parameter was only introduced in CUDA 11.1

See this issue for more details:
https://github.com/ggerganov/examples/whisper/whisper.cpp/issues/2007

14 months agomain : add command-style grammar (whisper/1998)
ulatekh [Thu, 28 Mar 2024 10:02:10 +0000 (03:02 -0700)]
main : add command-style grammar (whisper/1998)

* Implemented command-style grammar in the main example.

Mostly just copied the relevant parts from the command example.

* main : code style

---------

Co-authored-by: Georgi Gerganov <redacted>
14 months agoscripts : sync ggml-cuda folder
Georgi Gerganov [Sun, 7 Apr 2024 13:09:49 +0000 (16:09 +0300)]
scripts : sync ggml-cuda folder

14 months agoci : temporary build embeded metal library
Georgi Gerganov [Sat, 6 Apr 2024 15:21:33 +0000 (18:21 +0300)]
ci : temporary build embeded metal library

ggml-ci

14 months agosync : llama.cpp (skip)
Georgi Gerganov [Sat, 6 Apr 2024 14:50:21 +0000 (17:50 +0300)]
sync : llama.cpp (skip)

ggml-ci

14 months agoFixed minor bug when enabling FP16 for non intel targets (llama/6464)
Ouadie EL FAROUKI [Fri, 5 Apr 2024 13:35:06 +0000 (14:35 +0100)]
Fixed minor bug when enabling FP16 for non intel targets (llama/6464)

* moved INTEL_MKL guard from gemm_impl to gemm (wrapper)

* Update ggml-sycl.cpp

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

Co-authored-by: AidanBeltonS <redacted>
14 months agoggml : mul_mat_id use the same tensor for all the experts (llama/6387)
slaren [Wed, 3 Apr 2024 13:07:05 +0000 (15:07 +0200)]
ggml : mul_mat_id use the same tensor for all the experts (llama/6387)

* ggml : update mul_mat_id to use the same tensor for all the experts

* update cuda

* minor

* update metal

* update test-backend-ops

* fix cuda

* Update ggml-metal.m

Co-authored-by: Georgi Gerganov <redacted>
* update convert.py

* update convert-hf-to-gguf.py

* update convert.py for mixtral hf models

* Update convert-hf-to-gguf.py

Co-authored-by: Georgi Gerganov <redacted>
* cuda : support non-pow-2 number of experts

* allow quantize to work for split and merged experts models in the same way

* cleanup + disable mmap automatically with split tensors models

* update imatrix

* test-backend-ops : test qwen argsort

* update grok model loading

* llama : add merged experts tensors to the grok tensor map

* minor

* gguf : bump version

* fix quantizing of merged experts

* convert-hf-to-gguf.py : update grok (untested)

* make linter happy

* cuda/argsort : use shared memory instead of pool memory

* convert : fix grok tensor names

* metal : add support for non-pow-2 argsort

* llama : more loader cleanup, better error checking

* cuda : fix warning

* llama : still use mmap for loading old models, but copy the data to a host buffer

* add review note

* llama : remove ffn tensor counting + add sanity check

ggml-ci

* convert : fix handling of n_experts == None

ggml-ci

* imatrix : fix ncall counters

* llama : produce error if imatrix size does not match

* quantize : terminate on errors + trace logs

ggml-ci

* metal : pad shared memory to 16 bytes

---------

Co-authored-by: Georgi Gerganov <redacted>
14 months agoDisable iqx on windows as WA (llama/6435)
Meng, Hengyu [Wed, 3 Apr 2024 02:34:40 +0000 (10:34 +0800)]
Disable iqx on windows as WA (llama/6435)

* disable iqx on windows as WA

* array instead of global_memory

14 months agoVulkan k-quant mmq and ggml-backend offload functionality (llama/6155)
0cc4m [Fri, 29 Mar 2024 16:29:21 +0000 (17:29 +0100)]
Vulkan k-quant mmq and ggml-backend offload functionality (llama/6155)

* Fix Vulkan no kv offload incoherence

* Add k-quant mul mat mat shaders

* Rework working buffer allocation, reduces vram use noticeably

Clean up cpu assist code, replaced with ggml-backend offload function

* Default to all dedicated GPUs

* Add fallback for integrated GPUs if no dedicated GPUs are found

* Add debug info which device is allocating memory

* Fix Intel dequant issue

Fix validation issue

* Fix Vulkan GGML_OP_GET_ROWS implementation

* Clean up merge artifacts

* Remove Vulkan warning

14 months agofix set main gpu crash (llama/6339)
Neo Zhang Jianyu [Thu, 28 Mar 2024 00:55:24 +0000 (08:55 +0800)]
fix set main gpu crash (llama/6339)

14 months agoggml : fix bounds checking of zero size views (llama/6347)
slaren [Wed, 27 Mar 2024 14:07:50 +0000 (15:07 +0100)]
ggml : fix bounds checking of zero size views (llama/6347)

14 months agobackend : fix typo in scheduler documentation (#781)
Daniel Bevenius [Wed, 3 Apr 2024 20:57:20 +0000 (22:57 +0200)]
backend : fix typo in scheduler documentation (#781)

Signed-off-by: Daniel Bevenius <redacted>
14 months agocmake : add `GGML_METAL_EMBED_LIBRARY` (#780)
Andrei [Wed, 3 Apr 2024 20:56:58 +0000 (16:56 -0400)]
cmake : add `GGML_METAL_EMBED_LIBRARY` (#780)

14 months agologo : add files (#782)
Georgi Gerganov [Wed, 3 Apr 2024 19:59:55 +0000 (22:59 +0300)]
logo : add files (#782)

15 months agogguf : update type enum (#775)
JacobLinCool [Wed, 27 Mar 2024 17:48:56 +0000 (01:48 +0800)]
gguf : update type enum (#775)

* spec: add missing semicolons in GGUF structs

Co-Authored-By: 郝東彥 Arthur Hao <redacted>
* spec: update GGUF tensor types

---------

Co-authored-by: 郝東彥 Arthur Hao <redacted>
15 months agosync : whisper.cpp
Georgi Gerganov [Wed, 27 Mar 2024 17:01:42 +0000 (19:01 +0200)]
sync : whisper.cpp

15 months agosync : whisper.cpp
Georgi Gerganov [Wed, 27 Mar 2024 11:37:14 +0000 (13:37 +0200)]
sync : whisper.cpp

15 months agowhisper : improve handling of prompts (whisper/1981)
Georgi Gerganov [Mon, 25 Mar 2024 12:48:19 +0000 (14:48 +0200)]
whisper : improve handling of prompts (whisper/1981)

* whisper : improve handling of prompts

* whisper : add whisper_token_count helper

15 months agowhisper : improve support for distil-large-v3 (whisper/1982)
Sanchit Gandhi [Thu, 21 Mar 2024 16:53:30 +0000 (22:23 +0530)]
whisper : improve support for distil-large-v3 (whisper/1982)

15 months agowhisper : token-level timestamps with DTW (whisper/1485)
denersc [Wed, 20 Mar 2024 16:25:26 +0000 (13:25 -0300)]
whisper : token-level timestamps with DTW (whisper/1485)

* whisper.cpp: impl dtw algo

* WIP: producing and placing DTW timestamps on tokens

* Fix compile and assertion errors. Attempt to DTW timestamp with single_segment=false.

* Fix mistake causing incorrect alignment of dtw timestamps

* implement N_TOP_MOST and CUSTOM alignment heads setting

* whisper: fix typo on alignment heads enum

* Fix issues related to changes in whisper.cpp

* Fixed excessive memory use when using DTW timestamps. Other minor fixes to DTW timestamping function

* decoder: save cross QKs only if requested

* Calling median filter with ggml_map_custom1

* Reimpl aheads n_top_most and custom. Sanity checks on chosen aheads

* Copying cross QKs from decoder backend correctly

* dtw: cleanup

* Fix incorrect n_frames passed to dtw when near end of audio

* Fix aheads_masks_init for backend != CPU

* whisper : minor style

* main : add dtw (wip)

* whisper: fix invalid memory access in aheads_masks_init

* main : add dtw (cont)

* whisper : minor

---------

Co-authored-by: Georgi Gerganov <redacted>
15 months agoexamples : rename --audio-context to --audio-ctx per help text (whisper/1953)
Jo Liss [Mon, 18 Mar 2024 15:53:33 +0000 (15:53 +0000)]
examples : rename --audio-context to --audio-ctx per help text (whisper/1953)

15 months agoexamples : more CUDA leftovers (#0)
Georgi Gerganov [Wed, 27 Mar 2024 11:10:13 +0000 (13:10 +0200)]
examples : more CUDA leftovers (#0)

15 months agoexamples : fix CUBLAS leftovers (#0)
Georgi Gerganov [Wed, 27 Mar 2024 11:09:12 +0000 (13:09 +0200)]
examples : fix CUBLAS leftovers (#0)

ggml-ci

15 months agosync : adapt to CUDA changes (#0)
Georgi Gerganov [Wed, 27 Mar 2024 09:26:04 +0000 (11:26 +0200)]
sync : adapt to CUDA changes (#0)

ggml-ci

15 months agosync : llama.cpp
Georgi Gerganov [Wed, 27 Mar 2024 09:15:00 +0000 (11:15 +0200)]
sync : llama.cpp

ggml-ci

15 months agoFix batched impl for NVidia GPU (llama/6164)
AidanBeltonS [Wed, 27 Mar 2024 08:16:40 +0000 (08:16 +0000)]
Fix batched impl for NVidia GPU (llama/6164)

* Fix batched impl

* Maintain previous behaviour for igpu

* retrigger CI

---------

Co-authored-by: Abhilash Majumder <redacted>
15 months agoMake IQ1_M work for QK_K = 64 (llama/6327)
Kawrakow [Wed, 27 Mar 2024 07:44:27 +0000 (08:44 +0100)]
Make IQ1_M work for QK_K = 64 (llama/6327)

* iq1_m: make it work for QK_K = 64 (WIP)

* iq1_m: make it work for QK_K = 64 (scalar and AVX2)

* iq1_m: QK_K = 64 seems to work on Metal and ARM_NEON

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agollama : greatly reduce output buffer memory usage (llama/6122)
compilade [Tue, 26 Mar 2024 14:46:41 +0000 (10:46 -0400)]
llama : greatly reduce output buffer memory usage (llama/6122)

* llama : greatly reduce logits memory usage

* llama : more compact state saving and reloading

* llama : fix lctx.n_outputs not being set before building graph

* perplexity : adapt to the logits API changes

* perplexity : fix Winogrande, use correct logits for second choice start

The first logits used to evaluate the second choice were not from
the end of the common prefix; instead, they were the logits from the end
of the first choice. This has been corrected.

The previous implementation sometimes had outliers in the scores of
choices for some tasks, and the logic to skip choices words
in the log-likelihood evaluation probably was an attempt to reduce those,
but it was complex and didn't quite seem to be the right thing.

This is simpler now, and the outlier scores aren't there anymore.

* perplexity : normalize spaces and punctuation in Winogrande sentences

* llama : fix embedding conditions

* llama : fix llama_get_embeddings_ith when the resulting id is 0

* llama : fix wrong n_outputs in llama_set_inputs

A mismatch happened when using a smaller n_ubatch than n_batch and then using
llama_batch_get_one(). The decision of what n_outputs should be now almost
fully depends on how lctx.n_outputs is set in llama_decode_internal.
The conditions are simpler this way.

* llama : when saving the state, recalculate n_outputs

This ensures the correct number of outputs for the entire previous batch
is stored in the session file, even when n_ubatch is smaller than n_batch.

* llama : fix not-skipping outputs of non-causal models

* llama : fix running a batch with n_outputs == 0

It previously worked because lctx.inp_out_ids was not initialized,
so it pointed to some garbage address which was somehow still valid when I
ran my tests.

* llama : keep same graph topology even when n_outputs == 0

* ggml : saner ggml_can_repeat with empty tensors

*  ggml : future-proof ggml_is_empty by using GGML_MAX_DIMS - 1

* ggml : do not multi-thread ops returning empty tensors

* ggml : make ggml_is_empty public and work with views

* llama : use a vector for ctx->output_ids

* llama : rework reallocation logic for llama_output_reserve

Now comparing the actual size with the new total size of the output buffer
to allow more efficient enabling and disabling of the embeddings
and/or logits output in the future.

* ggml : skip empty tensors in all backends

* llama : fix llama_output_reserve nullptr deref when new_size is 0

* perplexity : make Winogrande work as it does on master

The problems with the Winogrande implementation will
need to be fixed in a separate PR to ease review.

* llama : clearer error messages for invalid logits or embeddings ids

* llama : assert all models that can have inp_out_ids

Since the graph topology is now constant, this presence check
can be done even when there are no outputs.

* llama : assert logits and embd buffers exist before writing to them

* llama : handle errors from llama_output_reserve at call sites

* perplexity : make hellaswag and multiple-choice outputs identical to master

Due to how the KV cache is updated, the logprobs for tokens in a batch
are very slightly affected by the other tokens present in the batch,
so to make hellaswag and multiple-choice return exactly the same results
as on master, the last token of each sequence needs to be evaluated
even though its output is not used at all.

This will probably be changed back in the future to make these benchmarks
a tiny bit faster.

* perplexity : fix division by zero when using less than 100 multiple-choice tasks

* llama : allow loading state saved with a different ctx size

When loading a session file, the context size is now only required to be
at least enough to load the KV cells contained in that session file,
instead of requiring to use exactly the same context size as when saving.

Doing this enables the use-case of extending or shrinking the context size
of a saved session.

This breaks existing session files because the meaning of kv_buf_size
is slightly changed (previously it was the size of the whole KV cache,
now it's only the size of the saved part of it). This allows for
finer-grained sanity checks when loading in an effort to keep kv_buf_size
useful even when the kv_size is changed.

* llama : minor

ggml-ci

* readme : update recent API changes, and warn about Vulkan

---------

Co-authored-by: Georgi Gerganov <redacted>
15 months agoIQ1_M: 1.75 bpw quantization (llama/6302)
Kawrakow [Tue, 26 Mar 2024 14:21:27 +0000 (15:21 +0100)]
IQ1_M: 1.75 bpw quantization (llama/6302)

* iq1_m: basics

* iq1_m: basics-2

* iq1_m: CUDA dequantize works

Very 1st shot I get PPL = 9.76 for LLaMA-v2-7B.

* iq1_m: separate shifts for each group of 8 in a block

We get
PPL(LLaMA-v2-7B ) = 9.2810
PPL(LLaMA-v2-13B) = 6.8105

Not bad, but slightly higher than
  sqrt(PPL(IQ1_S) * PPL(IQ2_XXS))
which is the expected outcome given that IQ1_M is
halfway between IQ1_S and IQ2_XXS in terms of bpw.
From this, we would expect
 PPL = 9.14 for LLaMA-v2-7B
 PPL = 6.63 for LLaMA-v2-13B

* iq1_m: go to 3-bit scales

There is slight increase in PPL, but the 0.0625 bpw reduction
in size is totally worth it.

We now have
PPL(LLaMA-v2-7B ) = 9.4469 at 1.96 bpw
PPL(LLaMA-v2-13B) = 6.8717 at 1.93 bpw
PPL(LLaMA-v2-70B) = 4.8568 at 1.85 bpw

* iq1_m: scalar dot product

* iq1_m: AVX2 dot product

* iq1_m: very slightly faster AVX2 dot product

* iq1_m: ARM_NEON dot product

Works, but very slow (10.5 t/s)

* iq1_m: Metal - dequantize works, dot product does not

* iq1_m: Metal now works

About the same performance as iq1_s.

* iq1_m: minor

* iq1_m: checking pure iq1_m quantization

It is pretty bad: PPL(LLaMA-v2-7B) = 34 if we quantize output.weight
with Q4_K.

* iiq1_m: slightly faster ARM_NEON dot product

10.5 t/s -> 11.65 t/s

* iq1_m: faster ARM_NEON dot product

11.65 t/s -> 14.9 t/s

* iq1_m: another minor ARM_NEON dot product improvement

14.9 -> 15.0 t/s

* iq1_m: small PPL improvement via super-block scale adjustment

After quantizing block scales redo the super-block scale fit.

PPL(LLaMA-v2-7B ) = 9.3346
PPL(LLaMA-v2-13B) = 6.8419
PPL(LLaMA-v2-70B) = 4.8294
PPL(Mistral-7B  ) = 8.1624

* iq1_m: adapt to CUDA refactoring

* iq1_m: remove unused variable

We have progressed to warnings being errors.

* iq1_m: add to backend-ops tests

* iq1_m: fix Windows ARM

* iq1_m: use common definition of iq1m_scale_t

* cuda: assert -> NO_DEVICE_CODE

* iq1_M: PR comments

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agocuda : rename build flag to LLAMA_CUDA (llama/6299)
slaren [Tue, 26 Mar 2024 00:16:01 +0000 (01:16 +0100)]
cuda : rename build flag to LLAMA_CUDA (llama/6299)

15 months agotests : include IQ2_XXS and IQ2_XS in test-quantize-fns (llama/6303)
Kawrakow [Mon, 25 Mar 2024 17:33:15 +0000 (18:33 +0100)]
tests : include IQ2_XXS and IQ2_XS in test-quantize-fns (llama/6303)

Co-authored-by: Iwan Kawrakow <redacted>
15 months agocuda : refactor into multiple files (llama/6269)
slaren [Mon, 25 Mar 2024 12:50:23 +0000 (13:50 +0100)]
cuda : refactor into multiple files (llama/6269)

15 months agoggml : support AVX512VNNI (llama/6280)
Justine Tunney [Mon, 25 Mar 2024 05:39:56 +0000 (01:39 -0400)]
ggml : support AVX512VNNI (llama/6280)

This change causes some quants (e.g. Q4_0, Q8_0) to go faster on some
architectures (e.g. AMD Zen 4).

15 months agoFix heap corruption from wmode out-of-bound writes on windows (llama/6272)
Rick G [Sun, 24 Mar 2024 21:45:56 +0000 (14:45 -0700)]
Fix heap corruption from wmode out-of-bound writes on windows (llama/6272)

* would throw error on VS2022 on GGML_FREE(wmode)
* wchar_t is usually 2 bytes, but malloc wants bytes
  * therefore `*wmode_p++ = (wchar_t)*mode;` could write off the end of the allocation
* Fixes error possibly introduced by https://github.com/ggerganov/llama.cpp/pull/6248

15 months agooffload op (llama/6217)
Meng, Hengyu [Sun, 24 Mar 2024 04:04:25 +0000 (12:04 +0800)]
offload op (llama/6217)

* remove no USM methods

* leave the schedule to ggml_backend_sched entirely

15 months agouse _wfopen instead of fopen on Windows (llama/6248)
Jared Van Bortel [Sat, 23 Mar 2024 22:48:02 +0000 (18:48 -0400)]
use _wfopen instead of fopen on Windows (llama/6248)

also fix missing #defines before windows.h, and BPE LF token on MSVC

15 months agocuda : add LLAMA_CUDA_NO_PEER_COPY to workaround broken ROCm p2p copy (llama/6208)
slaren [Fri, 22 Mar 2024 13:05:31 +0000 (14:05 +0100)]
cuda : add LLAMA_CUDA_NO_PEER_COPY to workaround broken ROCm p2p copy (llama/6208)

* cuda : add LLAMA_CUDA_NO_PEER_COPY to workaround broken ROCm p2p copy

* add LLAMA_CUDA_NO_PEER_COPY to HIP build

15 months agometal : proper assert for mat-mat memory alignment (llama/6225)
Georgi Gerganov [Fri, 22 Mar 2024 09:35:53 +0000 (11:35 +0200)]
metal : proper assert for mat-mat memory alignment (llama/6225)

* metal : proper assert for mat-mat memory alignment

ggml-ci

* readme : add notice about the bug fix

* metal : fix the fix

ggml-ci

15 months agometal : pad n_ctx by 32 (llama/6177)
Georgi Gerganov [Fri, 22 Mar 2024 07:36:03 +0000 (09:36 +0200)]
metal : pad n_ctx by 32 (llama/6177)

* metal : require ne00 >= 128 for mat-mat kernels

ggml-ci

* llama : pad n_ctx by 32

ggml-ci

15 months agocuda : disable host register by default (llama/6206)
slaren [Thu, 21 Mar 2024 18:54:28 +0000 (19:54 +0100)]
cuda : disable host register by default (llama/6206)

15 months agocuda : fix LLAMA_CUDA_F16 build (llama/6197)
slaren [Thu, 21 Mar 2024 12:59:53 +0000 (13:59 +0100)]
cuda : fix LLAMA_CUDA_F16 build (llama/6197)

15 months agoggml : same IQ4_NL quantization for CPU/CUDA/Metal (llama/6196)
Kawrakow [Thu, 21 Mar 2024 12:59:38 +0000 (13:59 +0100)]
ggml : same IQ4_NL quantization for CPU/CUDA/Metal (llama/6196)

* Make quantize_row_iq4_nl do the same thing is quantization on CUDA

* Make quantize_row_iq4_nl do the same thing is quantization on CUDA

This time for real. backend-ops tests pass.

* Now fix test-quantize-fns

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agoAdd ability to use Q5_0, Q5_1, and IQ4_NL for quantized K cache (llama/6183)
Kawrakow [Thu, 21 Mar 2024 07:27:57 +0000 (08:27 +0100)]
Add ability to use Q5_0, Q5_1, and IQ4_NL for quantized K cache (llama/6183)

* k_cache: be able to use Q5_0

* k_cache: be able to use Q5_1 on CODA

* k_cache: be able to use Q5_0 on Metal

* k_cache: be able to use Q5_1 on Metal

* k_cache: be able to use IQ4_NL - just CUDA for now

* k_cache: be able to use IQ4_NL on Metal

* k_cache: add newly added supported types to llama-bench and CUDA supports_op

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agoAdd nvidia and amd backends (llama/6157)
AidanBeltonS [Thu, 21 Mar 2024 06:10:52 +0000 (06:10 +0000)]
Add nvidia and amd backends (llama/6157)

15 months agocuda : fix conflict with std::swap (llama/6186)
slaren [Thu, 21 Mar 2024 00:47:46 +0000 (01:47 +0100)]
cuda : fix conflict with std::swap (llama/6186)

15 months agocuda : print the returned error when CUDA initialization fails (llama/6185)
slaren [Wed, 20 Mar 2024 20:03:26 +0000 (21:03 +0100)]
cuda : print the returned error when CUDA initialization fails (llama/6185)

15 months agocuda : refactor to remove global resources (llama/6170)
slaren [Wed, 20 Mar 2024 13:42:59 +0000 (14:42 +0100)]
cuda : refactor to remove global resources (llama/6170)

* cuda : refactor to remove global resources

15 months agoincrease igpu cluster limit (llama/6159)
Abhilash Majumder [Wed, 20 Mar 2024 02:58:49 +0000 (08:28 +0530)]
increase igpu cluster limit (llama/6159)

15 months agobackend : set max split inputs to GGML_MAX_SRC (llama/6137)
slaren [Mon, 18 Mar 2024 15:33:44 +0000 (16:33 +0100)]
backend : set max split inputs to GGML_MAX_SRC (llama/6137)

15 months agobackend : offload large batches to GPU (llama/6083)
slaren [Mon, 18 Mar 2024 10:03:04 +0000 (11:03 +0100)]
backend : offload large batches to GPU (llama/6083)

* backend : offload large batches to GPU

* fix hip

* code cleanup

* fix CUDA split buffers

* Update ggml-backend-impl.h

Co-authored-by: Johannes Gäßler <redacted>
* cuda : fix memset without set_device

* imatrix : remove sched affix from weight names

* sched : add a new split if the current one has too many inputs
reduce max inputs per split
more cleanup

* update backends

ggml-ci

---------

Co-authored-by: Johannes Gäßler <redacted>
15 months agoggml:fix finding transfer queue family index error (llama/6094)
GainLee [Sun, 17 Mar 2024 17:12:22 +0000 (01:12 +0800)]
ggml:fix finding transfer queue family index error (llama/6094)

Co-authored-by: GainLee <redacted>
15 months agoggml : add AVX512F SIMD (llama/6088)
AmirAli Mirian [Sat, 16 Mar 2024 15:52:02 +0000 (11:52 -0400)]
ggml : add AVX512F SIMD (llama/6088)

15 months agocuda : disable unused cudaLaunchHostFunc code (llama/6078)
slaren [Fri, 15 Mar 2024 12:24:03 +0000 (13:24 +0100)]
cuda : disable unused cudaLaunchHostFunc code (llama/6078)

15 months agofix set main gpu error (llama/6073)
Neo Zhang Jianyu [Fri, 15 Mar 2024 10:53:53 +0000 (18:53 +0800)]
fix set main gpu error (llama/6073)

15 months agoFix non-intel device selection (llama/6042)
AidanBeltonS [Fri, 15 Mar 2024 09:26:20 +0000 (09:26 +0000)]
Fix non-intel device selection (llama/6042)

* Fix non-intel device selection

* Update ggml-sycl.cpp

Co-authored-by: Neo Zhang Jianyu <redacted>
* Update ggml-sycl.cpp

Co-authored-by: Neo Zhang Jianyu <redacted>
---------

Co-authored-by: Abhilash Majumder <redacted>
Co-authored-by: Neo Zhang Jianyu <redacted>
15 months agogguf : add support for I64 and F64 arrays (llama/6062)
Ondřej Čertík [Fri, 15 Mar 2024 08:46:51 +0000 (02:46 -0600)]
gguf : add support for I64 and F64 arrays (llama/6062)

* gguf : add support for I64 and F64 arrays

GGML currently does not support I64 or F64 arrays and they are not often
used in machine learning, however if in the future the need arises, it
would be nice to add them now, so that the types are next to the other
types I8, I16, I32 in the enums, and it also reserves their type number.

Furthermore, with this addition the GGUF format becomes very usable for
most computational applications of NumPy (being compatible with the most
common NumPy dtypes: i8, i16, i32, i64, f32, f64), providing a faster,
and more versatile alternative to the `npz` format, and a simpler
alternative to the `hdf5` format.

The change in this PR seems small, not significantly increasing the
maintenance burden. I tested this from Python using GGUFWriter/Reader
and `gguf-dump`, as well as from C, everything seems to work.

* Fix compiler warnings

15 months agocmake : bump to 3.12 (#768)
Bryan Lozano [Fri, 22 Mar 2024 07:18:20 +0000 (00:18 -0700)]
cmake : bump to 3.12 (#768)

This CMakeLists.txt file uses Cmake command `add_compile_definitions()`.

That command did not exist until CMake 3.12.

Reference documentation:
https://cmake.org/cmake/help/latest/command/add_compile_definitions.html

I see a failure in CMake 3.10 with existing CMakeLists.txt

15 months agoexamples : fix simple (#770)
Bryan Lozano [Fri, 22 Mar 2024 07:17:34 +0000 (00:17 -0700)]
examples : fix simple (#770)

* Update README.md

Correcting matrix multiplication expected result.

* Update simple-ctx.cpp

Fix incorrect striding through output.

* simple : update readme

---------

Co-authored-by: Georgi Gerganov <redacted>
15 months agosync : whisper.cpp
Georgi Gerganov [Mon, 18 Mar 2024 09:03:55 +0000 (11:03 +0200)]
sync : whisper.cpp

15 months agowhisper : set outputs from conv graph (whisper/1959)
Georgi Gerganov [Sat, 16 Mar 2024 15:30:55 +0000 (17:30 +0200)]
whisper : set outputs from conv graph (whisper/1959)

15 months agoalloc : fix allocation data of pre-allocated leafs
slaren [Sat, 16 Mar 2024 14:47:14 +0000 (15:47 +0100)]
alloc : fix allocation data of pre-allocated leafs

15 months agosync : whisper.cpp
Georgi Gerganov [Fri, 15 Mar 2024 12:22:55 +0000 (14:22 +0200)]
sync : whisper.cpp

15 months agospec : add GGUF diagram (#765)
Georgi Gerganov [Fri, 15 Mar 2024 12:10:35 +0000 (14:10 +0200)]
spec : add GGUF diagram (#765)

15 months agoupdate examples and tests
slaren [Thu, 14 Mar 2024 15:45:27 +0000 (16:45 +0100)]
update examples and tests

15 months agoggml : add ggml-common.h
Georgi Gerganov [Thu, 14 Mar 2024 15:16:45 +0000 (17:16 +0200)]
ggml : add ggml-common.h

15 months agosync : llama.cpp
Georgi Gerganov [Thu, 14 Mar 2024 15:10:55 +0000 (17:10 +0200)]
sync : llama.cpp

15 months agoggml : designate enum vals for integer types (llama/6050)
Georgi Gerganov [Thu, 14 Mar 2024 10:38:37 +0000 (12:38 +0200)]
ggml : designate enum vals for integer types (llama/6050)

15 months agometal : build metallib + fix embed path (llama/6015)
Georgi Gerganov [Thu, 14 Mar 2024 09:55:23 +0000 (11:55 +0200)]
metal : build metallib + fix embed path (llama/6015)

* metal : build metallib + fix embed path

ggml-ci

* metal : fix embed build + update library load logic

ggml-ci

* metal : fix embeded library build

ggml-ci

* ci : fix iOS builds to use embedded library

15 months agollama : add pipeline parallelism support (llama/6017)
slaren [Wed, 13 Mar 2024 17:54:21 +0000 (18:54 +0100)]
llama : add pipeline parallelism support (llama/6017)

* llama : add pipeline parallelism support for batch processing with multiple CUDA GPUs

ggml-ci

* server : add -ub, --ubatch-size parameter

* fix server embedding test

* llama : fix Mamba inference for pipeline parallelism

Tested to work correctly with both `main` and `parallel` examples.

* llama : limit max batch size to n_batch

* add LLAMA_SCHED_MAX_COPIES to configure the number of input copies for pipeline parallelism
default increase to 4 (from 2)

changing this value may improve performance for some systems, but increases memory usage

* fix hip build

* fix sycl build (disable cpy_tensor_async)

* fix hip build

* llama : limit n_batch and n_ubatch to n_ctx during context creation

* llama : fix norm backend

* batched-bench : sync after decode

* swiftui : sync after decode

* ggml : allow ggml_get_rows to use multiple threads if they are available

* check n_ubatch >= n_tokens with non-casual attention

* llama : do not limit n_batch to n_ctx with non-casual attn

* server : construct batch with size of llama_n_batch

* ggml_backend_cpu_graph_compute : fix return value when alloc fails

* llama : better n_batch and n_ubatch comment

* fix merge

* small fix

* reduce default n_batch to 2048

---------

Co-authored-by: Francis Couture-Harpin <redacted>
Co-authored-by: Georgi Gerganov <redacted>
15 months agotest-backend-ops : skip CPU backend by default (llama/6028)
slaren [Wed, 13 Mar 2024 13:58:30 +0000 (14:58 +0100)]
test-backend-ops : skip CPU backend by default (llama/6028)

15 months agoUpdate get version (llama/6025)
AidanBeltonS [Wed, 13 Mar 2024 13:17:54 +0000 (13:17 +0000)]
Update get version (llama/6025)

15 months agoggml : reuse quantum structs across backends (llama/5943)
Georgi Gerganov [Tue, 12 Mar 2024 12:27:20 +0000 (14:27 +0200)]
ggml : reuse quantum structs across backends (llama/5943)

* ggml : reuse quant blocks across backends

ggml-ci

* ggml : define helper constants only for CUDA and SYCL

ggml-ci

* ggml : define helper quantum constants for SYCL

ggml-ci

15 months agoggml : fix UB in IQ2_S and IQ3_S (llama/6012)
Georgi Gerganov [Tue, 12 Mar 2024 11:49:55 +0000 (13:49 +0200)]
ggml : fix UB in IQ2_S and IQ3_S (llama/6012)

15 months agosycl : update IQ1_S kernels (WIP - not working!) (llama/5995)
Georgi Gerganov [Tue, 12 Mar 2024 09:15:05 +0000 (11:15 +0200)]
sycl : update IQ1_S kernels (WIP - not working!) (llama/5995)

* sycl : try to fix after IQ1_S changes

* sycl : iq1s_grid -> iq1s_grid_gpu

* sycl : fix grid type

15 months ago1.5 bit: we can do even better (llama/5999)
Kawrakow [Mon, 11 Mar 2024 15:53:15 +0000 (16:53 +0100)]
1.5 bit: we can do even better (llama/5999)

* iq1_s: we can do even better

Spent one of the 4 scale bits on a signs of a 0.125 shift.
I.e., quants are now -1 + delta, delta, 1 + delta, where delta
is +/- 0.125.

CUDA works, same performance as before.
PPL(LLaMA-v2-7B) is now 11.85!

* iq1_s: make scalar and AVX2 work with the new version

* iq1_s: make Neon work with new version.

~10% drop in performance, so will need some more work.

* iq1_s: make Metal work with new version

* iq1_s: very slightly faster dequantize on Metal

* iq1_s: fix dequantize on the CPU

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agoggml, ci : Windows ARM runner and build fixes (llama/5979)
Michael Podvitskiy [Mon, 11 Mar 2024 09:28:51 +0000 (10:28 +0100)]
ggml, ci : Windows ARM runner and build fixes (llama/5979)

* windows arm ci

* fix `error C2078: too many initializers` with ggml_vld1q_u32 macro for MSVC ARM64

* fix `warning C4146: unary minus operator applied to unsigned type, result still unsigned`

* fix `error C2065: '__fp16': undeclared identifier`

15 months agoBetter 1.5 bit quantization (llama/5971)
Kawrakow [Mon, 11 Mar 2024 06:51:49 +0000 (07:51 +0100)]
Better 1.5 bit quantization (llama/5971)

* Trying blocvks of 16 for IQ1_S - seems slightly better

* iq1s_blocks16: Adjust scale fudge factor to 1.125

* iq1s_blocks16: going to blocks of 32

with 2048 lattice points, so same bpw.
This is even better than blocks of 16.
Should I try blocks of 64? But to keep the same
bpw, when I go to 4096 lattice points, I need to
remove blocks alltogether and just have superblocks of
256 weights.

* iq1s_blocks16: Use 2*<x^2> as sigma2 in weight adjustment

* iq1s_blocks16: scalar and AVX2 dot products

* iq1s_blocks16: CUDA dot product

* iq1s_blocks16: Metal works, Neon does not

Metal works but TG is dog slow (35 t/s). PP is OKish (493 t/s).
Not seeing the bug in the Neon implementation for now.

* iq1s_blocks16: fixed Neon

* iq1s_blocks16: very slightly faster TG on Metal

Still pathetic at 37 t/s

* iq1s_blocks16: speedup Metal by packing codebook into uint32_t's

* Formatting

* iq1s_blocks16: uint32_t codebook is also better in CUDA

TG-128 is now 204 t/s up from 194 t/s.
PP-512 is 5890 t/s, so significantly better than other quants

* iq1s_blocks16: slightly faster Neon dot product

* iq1s_blocks16: faster AVX2 dot product

* iq1s_blocks16: adjust to ggml-common.h

---------

Co-authored-by: Iwan Kawrakow <redacted>
15 months agoAdd q3_s and q1_s (llama/5886)
Abhilash Majumder [Mon, 11 Mar 2024 04:57:56 +0000 (10:27 +0530)]
Add q3_s and q1_s (llama/5886)

* Add q3_s and q1_s

* fix compilation

* fix build

* fix build

* fix build

* enable ops

* rm macro

* increase grid space

15 months agometal : move mm_id indices to shared mem (llama/5982)
Georgi Gerganov [Sun, 10 Mar 2024 21:12:48 +0000 (23:12 +0200)]
metal : move mm_id indices to shared mem (llama/5982)

15 months agoggml : remove __constant__ specifier for CUDA tables (llama/5940)
Georgi Gerganov [Sun, 10 Mar 2024 18:09:24 +0000 (20:09 +0200)]
ggml : remove __constant__ specifier for CUDA tables (llama/5940)

15 months agoggml : fix unnecessary f32 -> f16 -> f32 casts (mmla) (llama/5951)
Georgi Gerganov [Sat, 9 Mar 2024 15:36:20 +0000 (17:36 +0200)]
ggml : fix unnecessary f32 -> f16 -> f32 casts (mmla) (llama/5951)

15 months agoggml : remove old quantization functions (llama/5942)
Georgi Gerganov [Sat, 9 Mar 2024 13:53:59 +0000 (15:53 +0200)]
ggml : remove old quantization functions (llama/5942)

* ggml : remove old quantization functions

ggml-ci

* ggml : simplify ggml_quantize_chunk

ggml-ci

* ggml : restrict correctness

ggml-ci

* ggml : remove hist data from the quantization API

ggml-ci

* tests : remove hist usage in test-backend-ops

ggml-ci

* vulkan : remove hist and fix typo

15 months agoggml : add ggml-common.h to deduplicate shared code (llama/5940)
Georgi Gerganov [Sat, 9 Mar 2024 10:47:57 +0000 (12:47 +0200)]
ggml : add ggml-common.h to deduplicate shared code (llama/5940)

* ggml : add ggml-common.h to shared code

ggml-ci

* scripts : update sync scripts

* sycl : reuse quantum tables

ggml-ci

* ggml : minor

* ggml : minor

* sycl : try to fix build

15 months agollama : support Mamba Selective State Space Models (llama/5328)
compilade [Fri, 8 Mar 2024 22:31:00 +0000 (17:31 -0500)]
llama : support Mamba Selective State Space Models (llama/5328)

* mamba : begin working on support for Mamba SSM

* mamba : begin figuring out how to (ab)use the kv cache for Mamba

* mamba : recurrent inference almost works, but incoherent

* mamba : recurrent inference WORKS!!!

* convert : optionally use d_conv and d_state from config.json for Mamba

* mamba : refactor recurrent conv, resulting in 20% perf increase

It's still slower than I'd like, but I did not really optimize `ggml_exp` yet.

I also refactored `ggml_exp` to work with tensors with more than 2 dimensions.

* ggml : parallelize ggml_exp

This results in 8% faster token generation for Mamba-130M.

* mamba : simplify the conv step with a self-overlapping view

Turns out the conv_state can be made smaller by one column.
Note that this breaks existing GGUFs of Mamba,
because the key_value_length field is tied to the conv_state size.

Convolution with a self-overlapping view is cool!
And it's much simpler than what I initially thought would be necessary
to make the convolution step work with more than 1 token at a time.

Next step is to make the SSM step work on batches of tokens too,
and thus I need to figure out a way to make a parallel selective scan
which will keep the ssm_state small and won't make it bigger
by a factor of (n_layer * batch_size).

* llama : fix Mamba KV self size wrongly displaying as f16 instead of f32

Relatedly, I also tried to see if other types than f32 worked for the states,
but they don't, because of the operators used.
It's probably better anyway to keep lots of precision there,
since the states are small anyway.

* mamba : fix self-overlapping view depth stride

* mamba : handle batches of more than 1 token

This means running Mamba no longer crashes when using the default settings!
And probably also slightly faster prompt processing.
Both batched and non-batched processing yield the same output.

Previously, the state was not cleared when starting a sequence.
Next step is to make the KV cache API work as expected for Mamba models.

* ggml: add ggml_ssm_scan to help with parallel selective scan

If the selective scan was implemented without a custom operator,
there would be waaay too many nodes in the graph. For example,
for Mamba-130M, with a batch size of 512 (the default),
a naive selective scan could add at least 24*512=12288 nodes,
which is more than LLAMA_MAX_NODES (8192),
and that's only for the smallest Mamba model.
So it's much cleaner with a custom operator.
Not sure about the name, though.

* ggml : in ggml_ssm_scan, merge multiple rows in the same vec operation

This will help with performance on CPU if ggml_vec_mul_f32
and ggml_vec_add_f32 are ever optimized with SIMD.

* mamba : very basic quantization support

Mostly works, but there is currently no difference
between the variants of a k-quant (e.g. Q4_K_S and Q4_K_M are the same).
Most of the SSM-specific weights can be kept in f32 without affecting
the size that much, since they are relatively small.
(the linear projection weights are responsible for most of Mamba's size)

Too much quantization seems to make the state degrade quite fast, and
the model begins to output gibberish.
It seems to affect bigger models to a lesser extent than small models,
but I'm not sure by how much.

Experimentation will be needed to figure out which weights are more important
for the _M (and _L?) variants of k-quants for Mamba.

* convert : fix wrong name for layer norm weight of offical Mamba models

I was using Q-bert/Mamba-* models before, which have a slighlty different
naming scheme for the weights.
(they start with "model.layers" instead of "backbone.layers")

* mamba : fuse more steps of the SSM scan in the ggml_ssm_scan operator

This increases performance on CPU by around 30% for prompt processing,
and by around 20% for text generation.

However, it also makes the ggml_exp and ggml_soft_plus operators unused.
Whether or not they should be kept will be decided later.

* convert : for Mamba, also consider the "MambaLMHeadModel" arch name

It's the name of the class of the official implementation,
though they don't use it (yet) in the "architectures" field of config.json

* mamba : fix vocab size problems with official models

The perplexity was waaaay to high for models with a non-round vocab size.
Not sure why, but it needed to be fixed in the metadata.

Note that this breaks existing GGUF-converted Mamba models,
but **only if** the vocab size was not already rounded.

* ggml : remove ggml_exp and ggml_soft_plus

They did not exist anyway outside of this branch,
and since ggml_ssm_scan fused operations together, they are unused.
It's always possible to bring them back if needed.

* mamba : remove some useless comments

No code change.

* convert : fix flake8 linter errors

* mamba : apply suggestions from code review

* mamba : remove unecessary branch for row-wise ssm_state and C multiplication

It was previously done to avoid permuting when only one token is processed
at a time (like when generating text), but permuting is cheap,
and dynamically changing the compute graph is not future-proof.

* ggml : in ggml_ssm_scan, use more appropriate asserts

* ggml : rename the destination pointer in ggml_compute_forward_ssm_scan_f32

* mamba : multiple sequences, but one at a time

This is a step towards making this Mamba implementation usable
with the server example (the way the system prompt is kept when clearing
the client slots will need to be changed before this can work, though).

The KV cache size for this kind of model is tied to the maximum number
of sequences kept at any single time.
For now, this number is obtained from n_parallel (plus one,
to have an extra sequence to dedicate to the system prompt),
but there might be a better way to do this which won't also
make the main example use 2 cells even if only 1 is really used.
(for this specific case, --parallel 0 helps)

Simultaneous sequence processing will probably require changes to
ggml_ssm_scan, and possibly a new operator for the conv step.

* mamba : support llama_kv_cache_seq_cp

This (mis)uses the logic around K shifts, because tokens in a state
can't be shifted anyway, and because inp_K_shift has the right shape and type.
Using ggml_get_rows is a nice way to do copies, but copy chains can't work.
Fortunately, copy chains don't really seem to be used in the examples.

Each KV cell is dedicated to the sequence ID corresponding to its own index.

* mamba : use a state mask

It's cleaner than the previous heuristic of
checking for the pos of the first token in the batch.

inp_KQ_mask could not be re-used for this, because it has the wrong shape
and because it seems more suited to the next step of
simultaneous sequence processing (helping with the problem of
remembering which token belongs to which sequence(s)/state(s)).

* llama : replace the usage of n_ctx with kv_self.size in many places

* mamba : use n_tokens directly instead of n_tok

* mamba : in comments, properly refer to KV cells instead of slots

* mamba : reduce memory usage of ggml_ssm_scan

From 290.37 MiB to 140.68 MiB of CPU compute buffer size
with Mamba 3B with a batch size of 512.

The result tensor of ggml_ssm_scan was previously a big part
of the CPU compute buffer size. To make it smaller,
it does not contain the intermediate ssm states anymore.
Both y and the last ssm state are combined in the result tensor,
because it seems only a single tensor can be returned by an operator
with the way the graph is built.

* mamba : simultaneous sequence processing

A batch can now contain tokens from multiple sequences.

This is necessary for at least the parallel example, the server example,
and the HellaSwag test in the perplexity example.

However, for this to be useful, uses of llama_kv_cache_seq_rm/cp
will need to be changed to work on whole sequences.

* ggml : add ggml_ssm_conv as a new operator for the conv step of Mamba

This operator makes it possible to use and update the correct states
for each token of the batch in the same way as ggml_ssm_scan.
Other solutions which use existing operators would need loops which would
add too many nodes to the graph (at least the ones I thought of).

Using this operator further reduces the size of the CPU compute buffer
from 140.68 MiB to 103.20 MiB with Mamba 3B with a batch size of 512.
And (at least on CPU), it's a bit faster than before.

Note that "ggml_ssm_conv" is probably not the most appropriate name,
and it could be changed if a better one is found.

* llama : add inp_s_seq as a new input tensor

The most convenient implementation to select the correct state (for Mamba)
for each token is to directly get the correct index from a tensor.
This is why inp_s_seq is storing int32_t and not floats.

The other, less convenient way to select the correct state would be
to have inp_KQ_mask contain 1.0f for each state used by a token
and 0.0f otherwise. This complicates quickly fetching the first used
state of a token, and is also less efficient because a whole row
of the mask would always need to be read for each token.

Using indexes makes it easy to stop searching when there are
no more sequences for a token, and the first sequence assigned
is always very quickly available (it's the first element of each row).

* mamba : support llama_kv_cache_seq_cp copy chains

* mamba : support shifting and dividing the kv cache pos

* mamba : make the server and parallel examples work with whole sequences

A seq_id is dedicated to the system prompt in both cases.

* llama : make llama_kv_cache_seq_rm return whether it succeeded or not

* mamba : dedicate an input tensor for state copy indices

This is cleaner and makes it easier to adapt when/if token positions
(and by extension, inp_K_shift) are no longer integers.

* mamba : adapt perplexity, batched, and batched-bench examples

* perplexity : limit the max number of sequences

This adapts to what the loaded model can provide.

* llama : add llama_n_max_seq to get the upper limit for seq_ids

Used by the perplexity example.

* batched : pass n_parallel to the model's context params

This should have been there already, but it wasn't.

* batched-bench : reserve sequences to support Mamba

* batched-bench : fix tokens being put in wrong sequences

Generation quality isn't what's measured in there anyway,
but at least using the correct sequences avoids using non-consecutive
token positions.

* mamba : stop abusing attention metadata

This breaks existing converted-to-GGUF Mamba models,
but will allow supporting mixed architectures like MambaFormer
without needing to break Mamba models.

This will also allow changing the size of Mamba's states
without having to reconvert models in the future.
(e.g. using something else than d_conv - 1 columns for the conv_states
 will not require breaking existing converted Mamba models again)

* gguf-py : add new KV metadata key-value pairs for Mamba

* llama : add new metadata key-value pairs for Mamba

* llama : guard against divisions by zero when n_head is 0

* mamba : rename "unlimited" KV cache property to "recurrent"

* mamba : more correctly update the "used" field of the KV cache

* ggml : in ggml_ssm_scan, use a threshold for soft_plus

This is how the official Mamba implementation does it,
and it's also what torch.nn.Softplus does.

* convert : for Mamba, fallback to internal NeoX tokenizer

The resulting models are exactly the same
as if the tokenizer.json and tokenizer_config.json of GPT-NeoX were there.

* mamba : support state saving and restoring

* ggml : implicitly pass src tensors through dst for Mamba-related ops

* mamba : clarify some comments

* server : fix cache_tokens not getting correctly resized

Otherwise, when the "we have to evaluate at least 1 token" special case
was triggered, an extra token was kept in cache_tokens even if it was
removed from the KV cache.

For Mamba, this caused useless prompt reprocessing when the previous
request triggered the above case.

* convert-hf : support new metadata keys for Mamba

For the models available at
https://huggingface.co/collections/state-spaces/transformers-compatible-mamba-65e7b40ab87e5297e45ae406

* mamba : rename metadata to be more similar to transformers library

This breaks existing converted-to-GGUF models,
but the metadata names are more "standard".

* mamba : support mamba-*-hf models

These models share their token_embd.weight with their output.weight

* mamba : add missing spaces

This is purely a formatting change.

* convert-hf : omit output.weight when identical with token_embd.weight

Only for Mamba for now, but it might be relevant for other models eventually.
Most Mamba models actually share these two tensors, albeit implicitly.

* readme : add Mamba to supported models, and add recent API changes

* mamba : move state_seq and state_mask views outside layer loop

A few tensors were also missing `struct` in front of `ggml_tensor`.