Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

CUDA : faster k-quant dot kernels #1862

Merged
merged 5 commits into from
Jun 16, 2023
Merged

CUDA : faster k-quant dot kernels #1862

merged 5 commits into from
Jun 16, 2023

Conversation

ikawrakow
Copy link
Contributor

After the CUDA inference speed up in PR #1827 I got motivated to improve the performance of the k-quants, and this PR is the result.

The following table compares token prediction times on current master (i.e., after merging #1827) and this PR for all k-quantization levels. Shown is the average of 3 runs in milliseconds predicting 128 tokens on an RTX-4080 GPU using the following command:

./bin/main -m ../models/7B/qXk.bin -p "I believe the meaning of life is" -c 2048 -n 512 --ignore-eos -n 128 -s 1234 -t 1 -ngl 100
Quantization Master This PR Speedup
Q2_K 11.94 10.62 12.4%
Q3_K_S 15.35 11.68 31.4%
Q4_K_S 12.09 10.98 10.1%
Q5_K_S 13.53 12.03 12.5%
Q6_K 14.51 13.33 8.9%
Q4_0 12.40 - -

I have added Q4_0, not touched by this PR, for comparison. It is now more than 10% slower than Q4_K, bit it should be faster as it needs to do less work, so there is a potential for improvement there (on Metal, Q4_0 is ~10% faster than Q4_K).

I have moved away from using a templated implementation. There are subtle differences between the different k-quants in what is the best way to spread the threads in a thread group within a quantization block, and these variations are difficult to capture in a template. The result is ~100 extra lines of code compared to the templated implementation, which I think is worth it, given the increase in performance.

Copy link
Collaborator

@JohannesGaessler JohannesGaessler left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I didn't look at the k-quant code until now because I was preoccupied with other things. The code added in this PR looks good to me assuming it works as it should. I will advise you to be cautious with GPU performance optimizations though. Optimal values for e.g. the block size are not the same across different cards; it's possible that this PR causes a performance regression for some GPUs.

ggml-cuda.cu Outdated
Comment on lines 1213 to 1214
const dim3 block_dims(32, 1, 1);
dequantize_mul_mat_vec_q2_k<<<nrows, block_dims, 0, stream>>>(vx, y, dst, ncols);
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You could just pass 32 directly instead of block_dims which I think would be a little cleaner since it avoids mixing scalars with three-dimensional vectors.

@JohannesGaessler
Copy link
Collaborator

JohannesGaessler commented Jun 14, 2023

I did some performance testing:

GPU Model ms/t master ms/t PR
RTX 3090 7b q2_K 17.24 16.49
RTX 3090 7b q3_K_S 21.31 18.42
RTX 3090 7b q4_K_S 17.74 17.11
RTX 3090 7b q5_K_S 18.84 17.98
RTX 3090 7b q6_K 18.76 17.62
RTX 3090 7b q4_0 15.76 -
GTX 1070 7b q2_K 91.87 127.91
GTX 1070 7b q3_K_S 91.59 81.96
GTX 1070 7b q4_K_S 129.12 134.02
GTX 1070 7b q5_K_S 144.76 136.66
GTX 1070 7b q6_K 81.73 143.29
GTX 1070 7b q4_0 62.61 -

On my RTX 3090 the new kernels are consistently faster. q4_0 is still faster though (maybe you didn't compile with LLAMA_CUDA_DMMV_X=64 LLAMA_CUDA_DMMV_Y=2?). On my GTX 1070 the performance is only better for two out of the five quantization formats. For q2_K and q6_K the performance is significantly worse.

Copy link
Collaborator

@slaren slaren left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be great if the performance of q2_K and q6_K on older GPUs can be fixed, but if you don't have a Pascal GPU for testing, I think it is ok prioritizing the performance on newer architectures.

@ikawrakow
Copy link
Contributor Author

@JohannesGaessler Thank you for taking the time to look into this PR and posting results for different cards. Do you have a hypothesis why the new version tuns slower on the 1070?

You are right, i did not use LLAMA_CUDA_DMMV_X=64 LLAMA_CUDA_DMMV_Y= 2 for the Q4_0 result. Using this, Q4_0 token prediction time becomes 11.5 ms on my 4080 card, so still slower than Q4_K_S.

@JohannesGaessler
Copy link
Collaborator

It would be great if the performance of q2_K and q6_K on older GPUs can be fixed, but if you don't have a Pascal GPU for testing, I think it is ok prioritizing the performance on newer architectures.

I personally would not merge an optimization that causes a small performance uplift for some users but a large performance malus for other users but I am neither the project lead nor the developer of k-quants so I'll leave the decision up to @ikawrakow

Do you have a hypothesis why the new version tuns slower on the 1070?

I don't. GPU performance optimizations depend very strongly on the hardware details but the difference between an RTX 3090 and a GTX 1070 is quite large. It oftentimes comes down to something like the number of warps that fit into the available registers on a streaming multiprocessor and it's much easier to just tweak parameters and test performance than to try to understand what's happening on the hardware side. It's also possible that you are now using some feature that just has bad performance on older cards.

@ggerganov
Copy link
Owner

I'm also leaning towards improving the performance on newer devices and optionally having compile flags / macros for better performance on older hardware given that this does not complicate or duplicate the code too much (i.e. similar to the GGML_CUDA_DMMV_X, GGML_CUDA_DMMV_Y defines).

But before making any decision, let's wait for one or two more reports that confirm the degraded performance on older cards, just in case

@ikawrakow
Copy link
Contributor Author

I did get my hands on a box with a GTX-1660 (released Oct 2019). Q3_K, Q4_K and Q5_K are faster with this PR, Q2_K and Q6_K are slower. So, I guess, I'll do something with ifdefs for Q2_K and Q6_K.

If I remember correctly, back then the 1660 was basically the cheapest option if one wanted to have a discrete GPU but didn't want to pay the premium Nvidia added with the introduction of the 2000 series. The card has 6 GB of VRAM and the 4-bit models fail to run on it with -ngl 100 on current master. They run with -ngl 32, but performance is lower compared to what we had before #1827 was merged, so it seems we merged a PR without being too much worried about the experience of users on older hardware.

I also got my hands on an even older PC gathering dust in the corner. It has a 1060 card in it, so I wanted to experiment with that. But after fighting with Nvidia drivers and CUDA toolkits for a while, I eventually gave up. Are we aware that none of the CUDA stuff will work with the default packages on Ubuntu 20.04 LTS? I any case, while at it, I ran Q4_K_S on the CPU (i9-9960X, released 2018) and it was faster than the numbers @JohannesGaessler posted for the 1070 card.

@johnson442
Copy link
Contributor

Some more data from the type of hardware in question - 1070ti and i5-8600K:

Master PR CPU (-t 5)
7B
Q2_K 84.92 116.34 83.68
Q3_K_S 83.07 75.21 86.10
Q4_K_S 120.97 122.55 102.77
Q5_K_S 136.51 127.96 122.62
Q6_K 74.18 131.62 143.83
Q4_0 57.50 - 103.09
13B (-ngl 24 -t 5)
Q6_K 196.96 255.84 273.61

I've included CPU only (with 5 threads on 6 cores) to compare, as well as 24 layers of a 13B model on the GPU (max in 8GB of VRAM).

Forgive me if its explained elsewhere, but is there a reason Q6_K is faster than the other new quants on this GPU?

@maddes8cht
Copy link
Contributor

With my rtx 3060 using the #1827 build performance got actually worse when using the "extra layers" with 13b models.
Generation times got about 40% slower, and prompt eval time raised about 25 times up ( from something like 80 ms/token to 2000 ms/token)
Going back to -ngl 40 with 13b models braught me back a lot of performance.

How about making block sizes configurable at run time?
And maybe have a small helper app that tests out optimal parameters?

@ikawrakow
Copy link
Contributor Author

Forgive me if its explained elsewhere, but is there a reason Q6_K is faster than the other new quants on this GPU?

The Q6_K quantization is simpler. The weight w is obtained from the quant q as w = d * s * q, where d is the scale of the super-block and s is the scale of the block inside the super-block in which q is located. In Q6_K the scale is 8 bit, so it can be taken and used directly. Q3_K is similar, but there the scale is 6 bit, so one needs more compute (bit manipulations) to recover it from the densely stored data. For the other k-quants we have w = d * s_q * q - m * s_m, where now m is another scale for the super-block and s_q, s_m are per-block 6-bit parameters (4-bit in the case of Q2_K). As a result one needs way more bit fiddling to recover the scales and the quants.

When the computation becomes very fast as it is the case on the latest generation Nvidia cards, memory bandwidth becomes important, and this results in Q6_K being slower than the others on my RTX-4080. To give more specific numbers, we have Q6_K taking ~13.5 ms. If I put return as the very first statement in the dot product kernel, the computation completes in ~3.5 ms/token. I.e., the dot product takes ~10 ms. The Q6_K 7B model is 5.5 GB, so this means ~550 GB/s are being read from memory while the dot products are being computed. The RTX-4080 spec claims 768 GB/s memory bandwidth, but that would be under ideal conditions with all threads/warps operating in sync and fetching data from memory in the optimum way. We don't have that here as we are reading single bytes and jumping around within the data of a super-block to fetch scales/quants, so my guess is that 550 GB/s is as good as it gets (or at least very close to). In comparison, Q4_K_S runs in ~11 ms, so ~7.5 ms for dot products. The model is 3.8 GB, so ~500 GB/s are being read from VRAM while running inference with a Q4_K model.

@JohannesGaessler
Copy link
Collaborator

@maddes8cht I saw some people complain about how with the newest NVIDIA drivers the program no longer crashes when you run out of VRAM. Instead data from VRAM gets moved to RAM and you get terrible performance. And it's a "feature" that cannot be disabled.

We now have a K_QUANTS_PER_ITERATION macro, which should be
set to 1 on older and to 2 on newer GPUs.
With this, we preserve the performance of the original
PR on RTX-4080, and are faster compared to master on
GTX-1660.
Using the same K_QUANTS_PER_ITERATION macro as last commit,
we preserve performance on RTX-4080 and speed up
Q6_K on a GTX-1660.
Allowed values are 1 or 2. 2 gives the best performance on
modern GPUs and is set as default. On older GPUs 1 may work
better.
@ikawrakow
Copy link
Contributor Author

I put some effort into improving the performance of Q2_K and Q6_K on older GPUs, while preserving performance on newer cards. The oldest GPU I have available is a GTX-1660 released in 2019, so I would appreciate if someone tested on the the GTX-1000 series. To deal with older cards, I have added LLAMA_CUDA_KQUANTS_ITER to CMakeLists.txt and Makefile. It is set by default to 2 (best for RTX-4080). On older cards performance is better with LLAMA_CUDA_KQUANTS_ITER=1.

Here are the per token prediction times for 128 tokens using the 7B model after the last commits on a GTX-1660 using LLAMA_CUDA_KQUANTS_ITER=1 (performance on the RTX-4080 is the same as given in the initial description of the PR):

Quantization Master This PR Speedup
Q2_K 52.2 47.4 10.1%
Q3_K_S 65.6 60.3 8.8%
Q4_K_S 68.0 65.5 3.8%
Q5_K_S 74.4 68.5 8.6%
Q6_K 63.0 62.5 0.9%
Q4_0 50.6 - -
Q5_1 87.6 - -

The GTX-1660 has 6 GB of VRAM, so only Q2_K and Q3_K fit completely on the GPU (including KV cache) and are run with -ngl 100 -t 1. Q4_K and Q5_K were run with -ngl 32 -t 4, Q6_K with -ngl 29 -t 4, and Q5_1 with -ngl 30 -t 4.

Makefile Outdated Show resolved Hide resolved
@johnson442
Copy link
Contributor

johnson442 commented Jun 16, 2023

On a 1070ti:

Master (ms/t) New PR New PR ITER=1
7B
Q2_K 84.92 91.13 76.54
Q3_K_S 83.07 74.76 73.63
Q4_K_S 120.97 122.39 121.69
Q5_K_S 136.51 127.34 126.65
Q6_K 74.18 133.55 73.28
13B (-ngl 24 -t 5)
Q6_K 196.96 260.25 198.73

Faster than master in most cases, amazing stuff!

And thanks as well for the detailed explanation of Q6_K.

@ikawrakow ikawrakow merged commit 3d01122 into master Jun 16, 2023
@ikawrakow ikawrakow deleted the ik/cuda-faster-k-quants branch June 16, 2023 17:08
@johnson442
Copy link
Contributor

Should something be added to the README maybe, either a small CUDA performance section or just added to the CUBLAS instructions - mentioning LLAMA_CUDA_DMMV_X/Y and LLAMA_CUDA_KQUANTS_ITER?

I guess there is no K-quant stuff yet either, LLAMA_CUDA_KQUANTS_ITER might be best put in a future section describing them.

@mirek190
Copy link

mirek190 commented Jun 16, 2023

I tested with RTX 3090 and 33B q4k_m molel with 63 layers on GPU.

Before your patch I had 89 ms/t and now 82ms/t

Great work!!

@dragonfyre13
Copy link

dragonfyre13 commented Jun 21, 2023

Dropping a comment here because I seem to have independently come across very similar results downstream in koboldcpp, was actually originally coming to submit a PR to llama.cpp to have a compile time check if compute capability <7, if so set K_QUANTS_PER_ITERATION to 1 vs 2. On my 3070 vs p40, modifying the compile flag LLAMA_CUDA_KQUANTS_ITER makes a drastic difference (2x) in ms/t for the p40. I'm headed over to open something in the kobold.cpp side, but figured it's worth dropping confirmation about the difference in this ticket as well.

Given it seems to make enough difference, does it make sense to have a check for pascal architecture somewhere? My CPP is a bit rusty at this stage, but I can cobble together a PR if it'd help. I'm not sure without digging, but seems like there's an opportunity for a runtime check even, if the system has both a compute capability <7 device and a >7 device (i.e. they have a 3070 and a p40)

@cmp-nct
Copy link
Contributor

cmp-nct commented Jun 23, 2023

I've tested the update with Falcon 40 q6_k and it is a solid 6-7% boost in performance 92ms/tk went down to 86ms/tk for that large 38GB model. Awesome work!

@mirek190
Copy link

strange ... I tested latest build with rtx 3090 and 33B model q4k_m before i had around 82 ms/t and now I have 59 mn/s ....

byroneverson added a commit to byroneverson/llm.cpp that referenced this pull request Jun 30, 2023
commit b8c8dda75fdf5fdea49c80af36818e7c30fe0ddf
Author: Howard Su <[email protected]>
Date:   Thu Jun 29 21:15:15 2023 +0800

    Use unsigned for random seed (#2006)

    * Use unsigned for random seed. Keep -1 as the value to use a time based seed.

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 96a712ca1b7f427e3bd7ffc0c70b2105cfc7fbf1
Author: LostRuins <[email protected]>
Date:   Thu Jun 29 11:56:43 2023 +0800

    Porting the improved K-Quant CUDA kernels to OpenCL (#1966)

    * Added broken new q4k quant

    * xx + ib0

    * Fix q2_k fast kernel

    * Use preprocessor for QK_K

    * Add q6_k fast matmul kernel

    * ported q3k speedup successfully

    * ported q2k and q5k speedups

    * remove old dot kernels and template

    * fixed global const struct types

    * fixing address spaces

    * fixed string too long CI issue

    ---------

    Co-authored-by: 0cc4m <[email protected]>

commit d3494bb86bf7ad5b0b60aae0220ea576f273b5c0
Author: m3ndax <[email protected]>
Date:   Wed Jun 28 20:39:08 2023 +0200

    llama : replacing auto &kv with const auto &kv (#2041)

    * Replacing auto &kv with const auto &kv

    * Create codacy.yml

    * Delete codacy.yml

commit 5b351e94d041742cd50ffcf2d44718d63bab398a
Author: Salvador E. Tropea <[email protected]>
Date:   Wed Jun 28 14:27:31 2023 -0300

    cuda : remove nchannels_x argument from mul_mat_vec_nc_f16_f32 (#2028)

    - Not used

commit 6432aabb6dc887436e4d57414b63116189c3b13b
Author: Salvador E. Tropea <[email protected]>
Date:   Wed Jun 28 14:26:26 2023 -0300

    cuda : fix missing const qualifier in casts (#2027)

commit b922bc351b69770cec2d35d2aa50fa052b95ca93
Author: Howard Su <[email protected]>
Date:   Wed Jun 28 10:13:02 2023 -0700

    llama : remove shards weight file support (#2000)

    * Remove multiple shards

    * Remove multiple file loaders

    * Remove llama_load_tensor_shard class

    * Simplify load logic

    * Remove dead code guess_n_parts function

    * Remove vocab_only from constructor of llama_model_loader

    * Remove alignment_prevents_mmap which is not more needed.

    * Remove useless check

commit 7f9753fa1263c4eded9a3de19778562f0e1093d7
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 28 18:35:54 2023 +0200

    CUDA GPU acceleration for LoRAs + f16 models (#1970)

commit cfa0750bc9dbc2d957a91b8ed09ab0035d8f3d4e
Author: ningshanwutuobang <[email protected]>
Date:   Wed Jun 28 23:53:37 2023 +0800

    llama : support input embeddings directly  (#1910)

    * add interface for float input

    * fixed inpL shape and type

    * add examples of input floats

    * add test example for embd input

    * fixed sampling

    * add free for context

    * fixed add end condition for generating

    * add examples for llava.py

    * add READMD for llava.py

    * add READMD for llava.py

    * add example of PandaGPT

    * refactor the interface and fixed the styles

    * add cmake build for embd-input

    * add cmake build for embd-input

    * Add MiniGPT-4 example

    * change the order of the args of llama_eval_internal

    * fix ci error

commit 9d23589d638dc74577d5ff880e6d4248b795f12e
Author: Erik Scholz <[email protected]>
Date:   Tue Jun 27 19:06:33 2023 +0200

    fix pthreads setaffinity usage on android (#2020)

commit 0be54f75a6c3e9a09ea71bdfcdabf9a996a0549b
Author: Howard Su <[email protected]>
Date:   Tue Jun 27 13:07:13 2023 +0800

    baby-llama : fix build after ggml_rope change (#2016)

commit 181e8d975528a4e27eabb8ae6e9865f9ceae4b37
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 27 00:37:13 2023 +0300

    llama : fix rope usage after ChatGLM change

commit d9779021bd59ed96daae75e820a5ac5da47ca8ff
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 27 00:06:51 2023 +0300

    ggml : add support for ChatGLM RoPE

commit d38e45157862b58a1824387e64860d68ca3533a7
Author: Roman Parykin <[email protected]>
Date:   Mon Jun 26 22:47:59 2023 +0300

    readme : add Scala 3 bindings repo (#2010)

commit eaa6ca5a61b8c9501df9ebe3d264f45b75a5f8aa
Author: David Yang <[email protected]>
Date:   Tue Jun 27 03:45:32 2023 +0800

    ggml : increase max tensor name + clean up compiler warnings in train-text (#1988)

    * Clean up compiler warnings in train-text

    Some brackets to disambiguate order of operations

    * Increase GGML_MAX_NAME

    Avoiding strncpy danger in train-text-from-scratch and reducing potential future name length issues

commit aa777abbb73655c4e1e9237b7c0ad66745e8e48c
Author: Gustavo Rocha Dias <[email protected]>
Date:   Mon Jun 26 16:34:45 2023 -0300

    readme : LD_LIBRARY_PATH complement for some Android devices when building with CLBlast inside Termux (#2007)

    * docs - Alternative way to build at Android, with CLBlast.

    * doc - LD_LIBRARY_PATH complement for some Android devices when building with CLBlast inside Termux.

    * doc- fix typo

commit c824d2e368d193d9f564ff29880a51cda9f90527
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 21:03:59 2023 +0300

    ggml : avoid conv 2d kernel round up

commit b853d456018b10820686362af41b2f2f75f1eec6
Author: zrm <[email protected]>
Date:   Mon Jun 26 13:57:59 2023 -0400

    ggml : add NUMA support (#1556)

    * detect NUMA systems and pin work threads to nodes (linux)

    * disable mmap prefetch/readahead for NUMA systems

    * avoid sending finalize op to thread pool if it does nothing

    * silence robot

    * fix args

    * make --numa a param

    * recommendation that n_nodes evenly divide n_threads did not warrant such aggressive enforcement

    * lower synchronization overhead

    * statically allocate

    * move numa state to g_state

    * add description for --numa

    * ggml : minor style changes

    * ggml : minor style + try fix sanitizer build

    * llama : allow to initialize backend with NUMA support

    * llama : avoid ggml include in llama-util.h

    * ggml : style / formatting

    * ggml : fix handling of ops with n_threads > n_tasks > 1

    * server : utilize numa parameter

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 9225baef71407d799a6f7f563b77fd7f82791416
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 20:10:52 2023 +0300

    k-quants : fix indentation

commit a84ab1da8dc6a59a5b67420ae1322f09503ffc72
Author: katsu560 <[email protected]>
Date:   Tue Jun 27 01:47:02 2023 +0900

    tests : fix quantize perf (#1990)

    * fix test quantize perf

    * avoid the global state

commit 5743ca80928d8410754ec64a5673d5c2dd6cfbb7
Author: katsu560 <[email protected]>
Date:   Tue Jun 27 01:46:07 2023 +0900

    k-quants : add AVX support to dot functions (#1916)

    * k_quants : add AVX support

    * k_quants : apply review comments

commit 412c60e4739367144e51e59add5dc7749d084115
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 19:45:09 2023 +0300

    readme : add link to new k-quants for visibility

commit 6769e944c727c63612dcafbef52009d21ae00fff
Author: Kawrakow <[email protected]>
Date:   Mon Jun 26 19:43:07 2023 +0300

    k-quants : support for super-block size of 64 (#2001)

    * k_quants: WIP super-blocks with 64 weights

    * k_quants: WIP super-blocks with 64 weights

    Q6_K scalar and AVX2 works

    * k_quants: WIP super-blocks with 64 weights

    Q4_K scalar and AVX2 works

    * k_quants: WIP super-blocks with 64 weights

    Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
    than the scalar implementation)

    * k_quants: WIP super-blocks with 64 weights

    Q3_K scalar and AVX2 works.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K scalar and AVX2 works, and with that all
    k_quants are done on AVX2 and scalar

    * k_quants: WIP super-blocks with 64 weights

    Q6_K working on CUDA. Cannot make it run quite as gast as
    with super-blocks with 256 weigths: 8% slower on 4080,
    20% slower on the 1660 (but there we fit 1 less layer on the
    GPU because pf the larger model size), so some fraction of
    these 20% is due to that,

    * k_quants: WIP super-blocks with 64 weights

    Q4_K working on CUDA. ~10% slower on GTX-1660,
    16% slower on 4080.

    * k_quants: WIP super-blocks with 64 weights

    Q2_K working on CUDA. ~3% slower on GTX-1660,
    10% slower on 4080.

    * k_quants: WIP super-blocks with 64 weights

    Q3_K working on CUDA.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K working on CUDA, and with this CUDA is done.

    * k_quants: WIP super-blocks with 64 weights

    Q6_K working on ARM_NEON

    * k_quants: WIP super-blocks with 64 weights

    Q4_K working on ARM_NEON, but quite a bit slower than 256 weights

    * k_quants: WIP super-blocks with 64 weights

    Q2_K working on ARM_NEON, but quite a bit slower than 256 weights

    * k_quants: WIP super-blocks with 64 weights

    Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.

    With that, we have full support for ARM_NEON, although
    performance is not quite there.

    * k_quants: WIP super-blocks with 64 weights

    Slightly more efficient Q3_K and Q5_K

    * k_quants: WIP super-blocks with 64 weights

    Another small improvement for Q3_K and Q5_K on ARM_NEON

    * k_quants: WIP super-blocks with 64 weights

    Yet another speedup for Q5_K on ARM_NEON.
    We are now within 10% of the QK_K = 256 version.

    * k_quants: WIP super-blocks with 64 weights

    * We are able to pass preprocessor macros to the Metal
      compiler
    * Q6_K works and is actually slightly more efficient than
      the QK_K = 256 version (25.2 ms vs 25.8 ms)

    * k_quants: WIP super-blocks with 64 weights

    Q4_K works on Metal and is actually slightly faster
    than QK_K = 256 (21.95 ms vs 24.0 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q2_K works on Metal and is very slightly faster
    than QK_K = 256 (23.8 ms vs 24.2 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q3_K works on Metal and is slightly faster
    than QK_K = 256 (26.6 ms vs 28.3 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q5_K works on Metal and is slightly faster
    than QK_K = 256 (23.7 ms vs 26.3 ms).

    * k_quants: call them _K, not _k, also on Metal

    * k_quants: correctly define QK_K in llama.cpp

    * Fixed bug in q4_K quantization added with the 64-block addition

    * Simplify via lambda

    * k_quants: swicth Q3_K to 4-bit scales when QK_K = 64

    Otherwise there isn't much benefit from this
    quantization type. There is some very slight loss
    in accuracy, but we reduce size by ~7%.
    E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
    8.6131 with 8-bit scales and 8.6352 with 4-bit,
    while file size decreases from 1.53G to 1.44G.

    * k_quants: switch Q4_K to 4-bit scales when QK_K = 64

     Here the loss in accuracy is greater than for Q3_K,
     but the Q4_K points still move further to the left on
     the perplexity vs size curve.

    * k_quants: forgot to add the Metal changes in last commit

    * k_quants: change Q5_K to be type 0 when QK_K = 64

    Still needs AVX2 implementation

    * k_quants: AVX2 implementation for new 64-weight Q5_K

    * k_quants: 10% faster ARM_NEON Q5_K dot product

    * k_quants: fixed issue caused by merging with master

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit cbebf61ca7584e9709265395f0127ae7fc0f1882
Author: Howard Su <[email protected]>
Date:   Mon Jun 26 23:15:47 2023 +0800

    Fix assert when free invalid cuda pointer (#2005)

    Fix assert via initializing extra structure always.
    CUDA error 1 at C:\GPT\llama.cpp\ggml-cuda.cu:2536: invalid argument

commit 447ccbe8c39332fcdd0d98a041b6e2ff6f06219d
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 16:08:12 2023 +0300

    readme : add new roadmap + manifesto

commit bd34cdde38f8fd661890ddd5f57ca30bf279877b
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 14:25:08 2023 +0300

    ggml : sync latest ggml (custom operators)

commit c2a08f87b8d180115d04b8688f383d1b2761b16d
Author: anon998 <[email protected]>
Date:   Sun Jun 25 08:48:36 2023 +0000

    fix server sampling: top k sampler first (#1977)

    Co-authored-by: anon <[email protected]>

commit 66a2555ba6cab954c56d653b29c27bfbbacfbfb1
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 09:07:03 2023 +0300

    readme : add Azure CI discussion link

commit e65ca7e14ac76c4046091da39d41a9017abaa9b3
Author: sjinzh <[email protected]>
Date:   Sun Jun 25 13:45:44 2023 +0800

    zig : upgrade build system support (#1981)

    * upgrade zig build system support

    * zig : add new line at the end of the file

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 5ec8dd5a3c6a9a109351d2257bb9d53869bd0a94
Author: Robyn <[email protected]>
Date:   Sun Jun 25 04:10:29 2023 +1000

    #1869 Fix null reference errors when training from scratch with CUDA (#1907)

    * #1869 Fix null reference errors when training from scratch with CUDA build

    Calling ggml_compute_forward when node->src0 was null was causing train-text-from-scratch.exe to terminate unexpectedly.

    * ggml : do not dereference src0 if NULL

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 65bdd52a867539691007f85c5508146d507f72c1
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 24 19:40:18 2023 +0300

    tests : sync test-grad0 from ggml

commit fdd18609113862dc6eb34dfc44a093d54c59ff1f
Author: Rowan Hart <[email protected]>
Date:   Sat Jun 24 04:07:08 2023 -0700

    flake : fix ggml-metal.metal path and run nixfmt (#1974)

commit c943d823c14cef33092205ca3944de6fdf7abf99
Author: AN Long <[email protected]>
Date:   Sat Jun 24 19:02:06 2023 +0800

    convert : fix invalid params in write_vocab_only (#1975)

commit f2c754e1c38936fdde74e4848ac468a696eb73c6
Author: slaren <[email protected]>
Date:   Sat Jun 24 12:57:18 2023 +0200

    ggml : improve ggml_graph_dump_dot, add ggml_format_name (#1978)

    * Improve ggml_graph_dump_dot, add ggml_format_name

    * add more automatic names to view ops

    * fix name of copies

commit 11da1a85cd69af84b5861134738c7e9e20907470
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 24 13:38:18 2023 +0300

    readme : fix whitespaces

commit 235b610d650cbfed6dbd5d671f750d35fc18cd7d
Author: Alberto <[email protected]>
Date:   Sat Jun 24 12:32:13 2023 +0200

    readme : fixed termux instructions (#1973)

commit b061ba9e2a7a2c335a200df8c11aed5e31e4ccbb
Author: Alex Renda <[email protected]>
Date:   Sat Jun 24 03:15:01 2023 -0700

    llama : fix top-p sampling to match the canonical definition (#1953)

    * Fix top-p sampling to match the standard definition (smallest set that has probability mass at least p, not largest set with probability mass less than p)

    * top-p: correct gt to gte

    * add test for correct top-p behavior

commit 527b6fba1d237befb324fd846bda7418c0fa394d
Author: Didzis Gosko <[email protected]>
Date:   Sat Jun 24 11:47:58 2023 +0300

    llama : make model stateless and context stateful (llama_state) (#1797)

    * llama : make model stateless and context stateful

    * llama : minor cleanup

    * llama : update internal API declaration

    * Apply suggestions from code review

    fix style

    Co-authored-by: Georgi Gerganov <[email protected]>

    * Missing model memory release

    * Fix style

    * Add deprecated warning for public API function llama_init_from_file

    * Update public API use cases: move away from deprecated llama_init_from_file

    * Deprecate public API function llama_apply_lora_from_file

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit d7b7484f74d486f77feb4c0b7af7e1718ed91651
Author: eiery <[email protected]>
Date:   Fri Jun 23 04:38:01 2023 -0400

    Add OpenLLaMA instructions to the README (#1954)

    * add openllama to readme

commit 7487137227eb32ed9b12156338b865cb29b2dfd1
Author: Erik Scholz <[email protected]>
Date:   Thu Jun 22 14:20:47 2023 +0200

    rework convert.py to read hyper-parameters from config.json (#1958)

    * Read hyper-parameters from HuggingFace-transformer config.json, if they exist, and fall back to guessing, like before otherwise.
      This allows converting open_llama 3B and other non-standard model designs.

commit bbca06e26949686d61a5126332680ba3cccf235c
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 21 23:49:25 2023 +0200

    cmake: revert CUDA arch default to 52, 61 if f16 (#1959)

commit fb98254f99d769fcbbf20966ef386abdb48ef601
Author: Rahul Vivek Nair <[email protected]>
Date:   Thu Jun 22 03:18:43 2023 +0530

    Fix typo in README.md (#1961)

commit 049aa16b8c5c6d086246e4e6b9feb18de4fbd663
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 20 19:05:54 2023 +0300

    readme : add link to p1

commit 2322ec223a21625dfe9bd73ee677444a98a24ac9
Author: Xiake Sun <[email protected]>
Date:   Tue Jun 20 05:42:40 2023 -0700

    Fix typo (#1949)

commit aacdbd40562684665b6f7b8ba6695b7a2088bbb0
Author: Ettore Di Giacinto <[email protected]>
Date:   Tue Jun 20 03:24:39 2023 +0200

    llama : fix params struct slignment (#1936)

    * Workaround struct misalignment during value-copy

    Signed-off-by: mudler <[email protected]>

    * Move booleans at the bottom of the structure

    Signed-off-by: mudler <[email protected]>

    * Add comment

    Signed-off-by: mudler <[email protected]>

    ---------

    Signed-off-by: mudler <[email protected]>

commit 20568fe60f00155fa25e92eb3a7f6b911d557967
Author: Henri Vasserman <[email protected]>
Date:   Tue Jun 20 01:12:39 2023 +0300

    [Fix] Reenable server embedding endpoint (#1937)

    * Add back embedding feature

    * Update README

commit 18b35625c3c19c64b7818a12460ba5ddb006dfdc
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 20:43:30 2023 +0300

    ggml : fix bug in LBFGS optimizer (found by ggml tests)

commit ba4e85a8339b9dd7cdffad31838235f2fe45a8ea
Author: l3utterfly <[email protected]>
Date:   Mon Jun 19 23:20:06 2023 +0800

    llama : use aligned memory during ggml_init call from loading saved sessions (#1934)

    * fixed issue: memory is not guaranteed to be aligned properly during ggml_init call from loading saved sessions

    * - removed commented out old code from fix
    - updated another instance of same issue below original

commit 23fc5c219a9aebd57c8af3fac454062cc4622980
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 18:18:34 2023 +0300

    cmake : fix trailing whitespaces

commit cb40dfca694b5cb849837548fd69932117c78362
Author: Kawrakow <[email protected]>
Date:   Mon Jun 19 18:17:03 2023 +0300

    llama : only use Q6_K for output weights if tensor size is multiple of 256 (#1932)

    * Only use Q6_K for output weights if tensor size is multiple of 256

    * Fixed copy/paste mistake

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit ca7c3f4da5d144d4cd1dd44903552e6ba49b8ec8
Author: Kawrakow <[email protected]>
Date:   Mon Jun 19 18:14:09 2023 +0300

    cuda : faster k-quants on older GPUs (#1930)

    * k_quants: hopefully much faster Q4_K on older GPUs

    On the GTX-1660 that I have available to represent
    "old GPUs", token prediction drops from 65.5 ms/tok
    to 41.5 ms/tok!

    * k_quants: hopefully much faster Q3_K on older GPUs

    On the GTX-1660 that I have available to represent
    "old GPUs", token prediction drops from 60.3 ms/tok
    to 41.0 ms/tok!

    * k_quants: faster Q2_K on older GPUs

    It looks like I didn't need to change anything
    compared to what we already had, so this is just
    adding clarifying comments. But I now measure
    36.3 ms/tok on the GTX-1660, instead fo the
    47.2 ms/tok that I have written in the faster
    k-quants PR.

    * k_quants: faster Q5_K on older GPUs

    68.5 ms/tok -> 62.0 ms/tok on GTX-1660.
    For some reason the same access pattern that leads
    to such resounding success for Q2_K to Q4_K did not
    work at all for Q5_K.

    It is also more difficult to measure because for Q5_K_S
    we only have 32 layers on the GTX-1660, so output, tok embeddings
    and kv cache are done on the CPU.

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit b97ca431db35ec96a339a721acb1219c1dd78bed
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 18:12:33 2023 +0300

    ggml : sync latest ggml repo (#1924)

    * ggml : sync latest ggml repo

    * ggml : remove unused comments

    * ggml : asserts

commit 1e3abfcef073e73c2b31e8570cb06c5cb2fd1f55
Author: Howard Su <[email protected]>
Date:   Mon Jun 19 23:10:37 2023 +0800

    cmake : fix build shared ggml when CUDA is enabled (#1929)

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 16b9cd193965769089881bb8ec012fccca7b37b6
Author: Johannes Gäßler <[email protected]>
Date:   Mon Jun 19 10:23:56 2023 +0200

    Convert vector to f16 for dequantize mul mat vec (#1913)

    * Convert vector to f16 for dmmv

    * compile option

    * Added compilation option description to README

    * Changed cmake CUDA_ARCHITECTURES from "OFF" to "native"

commit b24c3049d96557c24782e4d32feaae65f47277af
Author: Johannes Gäßler <[email protected]>
Date:   Sun Jun 18 17:41:26 2023 +0200

    Added tokens per second to info prints (#1928)

commit 0ede372a51fd8160688e01b587582666c14e94e5
Author: Johannes Gäßler <[email protected]>
Date:   Sun Jun 18 16:07:09 2023 +0200

    Fixed incorrectly applying RMS norm twice (#1925)

commit 8596af427722775f0df4a7c90b9af067ba90d4ef
Author: l3utterfly <[email protected]>
Date:   Sun Jun 18 19:19:16 2023 +0800

    ggml : fix bug in ggml_compute_forward_add_q_f32 (#1918)

commit e1886cf4fe0d0f31661dda52a4a9f34bd9b9009a
Author: Mike <[email protected]>
Date:   Sun Jun 18 16:28:26 2023 +0800

    readme : update Android build instructions (#1922)

    Add steps for using termux on android devices to prevent common errors.

commit 8ab8ba62eb27cc340be2edf3418e051b1d967416
Author: Kawrakow <[email protected]>
Date:   Sun Jun 18 11:13:43 2023 +0300

    llama : prevent usage of k-quants when tensor size is not a multiple of 256 (#1921)

    * Fix examples/metal

    * k-quants: prevent usage when tensor size is not divisible by 256

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 90cc59d6ab1363a5c69c60c4b94db647d3a54a18
Author: Kawrakow <[email protected]>
Date:   Sun Jun 18 10:52:10 2023 +0300

    examples : fix examples/metal (#1920)

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit ce2c7d72e2d06988b5ddec6811ab923254542077
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 18 09:09:47 2023 +0300

    metal : handle buffers larger than device's maxBufferLength (#1826)

    * metal : handle buffers larger than device's maxBufferLength

    * metal : print more verbose device info + handle errors

    * metal : fix prints for overlapping views

    * metal : minimize view overlap to try to utilize device memory better

commit 57cd69460f736031a3fc54af1e97c03f80128478
Author: Howard Su <[email protected]>
Date:   Sun Jun 18 12:29:47 2023 +0800

    cmake : add CUDA_ARCHITECTURES to new target ggml_static (#1917)

commit b2416493ab3ab21686d47c96669da6d6c6af08a4
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 20:55:03 2023 +0300

    make : do not print help for simple example

commit 4f9c43e3bd488b7561119785485e1155dba338d7
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 20:24:11 2023 +0300

    minor : warning fixes

commit 2c9380dd2f77e41149340f3ecb09764d793b16db
Author: Johannes Gäßler <[email protected]>
Date:   Sat Jun 17 19:15:02 2023 +0200

    Only one CUDA stream per device for async compute (#1898)

commit 051e1b0e6a6e3aee7d989b47760980e6fda5861c
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 19:30:22 2023 +0300

    llama : fix kv_cache `n` init (close #1903)

commit 86c7571864ff331f8cdb9e092f3abeb123729a56
Author: DaniAndTheWeb <[email protected]>
Date:   Sat Jun 17 18:17:22 2023 +0200

    make : update for latest Arch (#1701)

    With the upcoming change to the openblas package in arch the Makefile workaround is no longer needed.

commit 3d59ec5935ea1d33e9d51060a8dd737169b9b89b
Author: Howard Su <[email protected]>
Date:   Sat Jun 17 23:46:15 2023 +0800

    ggml : fix warnings under MSVC (#1908)

commit 0711a5f6dce7f04c2a791b14bc47f7d4cb545408
Author: Aaron Miller <[email protected]>
Date:   Sat Jun 17 07:37:49 2023 -0700

    metal : add norm, cpy f16->f16, alibi kernels (#1823)

commit fc45a81bc642b9ef33d9004f2b363d558438a6c9
Author: Faez Shakil <[email protected]>
Date:   Sat Jun 17 17:13:05 2023 +0500

    exposed modules so that they can be invoked by nix run github:ggerganov/llama.cpp#server etc (#1863)

commit 794db3e7b982fee37e3995db9c3a216a57ff65e3
Author: Randall Fitzgerald <[email protected]>
Date:   Sat Jun 17 07:53:04 2023 -0400

    Server Example Refactor and Improvements (#1570)

    A major rewrite for the server example.

    Note that if you have built something on the previous server API, it will probably be incompatible.
    Check out the examples for how a typical chat app could work.

    This took a lot of effort, there are 24 PR's closed in the submitter's repo alone, over 160 commits and a lot of comments and testing.

    Summary of the changes:

    - adds missing generation parameters: tfs_z, typical_p, repeat_last_n, repeat_penalty, presence_penalty, frequency_penalty, mirostat, penalize_nl, seed, ignore_eos
    - applies missing top k sampler
    - removes interactive mode/terminal-like behavior, removes exclude parameter
    - moves threads and batch size to server command-line parameters
    - adds LoRA loading and matches command line parameters with main example
    - fixes stopping on EOS token and with the specified token amount with n_predict
    - adds server timeouts, host, and port settings
    - adds expanded generation complete response; adds generation settings, stop reason, prompt truncated, model used, and final text
    - sets defaults for unspecified parameters between requests
    - removes /next-token endpoint and as_loop parameter, adds stream parameter and server-sent events for streaming
    - adds CORS headers to responses
    - adds request logging, exception printing and optional verbose logging
    - adds better stopping words handling when matching multiple tokens and while streaming, or when it finishes on a partial stop string
    - adds printing an error when it can't bind to the host/port specified
    - fixes multi-byte character handling and replaces invalid UTF-8 characters on responses
    - prints timing and build info on startup
    - adds logit bias to request parameters
    - removes embedding mode
    - updates documentation; adds streaming Node.js and Bash examples
    - fixes code formatting
    - sets server threads to 1 since the current global state doesn't work well with simultaneous requests
    - adds truncation of the input prompt and better context reset
    - removes token limit from the input prompt
    - significantly simplified the logic and removed a lot of variables

    ---------

    Co-authored-by: anon998 <[email protected]>
    Co-authored-by: Henri Vasserman <[email protected]>
    Co-authored-by: Felix Hellmann <[email protected]>
    Co-authored-by: Johannes Gäßler <[email protected]>
    Co-authored-by: Lesaun Harvey <[email protected]>

commit 5ddf7ea1fb42bac21026de2f77e0f9c069b92234
Author: Jiří Podivín <[email protected]>
Date:   Sat Jun 17 12:32:48 2023 +0200

    hooks : setting up flake8 and pre-commit hooks (#1681)

    Small, non-functional changes were made to non-compliant files.
    These include breaking up long lines, whitespace sanitation and
    unused import removal.

    Maximum line length in python files was set to a generous 125 chars,
    in order to minimize number of changes needed in scripts and general
    annoyance. The "txt" prompts directory is excluded from the checks
    as it may contain oddly formatted files and strings for a good reason.

    Signed-off-by: Jiri Podivin <[email protected]>

commit bac19927c302737465a1deb14ac0943a221863e8
Author: Gustavo Rocha Dias <[email protected]>
Date:   Sat Jun 17 06:01:06 2023 -0300

    readme :  alternative way to build for Android with CLBlast. (#1828)

commit b4c6f46f17b6e02f1cd55a81339e7e64f3aaa688
Author: Kerfuffle <[email protected]>
Date:   Sat Jun 17 01:49:42 2023 -0600

    Allow cmake to build ggml as a library (#1896)

    * Allow cmake to build ggml as a library

    * A ggml_static library will be created

    * When BUILD_SHARED_LIBS is enabled, ggml_shared will also be built

commit 92f20d9942c86daeb78637bdad7296a572f4da28
Author: David Yang <[email protected]>
Date:   Sat Jun 17 14:51:54 2023 +0800

    train : get raw text instead of page with html (#1905)

    We probably want to train using just the text of Shakespeare instead of the html of the page displaying his work.

commit d411968e990c37f51328849c96a743dd78f3c3dd
Author: 0cc4m <[email protected]>
Date:   Fri Jun 16 20:59:49 2023 +0200

    opencl : support k-quants (#1836)

    * Porting q2_k kernel to OpenCL

    * Set global and local sizes for kernel calls for dequantizing k-quants

    * Added q6_k kernel

    * Fix q4_k opencl struct order

    * Replace uchar with uint8_t

    * Finish dequant kernels

    * Added OpenCL DMMV kernels

    * Fix q2_k, improve code

    * Fix q3_k

    * Shorten switch statements

    * Improve code formatting

    ---------

    Co-authored-by: Concedo <[email protected]>

commit b41b4cad6f956b5f501db0711dd7007c32b5eee5
Author: SuperUserNameMan <[email protected]>
Date:   Fri Jun 16 20:58:09 2023 +0200

    examples : add "simple" (#1840)

    * Create `simple.cpp`

    * minimalist example `CMakeLists.txt`

    * Update Makefile for minimalist example

    * remove 273: Trailing whitespace

    * removed trailing white spaces simple.cpp

    * typo and comments simple.cpp

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 13fe9d2d84f30cab613c960bf66ac83916006694
Author: Zenix <[email protected]>
Date:   Sat Jun 17 03:53:04 2023 +0900

    cmake : add auto detection of BLAS_INCLUDE_DIRS (#1886)

commit ac3b8869538c7fbdb48ff141d78c4dea091789f0
Author: Johannes Gäßler <[email protected]>
Date:   Fri Jun 16 20:25:51 2023 +0200

    llama : fix embd when offloading non-repeating layers (#1891)

commit 5b9ccaf104cc1054d4f8f17bc8a4b8dc949e5527
Author: FrankHB <[email protected]>
Date:   Sat Jun 17 02:25:01 2023 +0800

    Fixed possible macro redefinition (#1892)

    MinGW libstdc++ may define `NOMINMAX` unconditionally. This fixes the case when it is already defined.

commit 9cbf50c041a525d781c7764f493a5443924e4e38
Author: Borislav Stanimirov <[email protected]>
Date:   Fri Jun 16 21:23:53 2023 +0300

    build : fix and ignore MSVC warnings (#1889)

commit 3d0112261042b356621e93db3fa4c6798a5d098f
Author: Kawrakow <[email protected]>
Date:   Fri Jun 16 20:08:44 2023 +0300

    CUDA : faster k-quant dot kernels (#1862)

    * cuda : faster k-quant dot kernels

    * Imrove Q2_K dot kernel on older GPUs

    We now have a K_QUANTS_PER_ITERATION macro, which should be
    set to 1 on older and to 2 on newer GPUs.
    With this, we preserve the performance of the original
    PR on RTX-4080, and are faster compared to master on
    GTX-1660.

    * Imrove Q6_K dot kernel on older GPUs

    Using the same K_QUANTS_PER_ITERATION macro as last commit,
    we preserve performance on RTX-4080 and speed up
    Q6_K on a GTX-1660.

    * Add LLAMA_CUDA_KQUANTS_ITER to CMakeLists.txt and Makefile

    Allowed values are 1 or 2. 2 gives the best performance on
    modern GPUs and is set as default. On older GPUs 1 may work
    better.

    * PR comments

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 602c748863e15270d80d74aa2c3bf86ab8139e07
Author: Borislav Stanimirov <[email protected]>
Date:   Fri Jun 16 09:58:11 2023 +0300

    gitignore : add several entries specific to Visual Studio (#1888)

commit a09f9195be39afb4b023b646c0a6ec8a86915174
Author: Johannes Gäßler <[email protected]>
Date:   Thu Jun 15 21:49:08 2023 +0200

    Fixed CUDA runtime version check (#1879)

commit bed92756172d4514b23aaf9744cf8e2dc892fc7b
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 15 21:56:50 2023 +0300

    cmake : remove whitespaces

commit c36e81da62ebfe09a768201cc44fa8d712dd00ed
Author: yangli2 <[email protected]>
Date:   Thu Jun 15 11:05:53 2023 -0700

    examples : add chat-vicuna.sh (#1854)

    Co-authored-by: Yang Li <[email protected]>

commit 3559433fecedf365e7aba2fe3d5f89d9abb817c1
Author: Igor Okulist <[email protected]>
Date:   Thu Jun 15 12:51:26 2023 -0500

    cmake : set include path for OpenBlas (#1830)

commit 69b34a0e80300bfb3e996983ac3ea075f5526675
Author: Frederik Vogel <[email protected]>
Date:   Fri Jun 16 02:47:04 2023 +0900

    swift : Package compile breaks due to ggml-metal.metal (#1831)

    * Ignore metal file in spm

    * Add ggml.h to spm public Headers

    ---------

    Co-authored-by: Vogel Frederik <[email protected]>

commit cf267d1c71a781700698f8518e903239c3bcc929
Author: daboe01 <[email protected]>
Date:   Thu Jun 15 19:42:48 2023 +0200

    make : add train-text-from-scratch (#1850)

    * make finetuning example accessible

    * fixed: targed was in wrong line

    * fixed: name of executable was wrong

    * fixed: naming of binary

    * fixed: model path was wrong

    * fixed clean target

    * Update examples/train-text-from-scratch/README.md

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 9dda13e5e1f70bdfc25fbc0f0378f27c8b67e983
Author: Srinivas Billa <[email protected]>
Date:   Thu Jun 15 18:36:38 2023 +0100

    readme : server compile flag (#1874)

    Explicitly include the server make instructions for C++ noobsl like me ;)

commit 37e257c48e350cf03c353c10d31e777f8d00123d
Author: sandyiscool <[email protected]>
Date:   Thu Jun 15 23:06:06 2023 +0530

    make : clean *.so files (#1857)

commit 64cc19b4fe3df03bc20e520aa111c30cff3a655e
Author: Howard Su <[email protected]>
Date:   Fri Jun 16 01:29:59 2023 +0800

    Fix the validation of main device (#1872)

commit 4bfcc855abdb2c9fcc3c5a84747974521909fa41
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 15 20:29:48 2023 +0300

    metal : parallel command buffer encoding (#1860)

    * metal : parallel command buffer encoding

    * metal : determine number of command buffers based on gf->n_threads

commit 6b8312e7979b852f6b6ac9d29cd51fda16c17948
Author: Johannes Gäßler <[email protected]>
Date:   Thu Jun 15 19:06:46 2023 +0200

    Better error when using both LoRA + GPU layers (#1861)

commit 254a7a7a5ff4c874ff8488f1f5cbdd7e9c89d682
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 14 19:47:19 2023 +0200

    CUDA full GPU acceleration, KV cache in VRAM (#1827)

    * Fixed CUDA RoPE

    * ggml_cuda_mul_mat_vec_p021

    * ggml_cuda_scale

    * ggml_cuda_diag_mask_inf

    * ggml_is_permuted

    * ggml_cuda_cpy

    * flatten rows for ggml_cuda_op

    * Added a --low-vram option

    * Fixed Windows performance

    * Fixed LLAMA_CUDA_DMMV_Y > 1 for WizardLM

commit 92549202659fc23ba9fec5e688227d0da9b06b40
Author: 0xspringtime <[email protected]>
Date:   Tue Jun 13 15:37:54 2023 -0400

    baby-llama : fix operator!= (#1821)

    * Update baby-llama.cpp

    Seems to be an error in the implementation of the operator!= function. It attempts to compare the this pointer (a llama_hparams_lora object) with the other pointer (a llama_hparams object) using memcmp. This can lead to incorrect results because the sizes of the objects being compared (sizeof(llama_hparams) and sizeof(llama_hparams_lora)) are different, should now be able to compare two llama_hparams_lora objects for inequality.

    * Update baby-llama.cpp

    * Update baby-llama.cpp

commit e32089b2c20b1b87b22912f4a8b93fe01647d5b9
Author: xaedes <[email protected]>
Date:   Tue Jun 13 21:04:40 2023 +0200

    train : improved training-from-scratch example (#1652)

    * add python wrapper

    https://gist.github.com/abetlen/2b90e5f153f6efd00931d098de5c73ce

    * fix decoding error. adds errors=ignore parameter

    * add python bindings for functions to get and set the whole llama state
    (rng, logits, embedding and kv_cache)

    * update python bindings

    * add text generating baby-llama from scratch example

    * fix race condition bug in ggml_compute_forward_diag_mask_f32

    * implement ggml_soft_max_back for more performant backward pass of soft_max

    avoids creating big intermediate matrices of size n_embd x n_embd for llama layers and n_vocab x n_vocab for cross entropy loss

    * improve softmax backward pass

    go from quadratic runtime to linear runtime by simplifying the formulas

    * fix race condition bug in non-inplace ggml_compute_forward_diag_mask_f32

    memcpy needs to be synchronized across threads to avoid race conditions.
    => do it in INIT phase

    * fix bug in ggml_compute_forward_soft_max_back_f32 on DEBUG build

    * improve performance of mul_mat backward pass

    avoid transpose by using mul_mat with swapped arguments

    * avoid printing too much newlines in baby-llama-text

    * activate threading in baby-llama-text

    * add ggml_out_prod and use it for mul_mat backward pass for improved performance

    performance stats report improvement from 37 seconds to 16 seconds runtime during my training tests

    * better weight initialization improves training convergence at start

    * better weight initialization improves training convergence at start

    * improve ggml_out_prod performance

    - change iteration order (>15s -> 10s runtime)
    - parallelize over one more dimension: over dst matrix rows (10s -> <5s runtime)

    * add llama sampler, shuffle samples and constrain sampling to tokens occurring in train data

    * fix get_samples call, add model tensor names, increase model size, start training samples after newline

    * save train trained model to checkpoint and load model to be trained from checkpoint

    * use inplace functions where possible

    * initialize rng with srand

    * use different arguments for input and output checkpoint

    * ggml fixes to support backward pass on inplace operations

    * remove duplicate include

    * fix cross entropy loss

    - add target probabilities for each sample which is then used in cross entropy loss

    * print used memory before and after optimization

    * sample with non-greedy sampling parameters at the end of training

    * add cmake target for baby-llama-text

    * add ggml_add1_inplace to header

    * enable gradient propagation for inplace add1 and scale operations

    those functions backward passes don't need the original src0, so they also work when forward is inplace

    * implement AdamW in ggml_opt_adam by adding weight decay parameter (default 0.001f)

    also add a schedule parameter (default 1.0f) that can be used to scale alpha and decay according to learning schedule.
    setting the decay parameter to zero disables AdamW resulting in normal Adam optimizer.

    since the difference between Adam and AdamW is minimal it is not implemented as another optimizer, but integrated into the existing Adam optimizer.

    * use inplace operations in cross_entropy_loss

    * fix random weight initialization scale

    * add missing default parameters for adam optimizer

    * add ggml_opt_context, so that we can properly resume training

    otherwise the optimizer states, tracking statistics about the error function and its derivates,
    will reset to zero each time ggml_opt is called, hindering convergence on resumed training.

    now the optimizer context and all its memory is stored in a separate struct.

    * fix bug in llama_sample_token_mirostat_v2

    when all candidates are filtered out through mu threshold, the following soft_max operation will fail.
    so keep at least one.

    * add forward function without using cache, for more performant training

    during training on whole samples no cache is required.
    removing the cache and simplifying the remaining code results in performance and memory usage improvement.

    * print suppressed newline tokens as string "\n"

    printing too much actual newlines is suppressed to avoid flooding the console.

    * store optimizer state in training checkpoint and add learning schedule

    persistent optimizer state allows to resume training without resetting the optimizer
    learning schedule consists of linear warmup ramp followed by cosine decay with restarts

    * remove unused functions

    * fix bug in get_samples which corrupted training targets

    * save checkpoint only when it was trained

    * simplify code

    * remove trailing whitespace

    * simplify backward pass for SQRT

    * replace inefficient repeat backward pass with dedicated repeat_back operation

    * add ggml_cross_entropy_loss with backward pass for faster training

    cross entropy loss can also be implemented using softmax and log, but as dedicated operation it is faster and especially avoids unnecessary memory overhead.

    * add tests for cross_entropy_loss backward pass

    finite differences regularly results in estimated gradient of zero, despite the backward pass giving non zero gradient.
    _probably_ the finite differences fails due to numerical issues

    * use ggml_cross_entropy_loss in text training example

    * remove trailing whitespace

    * slightly improve how cross entropy loss is compute

    btw: directly implemented cross entropy loss seems to have way lower magnitudes than when implemented with softmax and log.
    probably the input to log gets closer to zero due to float numerics.
    maybe the multiplication by (1.0-eps)/sum is more accurate..

    * add llama_get_vocab to get the vocabulary as output parameters

    * set default model.type for unknown models with few layers

    * add export of training checkpoint to llama compatible model file

    * get vocabulary for exporting training checkpoint to llama compatible model file

    * implement backward pass of flash attention

    * bugfixes for backward pass of flash attention

    * test flash attention backward pass

    need to set loose error bounds to pass.
    the finitie differences are close to numeric limits and often return quite different values than the backward pass.
    reducing eps further lets the gradients vanish completely.
    likewise setting eps to big results in wronger values.
    the softmax in the middle of the function is probably the most responsible for the numeric issues using finite differences.

    * add option to train with flash attention and move options to the top of the main function

    training from scratch also works with flash attention
    training convergence and generation results after fix number of iterations are worse than when not using flash attention.
    maybe there still lingers a bug in the flash attention backward pass?
    but training works, just with slower convergence.

    flash attention is still worth to use, because it requires way less memory and is faster with high n_ctx

    * add train_params and command line option parser

    * remove unnecessary comments

    * add train params to specify memory size

    * remove python bindings

    * rename baby-llama-text to train-text-from-scratch

    * replace auto parameters in lambda function

    * add #include <climits>

    * add explicit cast to fix compile error

    "error: non-constant-expression cannot be narrowed from type 'int64_t' (aka 'long long') to 'uint32_t' (aka 'unsigned int') in initializer list [-Wc++11-narrowing]"

    * remove trailing whitespace

    * add ggml_opt_resume_g which accepts forward and backward cgraphs

    * fix formulas in comments

    * bug fix for ggml_compute_forward_get_rows_back_f32

    the result should be set to zero, not to whatever data is in opt0

    * improve training memory usage with scratch buffers

    instead of relying on the automatic backward pass, we manually create the graph for the backward pass.
    it turns out that all backward pass operations need only temporary memory which can be reused after each layer.

    will compute backward pass for ALL model parameters

    * add option to use scratch buffers in training or not

    make it configurable because currently training with scratch buffers implies flash attention and optimization over all parameters.

    * ci : disable temporary

    * store view offset and permute axes in opt[0] instead of storing it in padding

    use memcpy to store offset, because offset is of type size_t.
    when storing it as int32_t offset would have to be smaller than 2^31 which is not necessarily true.

    * minor : fix compile warnings + minor style changes

    * fix bug in threaded indices calculation of ggml_compute_forward_flash_attn_back_f32

    * store view offset like in master branch

    * bug fix in forward_batch_wo_cache_flash_attn_train

    * scratch buffer bug fixes in forward_batch_wo_cache_flash_attn_train

    data of permute and reshape is the same as their input.
    if we want to preserve the output of permute/reshape, we also need to preserve their inputs.

    replace reshape(src0, src1) with reshape_nd calls so that we don't need src1.

    replace (temporary) t03 with ggml_repeat(ctx0, layer.attention_norm, t02).
    in the future we could also use the new broadcasting ggml_mul to avoid these repeat calls.
    for this we need backward pass of broadcasting ggml_mul.

    * remove unnecessary scratch buffer 0

    buf 0 is persistent memory, so we can just disable scratch for this by using buf -1

    * avoid creating unnecessary grad tensors

    previously we need to create grads for model parameters, so that expand(..) correctly populates cgraph->leafs & cgraph->grads
    this wasted memory, because unnecessary grad for each op were automatically created:
    the automatically generated grad was unnecessary because we later manually set the grad (e.g. t35->grad = expand(gb, ...) ).
    this discarded the automatically generated grad resulting in wasted memory.

    improved this by changing expand(..) to not use ggml_build_forward_expand.
    expand set cgraph->nodes but not the leafs.
    cgraph->leafs & cgraph->grads are set in another pass after the last expand call.

    * print used training seed

    * zero initialize gfbuf and gbbuf

    * ci : re-enable workflows + add README for training

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 2347e45e7bdb09c9a7d74b2c0bc86c2b65f0c343
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 13 20:20:07 2023 +0300

    llama : do a warm-up eval at start for better timings (#1824)

commit 74d4cfa3438cb58bd177eed30014e6588694aaa8
Author: Kerfuffle <[email protected]>
Date:   Tue Jun 13 04:23:23 2023 -0600

    Allow "quantizing" to f16 and f32 (#1787)

    * Allow "quantizing" to f16 and f32

    Fix an issue where quantizing didn't respect LLAMA_NO_K_QUANTS

    Add brief help to the list of quantization types in the quantize tool

    Ignore case for quantization type arguments in the quantize tool

commit 74a6d922f12ccfe16b0c265f43be8978c6f25e98
Author: Kawrakow <[email protected]>
Date:   Mon Jun 12 22:39:21 2023 +0300

    Metal implementation for all k_quants (#1807)

    * metal : improve q4_K

    28.3 -> 26.0 ms/token by avoiding a branch in the
    calculation of the scales.

    * metal : small improvement for Q4_K

    * metal : still optimizing Q4_K

    This commit pushes it down to 25.3 ms / token.

    The crazy idea of using 6 bits for the scales is really costly on
    Metal: if I remove the bit fiddling necessary to make the block
    scales, time goes almost to the Q4_0 23 ms/token.

    Before pushing the k-quants upstream I had a Q4_K variant that
    had used 8-bit scales. It wasn't more accurate, used 0.125 bits more per weight,
    was running slightly slower on the CPU (due to the larger model size
    and being memory bound there), and the difference was entirely
    negligible under CUDA. So, I decided to publish the version with 6-bit
    scales. Perhaps I should re-consider and change to 8-bit scales?

    * metal : some more optimizations

    Q2_K: 25.4 ms/token
    Q6_K: 27.3 ms/token
    Q4_0: 22.8 ms/token
    Q4_1: 23.1 ms/token

    * metal : Q3_K support

    Something is not quite right yet.

    * metal : Q5_K support

    Initial version achieves 31.2 ms/token, 210 GB/s

    * metal : still not able to figure out why q3_K does not work

    * Minor

    * metal : yet another failed attempt to make q3_K work

    * metal : optimize Q5_K

    31.2 ms -> 27.8 ms.
    250 GB/s.

    * metal : q3_K still not working

    Adding a heavily commented q3_K metal kernel to explain
    my obviously faulty logic. Perhaps someone could spot the issue?

    * metal : q3_K finally working

    Not optimized at all.

    What was the issue? The scales are not 4-bytes aligned,
    and I was accessing them with a uint32_t pointer.
    When I tried that on CUDA, I got an error (illegal memory access)
    and added a memcpy to a local array of 3 uint32_t's.
    But on Metal it told me there is no memcpy, so I tried
    accessing directly. There is no error, just garbage results.
    At some point I did try accessing the scales with an uint16_t
    pointer (the scales are for sure 2-byte aligned), but was
    still getting garbage. I guess, there must have been another bug.

    No access to scales is via a uint16_t pointer and, after starting
    from scratch from the C dequantize function, it finally works.

    * metal : Q3_K 1st optimization pass

    * metal : Q3_K second optimization pass - 29.6 ms/token

    * metal : Q3_K cleanup

    * metal : fixed accidentally broken Q2_K

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit e4caa8da59c1c97dc23fa336f4d726984a20560f
Author: slaren <[email protected]>
Date:   Mon Jun 12 19:12:47 2023 +0200

    ci : run when changing only the CUDA sources (#1800)

commit 58970a4c39124a647ac2a640d9e178ea6c961e65
Author: Howard Su <[email protected]>
Date:   Mon Jun 12 20:44:16 2023 +0800

    Leverage mmap for offloading tensors to GPU (#1597)

    * Rebase to latest

    * Show progress

    * Add assert to make sure we only allocate temp buffer for non-CPU backend tensor

    Co-authored-by: Johannes Gäßler <[email protected]>

    ---------

    Co-authored-by: Johannes Gäßler <[email protected]>

commit 8c0a10e64dbf60fd9946c0cd5e6f59690800b123
Author: Kawrakow <[email protected]>
Date:   Mon Jun 12 14:31:36 2023 +0300

    metal : fix failure to load model (#1817)

    The number of buffers in the ggml context was left unitialized.
    This leads to sporadic failures to load the model on
    startup. It is actually strange that the failure occurred so
    infrequantly.

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit fa84c4b3e80199a5683438f062009c031a06c4fa
Author: Kerfuffle <[email protected]>
Date:   Sun Jun 11 08:19:17 2023 -0600

    Fix issue where interactive mode crashes when input exceeds ctx size (#1789)

    * Fix issue where interactive mode in the main example crashes when input exceeds ctx size

    * Ensure the context size is at least 8 tokens in the main example.

    Closes #1768

commit 12b063f0ecf280e98028e444fc492ee6222cdcdc
Author: Kyle Liang <[email protected]>
Date:   Sun Jun 11 21:20:52 2023 +0800

    Fixed WSL cuda's OOM error (#1594)

    * In the function , add the cuda error bypass.

    * remove excessive codes and prints

    ---------

    Co-authored-by: liang <[email protected]>

commit 31d2b5f4a4bae081e59b36ab37c6ff6f5b5940ad
Author: Ryan Landay <[email protected]>
Date:   Sun Jun 11 17:38:53 2023 +0800

    Update SHA256SUMS with current hashes for models quantized using q4_0 (#1798)

commit 4de0334f5cabf4696eced2e5d6e279fdfaa6c0f2
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 10 22:56:53 2023 +0300

    cmake : fix Metal build (close #1791)

commit 3f1223155a462477ac933474ebc4eab0ce3ca264
Author: Artyom Lebedev <[email protected]>
Date:   Sat Jun 10 22:51:36 2023 +0300

    k-quants : GCC12 compilation fix (#1792)

commit 303f5809f1b4ec49823dbe70cacd2124ec1d0df0
Author: Andrei <[email protected]>
Date:   Sat Jun 10 10:47:34 2023 -0400

    metal : fix issue with ggml-metal.metal path. Closes #1769 (#1782)

    * Fix issue with ggml-metal.metal path

    * Add ggml-metal.metal as a resource for llama target

    * Update flake.nix metal kernel substitution

commit 059e99066d95d73d1ca26c3375d47c0e35596229
Author: Aisuko <[email protected]>
Date:   Sun Jun 11 00:08:11 2023 +1000

    doc : fix wrong address of BLIS.md (#1772)

    Signed-off-by: Aisuko <[email protected]>

commit 17c10acfb44ecb7af25e37fb67b9501cbc0034d2
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 10 12:06:45 2023 +0300

    ggml : force no_alloc == false when creating opt tensors (close #1699)

    This is needed to make operators like ggml_view() be able to store their
    parameters in the ggml context's memory and not get discarded when
    no_alloc is true

commit e9b66ee9829039d4ab54550d6222e42a0b31e52a
Author: Kawrakow <[email protected]>
Date:   Sat Jun 10 11:28:11 2023 +0300

    metal : add Q4_1 implementation (#1785)

    23.3 ms / token, so just ~1% slower than q4_0.
    Achieves 290 GB/s memory throughput.

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 4f0154b0bad775ac4651bf73b5c216eb43c45cdc
Author: Kerfuffle <[email protected]>
Date:   Sat Jun 10 01:59:17 2023 -0600

    llama : support requantizing models instead of only allowing quantization from 16/32bit (#1691)

    * Add support for quantizing already quantized models

    * Threaded dequantizing and f16 to f32 conversion

    * Clean up thread blocks with spares calculation a bit

    * Use std::runtime_error exceptions.

commit ef3171d16241c18581d4d08374f0b9e396ade6b7
Author: Xingchen Song(宋星辰) <[email protected]>
Date:   Sat Jun 10 15:49:40 2023 +0800

    ggml : workaround for missing _mm256_setr_m128i in GCC < 8 (#1638)

commit 555275a693843273759230547001f9ae07fb537e
Author: rankaiyx <[email protected]>
Date:   Sat Jun 10 14:41:59 2023 +0800

    make : add SSSE3 compilation use case (#1659)

commit 98ed16557432d7a5179c57eddcc3a08a7ae6d54d
Author: Robert Sung-wook Shin <[email protected]>
Date:   Sat Jun 10 01:24:40 2023 +0900

    OpenCL: Add release memory (#1741)

    * Add opencl release memory

    * Rename function name

commit ae9663f1887513e152839e91f61c513075a19422
Author: Johannes Gäßler <[email protected]>
Date:   Fri Jun 9 13:58:15 2023 +0200

    Windows nvcc workaround (#1753)

    Fix gibberish output on Windows when using CUDA

commit b33dee282f5d8032b5f780152732dc45cbf2d349
Author: Georgi Gerganov <[email protected]>
Date:   Fri Jun 9 11:11:04 2023 +0300

    metal : fix build "tanhf" -> "tanh"

commit 92f44ff7f778ef1b94028b2ba6d39943b5ca0ada
Author: AT <[email protected]>
Date:   Fri Jun 9 04:00:51 2023 -0400

    metal : add GELU implementation (#1770)

    Co-authored-by: Adam Treat <[email protected]>

commit 245fc3c37da5ac5963f9f11a9f4f2ac08d96afc6
Author: Kawrakow <[email protected]>
Date:   Fri Jun 9 10:39:59 2023 +0300

    metal : faster q4_0 (#1775)

    * metal : 8% faster q4_0

    Avoid copying into local uchar4 anf float4.

    * metal : 17% faster Q4_0

    Use 64 threads in a thread group.

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 72ff5282bf0388c60821f504c4c8cc2b1f491aa6
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 22:28:21 2023 +0300

    metal : add Q2_K implementation (#1762)

    * metal : add Q2_K implementation

    27.1 ms / token on M2 Max 30-core GPU, so about the
    same speed as Q4_0. Memory throughput is ~156 GB/s.

    The access pattern used in the Q2_K
    CUDA implementation resulted in significantly lower
    performance (~31 ms/token).

    * Fixing merge conflicts

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 0bf7cf1b296fc9fca05411b37afdf08a531487d2
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 8 20:48:14 2023 +0300

    Revert "ggml : load data into int8x16x4_t using vld4q_s8 on arm64 (#1738)"

    This reverts commit 8432d4d9f716b25133e3ed671d91e21f6f3be867.

commit 8432d4d9f716b25133e3ed671d91e21f6f3be867
Author: le.chang <[email protected]>
Date:   Fri Jun 9 00:47:56 2023 +0800

    ggml : load data into int8x16x4_t using vld4q_s8 on arm64 (#1738)

commit 0f291e1f65c1d68201e71ce99c89562a36686b6d
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 19:46:22 2023 +0300

    metal : Q6_K implementation (#1752)

    * Metal implementation for Q4_K

    Very slow for now:
    42 ms / token, Q4_0 runs in 28 ms/token on my
    30-core M2 Max GPU.

    * Optimizing Q4_K on metal

    The first token always takes longer, I guess because
    the metal kernel is being jit-compiled.
    So, using n = 128 to measure time.

    At this point Q4_K takes 29.5 ms / token
    compared to 27.2 ms / token for Q4_0.
    Quite a bit better than the initial attempt,
    but still not good enough.

    * Optimizing q4_K metal dot some more

    For n = 256 it is now 28.1 ms/token compared to
    27 ms/token for q4_0.

    * Fix after merge with master

    * Metal implementation for Q6_K

    Similar to the CUDA implementation.
    No idea if this is the optimum for Metal, but the few
    alternative variants I tried all had a lower performance.

    We get 36.5 ms / token on M2 Max with 30 GPU cores.
    This corresponds to ~200 GB/second throughput.

    * clang-tidy : add config back

    * Much better Q6_K implementation for metal

    28.3 ms / token for 7B. Subtracting ~9 ms that is spent in
    other compute graph operations, we are left with ~19 ms
    for the matrix multiplications. The model is ~5.5 GB,
    so we are getting 1000 / 19 * 5.5 = 290 GB/s!

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 8fc8179919a11738910db07a800f2b176f8adf09
Author: qingfengfenga <[email protected]>
Date:   Thu Jun 8 15:58:53 2023 +0800

    Add llama.cpp docker support for non-latin languages (#1673)

    * Modify Dockerfile default character set to improve compatibility (#1673)

commit b50b570ed9d699d3d126d72fc02de92926bcd937
Author: Steven Roussey <[email protected]>
Date:   Thu Jun 8 00:12:28 2023 -0700

    ggml : fix fprintf warnings (#1720)

commit 53aba3f393f2e02a78ddaba2e934893a8bbf3246
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 8 10:09:08 2023 +0300

    clang-tidy : restore dot file from accidental deletion

commit 4161bdc04debb70bf5f275492b4d89fd9330087c
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 10:08:23 2023 +0300

    metal : add Q4_K implementation (#1733)

    * Metal implementation for Q4_K

    Very slow for now:
    42 ms / token, Q4_0 runs in 28 ms/token on my
    30-core M2 Max GPU.

    * Optimizing Q4_K on metal

    The first token always takes longer, I guess because
    the metal kernel is being jit-compiled.
    So, using n = 128 to measure time.

    At this point Q4_K takes 29.5 ms / token
    compared to 27.2 ms / token for Q4_0.
    Quite a bit better than the initial attempt,
    but still not good enough.

    * Optimizing q4_K metal dot some more

    For n = 256 it is now 28.1 ms/token compared to
    27 ms/token for q4_0.

    * Fix after merge with master

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 0035858273ebe0694926bf4414d279f3e1cd109d
Author: johnson442 <[email protected]>
Date:   Thu Jun 8 08:02:48 2023 +0100

    k-quants : add missing compile definition to CMakeLists (#1748)

commit 5c64a0952ee58b2d742ee84e8e3d43cce5d366db
Author: Georgi Gerganov <[email protected]>
Date:   Wed Jun 7 10:59:52 2023 +0300

    k-quants : allow to optionally disable at compile time (#1734)

    * k-quants : put behind optional compile flag LLAMA_K_QUANTS

    * build : enable k-quants by default

commit 5b57a5b72676540b6a45a3f527126299969ad241
Author: jacobi petrucciani <[email protected]>
Date:   Wed Jun 7 00:15:31 2023 -0400

    flake : update to support metal on m1/m2 (#1724)

commit 4dc62c545df0af60635d579e9e4dd91bc5afff51
Author: Georgi Gerganov <[email protected]>
Date:   Wed Jun 7 07:15:08 2023 +0300

    readme : add June roadmap

commit 35a84916fb029905c44746127026079268216e7a
Author: Willy Tarreau <[email protected]>
Date:   Wed Jun 7 04:10:17 2023 +0200

    main: add the possibility to open the prompt cache read-only (#1640)

    The prompt cache constitutes a nice speed up when using the same prompt
    prefix across multiple evaluations, but when using it, it will also be
    updated, which is not always desirable. One use case is to have a large
    prompt containing some context and usage rules, and a second part
    containing variable data of the problem being studied. In this case it's
    desirable to be able to save the first part once, and to always reuse it
    as-is without updating it with the second part.

    The new argument --prompt-cache-ro enables this read-only mode on the
    prompt cache. The prompt's contents that match the cache are loaded
    from the cache but the rest is not modified. This allowed to reduce a
    total analysis time from 112s to 49.7s here, without having to backup
    and restore a copy of the prompt, which takes significant time at 500
    MB.

    Signed-off-by: Willy Tarreau <[email protected]>

commit 2d7bf110edd8c49209401a16132052cba706ffd0
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 22:54:39 2023 +0300

    llama : fix vram_scratch var

commit 2a4e41a086ce80da68c402457c75c77e52dcc698
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 22:41:53 2023 +0300

    llama : fix compile warnings

commit 17366df842e358768c0df7024484fffecfc7865b
Author: Johannes Gäßler <[email protected]>
Date:   Tue Jun 6 21:33:23 2023 +0200

    Multi GPU support, CUDA refactor, CUDA scratch buffer (#1703)

    * CUDA multi GPU + scratch

    ggml_cuda_compute_forward

    Tensor parallelism

    ggml_cuda_add

    ggml_cuda_rms_norm

    ggml_cuda_silu

    CUDA scratch buffer

    --main-gpu CLI option

commit 44f906e8537fcec965e312d621c80556d6aa9bec
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 20:16:57 2023 +0300

    metal : add f16 support

commit d5b111f53d14972669eb52055f9df2567663ad8b
Author: LostRuins <[email protected]>
Date:   Wed Jun 7 01:00:01 2023 +0800

    Clblast fixes + enhancements to save VRAM and offload more layers (#1675)

    * Use events instead of clFinish, where possible

    * OpenCL: Don't load gpu layers into RAM, add mul_f32 kernel

    * Reduce queueing overhead for contiguous tensors by using single mul kernel call

    * Adapt to #1612 cl_mem malloc changes

    * Reduce code duplication between cuda and opencl branches

    * Improve implementation

    * Clblast fixes + enhancements to save VRAM:

    1. Change all Clblast buffers to CL_MEM_READ_WRITE, as the pool malloc currently doesn't properly handle them.
    2. When recycling buffers in pool malloc, always assign the SMALLEST available buffer that fits, instead of the FIRST available buffer
    3. When failing to recycle a buffer in pool malloc (all too small), instead recycle the largest available free buffer by resizing it.

    * change max value size_t to use limits

    * removed flags from the CL pool malloc, apply code tidying suggestions.

commit 2d43387dafe9c60f15f57aa23ee0b37864b98b32
Author: Georgi Gerganov <ggerga…
byroneverson added a commit to byroneverson/llm.cpp that referenced this pull request Jun 30, 2023
commit b8c8dda75fdf5fdea49c80af36818e7c30fe0ddf
Author: Howard Su <[email protected]>
Date:   Thu Jun 29 21:15:15 2023 +0800

    Use unsigned for random seed (#2006)

    * Use unsigned for random seed. Keep -1 as the value to use a time based seed.

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 96a712ca1b7f427e3bd7ffc0c70b2105cfc7fbf1
Author: LostRuins <[email protected]>
Date:   Thu Jun 29 11:56:43 2023 +0800

    Porting the improved K-Quant CUDA kernels to OpenCL (#1966)

    * Added broken new q4k quant

    * xx + ib0

    * Fix q2_k fast kernel

    * Use preprocessor for QK_K

    * Add q6_k fast matmul kernel

    * ported q3k speedup successfully

    * ported q2k and q5k speedups

    * remove old dot kernels and template

    * fixed global const struct types

    * fixing address spaces

    * fixed string too long CI issue

    ---------

    Co-authored-by: 0cc4m <[email protected]>

commit d3494bb86bf7ad5b0b60aae0220ea576f273b5c0
Author: m3ndax <[email protected]>
Date:   Wed Jun 28 20:39:08 2023 +0200

    llama : replacing auto &kv with const auto &kv (#2041)

    * Replacing auto &kv with const auto &kv

    * Create codacy.yml

    * Delete codacy.yml

commit 5b351e94d041742cd50ffcf2d44718d63bab398a
Author: Salvador E. Tropea <[email protected]>
Date:   Wed Jun 28 14:27:31 2023 -0300

    cuda : remove nchannels_x argument from mul_mat_vec_nc_f16_f32 (#2028)

    - Not used

commit 6432aabb6dc887436e4d57414b63116189c3b13b
Author: Salvador E. Tropea <[email protected]>
Date:   Wed Jun 28 14:26:26 2023 -0300

    cuda : fix missing const qualifier in casts (#2027)

commit b922bc351b69770cec2d35d2aa50fa052b95ca93
Author: Howard Su <[email protected]>
Date:   Wed Jun 28 10:13:02 2023 -0700

    llama : remove shards weight file support (#2000)

    * Remove multiple shards

    * Remove multiple file loaders

    * Remove llama_load_tensor_shard class

    * Simplify load logic

    * Remove dead code guess_n_parts function

    * Remove vocab_only from constructor of llama_model_loader

    * Remove alignment_prevents_mmap which is not more needed.

    * Remove useless check

commit 7f9753fa1263c4eded9a3de19778562f0e1093d7
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 28 18:35:54 2023 +0200

    CUDA GPU acceleration for LoRAs + f16 models (#1970)

commit cfa0750bc9dbc2d957a91b8ed09ab0035d8f3d4e
Author: ningshanwutuobang <[email protected]>
Date:   Wed Jun 28 23:53:37 2023 +0800

    llama : support input embeddings directly  (#1910)

    * add interface for float input

    * fixed inpL shape and type

    * add examples of input floats

    * add test example for embd input

    * fixed sampling

    * add free for context

    * fixed add end condition for generating

    * add examples for llava.py

    * add READMD for llava.py

    * add READMD for llava.py

    * add example of PandaGPT

    * refactor the interface and fixed the styles

    * add cmake build for embd-input

    * add cmake build for embd-input

    * Add MiniGPT-4 example

    * change the order of the args of llama_eval_internal

    * fix ci error

commit 9d23589d638dc74577d5ff880e6d4248b795f12e
Author: Erik Scholz <[email protected]>
Date:   Tue Jun 27 19:06:33 2023 +0200

    fix pthreads setaffinity usage on android (#2020)

commit 0be54f75a6c3e9a09ea71bdfcdabf9a996a0549b
Author: Howard Su <[email protected]>
Date:   Tue Jun 27 13:07:13 2023 +0800

    baby-llama : fix build after ggml_rope change (#2016)

commit 181e8d975528a4e27eabb8ae6e9865f9ceae4b37
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 27 00:37:13 2023 +0300

    llama : fix rope usage after ChatGLM change

commit d9779021bd59ed96daae75e820a5ac5da47ca8ff
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 27 00:06:51 2023 +0300

    ggml : add support for ChatGLM RoPE

commit d38e45157862b58a1824387e64860d68ca3533a7
Author: Roman Parykin <[email protected]>
Date:   Mon Jun 26 22:47:59 2023 +0300

    readme : add Scala 3 bindings repo (#2010)

commit eaa6ca5a61b8c9501df9ebe3d264f45b75a5f8aa
Author: David Yang <[email protected]>
Date:   Tue Jun 27 03:45:32 2023 +0800

    ggml : increase max tensor name + clean up compiler warnings in train-text (#1988)

    * Clean up compiler warnings in train-text

    Some brackets to disambiguate order of operations

    * Increase GGML_MAX_NAME

    Avoiding strncpy danger in train-text-from-scratch and reducing potential future name length issues

commit aa777abbb73655c4e1e9237b7c0ad66745e8e48c
Author: Gustavo Rocha Dias <[email protected]>
Date:   Mon Jun 26 16:34:45 2023 -0300

    readme : LD_LIBRARY_PATH complement for some Android devices when building with CLBlast inside Termux (#2007)

    * docs - Alternative way to build at Android, with CLBlast.

    * doc - LD_LIBRARY_PATH complement for some Android devices when building with CLBlast inside Termux.

    * doc- fix typo

commit c824d2e368d193d9f564ff29880a51cda9f90527
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 21:03:59 2023 +0300

    ggml : avoid conv 2d kernel round up

commit b853d456018b10820686362af41b2f2f75f1eec6
Author: zrm <[email protected]>
Date:   Mon Jun 26 13:57:59 2023 -0400

    ggml : add NUMA support (#1556)

    * detect NUMA systems and pin work threads to nodes (linux)

    * disable mmap prefetch/readahead for NUMA systems

    * avoid sending finalize op to thread pool if it does nothing

    * silence robot

    * fix args

    * make --numa a param

    * recommendation that n_nodes evenly divide n_threads did not warrant such aggressive enforcement

    * lower synchronization overhead

    * statically allocate

    * move numa state to g_state

    * add description for --numa

    * ggml : minor style changes

    * ggml : minor style + try fix sanitizer build

    * llama : allow to initialize backend with NUMA support

    * llama : avoid ggml include in llama-util.h

    * ggml : style / formatting

    * ggml : fix handling of ops with n_threads > n_tasks > 1

    * server : utilize numa parameter

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 9225baef71407d799a6f7f563b77fd7f82791416
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 20:10:52 2023 +0300

    k-quants : fix indentation

commit a84ab1da8dc6a59a5b67420ae1322f09503ffc72
Author: katsu560 <[email protected]>
Date:   Tue Jun 27 01:47:02 2023 +0900

    tests : fix quantize perf (#1990)

    * fix test quantize perf

    * avoid the global state

commit 5743ca80928d8410754ec64a5673d5c2dd6cfbb7
Author: katsu560 <[email protected]>
Date:   Tue Jun 27 01:46:07 2023 +0900

    k-quants : add AVX support to dot functions (#1916)

    * k_quants : add AVX support

    * k_quants : apply review comments

commit 412c60e4739367144e51e59add5dc7749d084115
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 26 19:45:09 2023 +0300

    readme : add link to new k-quants for visibility

commit 6769e944c727c63612dcafbef52009d21ae00fff
Author: Kawrakow <[email protected]>
Date:   Mon Jun 26 19:43:07 2023 +0300

    k-quants : support for super-block size of 64 (#2001)

    * k_quants: WIP super-blocks with 64 weights

    * k_quants: WIP super-blocks with 64 weights

    Q6_K scalar and AVX2 works

    * k_quants: WIP super-blocks with 64 weights

    Q4_K scalar and AVX2 works

    * k_quants: WIP super-blocks with 64 weights

    Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
    than the scalar implementation)

    * k_quants: WIP super-blocks with 64 weights

    Q3_K scalar and AVX2 works.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K scalar and AVX2 works, and with that all
    k_quants are done on AVX2 and scalar

    * k_quants: WIP super-blocks with 64 weights

    Q6_K working on CUDA. Cannot make it run quite as gast as
    with super-blocks with 256 weigths: 8% slower on 4080,
    20% slower on the 1660 (but there we fit 1 less layer on the
    GPU because pf the larger model size), so some fraction of
    these 20% is due to that,

    * k_quants: WIP super-blocks with 64 weights

    Q4_K working on CUDA. ~10% slower on GTX-1660,
    16% slower on 4080.

    * k_quants: WIP super-blocks with 64 weights

    Q2_K working on CUDA. ~3% slower on GTX-1660,
    10% slower on 4080.

    * k_quants: WIP super-blocks with 64 weights

    Q3_K working on CUDA.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K working on CUDA, and with this CUDA is done.

    * k_quants: WIP super-blocks with 64 weights

    Q6_K working on ARM_NEON

    * k_quants: WIP super-blocks with 64 weights

    Q4_K working on ARM_NEON, but quite a bit slower than 256 weights

    * k_quants: WIP super-blocks with 64 weights

    Q2_K working on ARM_NEON, but quite a bit slower than 256 weights

    * k_quants: WIP super-blocks with 64 weights

    Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.

    * k_quants: WIP super-blocks with 64 weights

    Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.

    With that, we have full support for ARM_NEON, although
    performance is not quite there.

    * k_quants: WIP super-blocks with 64 weights

    Slightly more efficient Q3_K and Q5_K

    * k_quants: WIP super-blocks with 64 weights

    Another small improvement for Q3_K and Q5_K on ARM_NEON

    * k_quants: WIP super-blocks with 64 weights

    Yet another speedup for Q5_K on ARM_NEON.
    We are now within 10% of the QK_K = 256 version.

    * k_quants: WIP super-blocks with 64 weights

    * We are able to pass preprocessor macros to the Metal
      compiler
    * Q6_K works and is actually slightly more efficient than
      the QK_K = 256 version (25.2 ms vs 25.8 ms)

    * k_quants: WIP super-blocks with 64 weights

    Q4_K works on Metal and is actually slightly faster
    than QK_K = 256 (21.95 ms vs 24.0 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q2_K works on Metal and is very slightly faster
    than QK_K = 256 (23.8 ms vs 24.2 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q3_K works on Metal and is slightly faster
    than QK_K = 256 (26.6 ms vs 28.3 ms).

    * k_quants: WIP super-blocks with 64 weights

    Q5_K works on Metal and is slightly faster
    than QK_K = 256 (23.7 ms vs 26.3 ms).

    * k_quants: call them _K, not _k, also on Metal

    * k_quants: correctly define QK_K in llama.cpp

    * Fixed bug in q4_K quantization added with the 64-block addition

    * Simplify via lambda

    * k_quants: swicth Q3_K to 4-bit scales when QK_K = 64

    Otherwise there isn't much benefit from this
    quantization type. There is some very slight loss
    in accuracy, but we reduce size by ~7%.
    E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
    8.6131 with 8-bit scales and 8.6352 with 4-bit,
    while file size decreases from 1.53G to 1.44G.

    * k_quants: switch Q4_K to 4-bit scales when QK_K = 64

     Here the loss in accuracy is greater than for Q3_K,
     but the Q4_K points still move further to the left on
     the perplexity vs size curve.

    * k_quants: forgot to add the Metal changes in last commit

    * k_quants: change Q5_K to be type 0 when QK_K = 64

    Still needs AVX2 implementation

    * k_quants: AVX2 implementation for new 64-weight Q5_K

    * k_quants: 10% faster ARM_NEON Q5_K dot product

    * k_quants: fixed issue caused by merging with master

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit cbebf61ca7584e9709265395f0127ae7fc0f1882
Author: Howard Su <[email protected]>
Date:   Mon Jun 26 23:15:47 2023 +0800

    Fix assert when free invalid cuda pointer (#2005)

    Fix assert via initializing extra structure always.
    CUDA error 1 at C:\GPT\llama.cpp\ggml-cuda.cu:2536: invalid argument

commit 447ccbe8c39332fcdd0d98a041b6e2ff6f06219d
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 16:08:12 2023 +0300

    readme : add new roadmap + manifesto

commit bd34cdde38f8fd661890ddd5f57ca30bf279877b
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 14:25:08 2023 +0300

    ggml : sync latest ggml (custom operators)

commit c2a08f87b8d180115d04b8688f383d1b2761b16d
Author: anon998 <[email protected]>
Date:   Sun Jun 25 08:48:36 2023 +0000

    fix server sampling: top k sampler first (#1977)

    Co-authored-by: anon <[email protected]>

commit 66a2555ba6cab954c56d653b29c27bfbbacfbfb1
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 25 09:07:03 2023 +0300

    readme : add Azure CI discussion link

commit e65ca7e14ac76c4046091da39d41a9017abaa9b3
Author: sjinzh <[email protected]>
Date:   Sun Jun 25 13:45:44 2023 +0800

    zig : upgrade build system support (#1981)

    * upgrade zig build system support

    * zig : add new line at the end of the file

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 5ec8dd5a3c6a9a109351d2257bb9d53869bd0a94
Author: Robyn <[email protected]>
Date:   Sun Jun 25 04:10:29 2023 +1000

    #1869 Fix null reference errors when training from scratch with CUDA (#1907)

    * #1869 Fix null reference errors when training from scratch with CUDA build

    Calling ggml_compute_forward when node->src0 was null was causing train-text-from-scratch.exe to terminate unexpectedly.

    * ggml : do not dereference src0 if NULL

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 65bdd52a867539691007f85c5508146d507f72c1
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 24 19:40:18 2023 +0300

    tests : sync test-grad0 from ggml

commit fdd18609113862dc6eb34dfc44a093d54c59ff1f
Author: Rowan Hart <[email protected]>
Date:   Sat Jun 24 04:07:08 2023 -0700

    flake : fix ggml-metal.metal path and run nixfmt (#1974)

commit c943d823c14cef33092205ca3944de6fdf7abf99
Author: AN Long <[email protected]>
Date:   Sat Jun 24 19:02:06 2023 +0800

    convert : fix invalid params in write_vocab_only (#1975)

commit f2c754e1c38936fdde74e4848ac468a696eb73c6
Author: slaren <[email protected]>
Date:   Sat Jun 24 12:57:18 2023 +0200

    ggml : improve ggml_graph_dump_dot, add ggml_format_name (#1978)

    * Improve ggml_graph_dump_dot, add ggml_format_name

    * add more automatic names to view ops

    * fix name of copies

commit 11da1a85cd69af84b5861134738c7e9e20907470
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 24 13:38:18 2023 +0300

    readme : fix whitespaces

commit 235b610d650cbfed6dbd5d671f750d35fc18cd7d
Author: Alberto <[email protected]>
Date:   Sat Jun 24 12:32:13 2023 +0200

    readme : fixed termux instructions (#1973)

commit b061ba9e2a7a2c335a200df8c11aed5e31e4ccbb
Author: Alex Renda <[email protected]>
Date:   Sat Jun 24 03:15:01 2023 -0700

    llama : fix top-p sampling to match the canonical definition (#1953)

    * Fix top-p sampling to match the standard definition (smallest set that has probability mass at least p, not largest set with probability mass less than p)

    * top-p: correct gt to gte

    * add test for correct top-p behavior

commit 527b6fba1d237befb324fd846bda7418c0fa394d
Author: Didzis Gosko <[email protected]>
Date:   Sat Jun 24 11:47:58 2023 +0300

    llama : make model stateless and context stateful (llama_state) (#1797)

    * llama : make model stateless and context stateful

    * llama : minor cleanup

    * llama : update internal API declaration

    * Apply suggestions from code review

    fix style

    Co-authored-by: Georgi Gerganov <[email protected]>

    * Missing model memory release

    * Fix style

    * Add deprecated warning for public API function llama_init_from_file

    * Update public API use cases: move away from deprecated llama_init_from_file

    * Deprecate public API function llama_apply_lora_from_file

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit d7b7484f74d486f77feb4c0b7af7e1718ed91651
Author: eiery <[email protected]>
Date:   Fri Jun 23 04:38:01 2023 -0400

    Add OpenLLaMA instructions to the README (#1954)

    * add openllama to readme

commit 7487137227eb32ed9b12156338b865cb29b2dfd1
Author: Erik Scholz <[email protected]>
Date:   Thu Jun 22 14:20:47 2023 +0200

    rework convert.py to read hyper-parameters from config.json (#1958)

    * Read hyper-parameters from HuggingFace-transformer config.json, if they exist, and fall back to guessing, like before otherwise.
      This allows converting open_llama 3B and other non-standard model designs.

commit bbca06e26949686d61a5126332680ba3cccf235c
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 21 23:49:25 2023 +0200

    cmake: revert CUDA arch default to 52, 61 if f16 (#1959)

commit fb98254f99d769fcbbf20966ef386abdb48ef601
Author: Rahul Vivek Nair <[email protected]>
Date:   Thu Jun 22 03:18:43 2023 +0530

    Fix typo in README.md (#1961)

commit 049aa16b8c5c6d086246e4e6b9feb18de4fbd663
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 20 19:05:54 2023 +0300

    readme : add link to p1

commit 2322ec223a21625dfe9bd73ee677444a98a24ac9
Author: Xiake Sun <[email protected]>
Date:   Tue Jun 20 05:42:40 2023 -0700

    Fix typo (#1949)

commit aacdbd40562684665b6f7b8ba6695b7a2088bbb0
Author: Ettore Di Giacinto <[email protected]>
Date:   Tue Jun 20 03:24:39 2023 +0200

    llama : fix params struct slignment (#1936)

    * Workaround struct misalignment during value-copy

    Signed-off-by: mudler <[email protected]>

    * Move booleans at the bottom of the structure

    Signed-off-by: mudler <[email protected]>

    * Add comment

    Signed-off-by: mudler <[email protected]>

    ---------

    Signed-off-by: mudler <[email protected]>

commit 20568fe60f00155fa25e92eb3a7f6b911d557967
Author: Henri Vasserman <[email protected]>
Date:   Tue Jun 20 01:12:39 2023 +0300

    [Fix] Reenable server embedding endpoint (#1937)

    * Add back embedding feature

    * Update README

commit 18b35625c3c19c64b7818a12460ba5ddb006dfdc
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 20:43:30 2023 +0300

    ggml : fix bug in LBFGS optimizer (found by ggml tests)

commit ba4e85a8339b9dd7cdffad31838235f2fe45a8ea
Author: l3utterfly <[email protected]>
Date:   Mon Jun 19 23:20:06 2023 +0800

    llama : use aligned memory during ggml_init call from loading saved sessions (#1934)

    * fixed issue: memory is not guaranteed to be aligned properly during ggml_init call from loading saved sessions

    * - removed commented out old code from fix
    - updated another instance of same issue below original

commit 23fc5c219a9aebd57c8af3fac454062cc4622980
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 18:18:34 2023 +0300

    cmake : fix trailing whitespaces

commit cb40dfca694b5cb849837548fd69932117c78362
Author: Kawrakow <[email protected]>
Date:   Mon Jun 19 18:17:03 2023 +0300

    llama : only use Q6_K for output weights if tensor size is multiple of 256 (#1932)

    * Only use Q6_K for output weights if tensor size is multiple of 256

    * Fixed copy/paste mistake

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit ca7c3f4da5d144d4cd1dd44903552e6ba49b8ec8
Author: Kawrakow <[email protected]>
Date:   Mon Jun 19 18:14:09 2023 +0300

    cuda : faster k-quants on older GPUs (#1930)

    * k_quants: hopefully much faster Q4_K on older GPUs

    On the GTX-1660 that I have available to represent
    "old GPUs", token prediction drops from 65.5 ms/tok
    to 41.5 ms/tok!

    * k_quants: hopefully much faster Q3_K on older GPUs

    On the GTX-1660 that I have available to represent
    "old GPUs", token prediction drops from 60.3 ms/tok
    to 41.0 ms/tok!

    * k_quants: faster Q2_K on older GPUs

    It looks like I didn't need to change anything
    compared to what we already had, so this is just
    adding clarifying comments. But I now measure
    36.3 ms/tok on the GTX-1660, instead fo the
    47.2 ms/tok that I have written in the faster
    k-quants PR.

    * k_quants: faster Q5_K on older GPUs

    68.5 ms/tok -> 62.0 ms/tok on GTX-1660.
    For some reason the same access pattern that leads
    to such resounding success for Q2_K to Q4_K did not
    work at all for Q5_K.

    It is also more difficult to measure because for Q5_K_S
    we only have 32 layers on the GTX-1660, so output, tok embeddings
    and kv cache are done on the CPU.

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit b97ca431db35ec96a339a721acb1219c1dd78bed
Author: Georgi Gerganov <[email protected]>
Date:   Mon Jun 19 18:12:33 2023 +0300

    ggml : sync latest ggml repo (#1924)

    * ggml : sync latest ggml repo

    * ggml : remove unused comments

    * ggml : asserts

commit 1e3abfcef073e73c2b31e8570cb06c5cb2fd1f55
Author: Howard Su <[email protected]>
Date:   Mon Jun 19 23:10:37 2023 +0800

    cmake : fix build shared ggml when CUDA is enabled (#1929)

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 16b9cd193965769089881bb8ec012fccca7b37b6
Author: Johannes Gäßler <[email protected]>
Date:   Mon Jun 19 10:23:56 2023 +0200

    Convert vector to f16 for dequantize mul mat vec (#1913)

    * Convert vector to f16 for dmmv

    * compile option

    * Added compilation option description to README

    * Changed cmake CUDA_ARCHITECTURES from "OFF" to "native"

commit b24c3049d96557c24782e4d32feaae65f47277af
Author: Johannes Gäßler <[email protected]>
Date:   Sun Jun 18 17:41:26 2023 +0200

    Added tokens per second to info prints (#1928)

commit 0ede372a51fd8160688e01b587582666c14e94e5
Author: Johannes Gäßler <[email protected]>
Date:   Sun Jun 18 16:07:09 2023 +0200

    Fixed incorrectly applying RMS norm twice (#1925)

commit 8596af427722775f0df4a7c90b9af067ba90d4ef
Author: l3utterfly <[email protected]>
Date:   Sun Jun 18 19:19:16 2023 +0800

    ggml : fix bug in ggml_compute_forward_add_q_f32 (#1918)

commit e1886cf4fe0d0f31661dda52a4a9f34bd9b9009a
Author: Mike <[email protected]>
Date:   Sun Jun 18 16:28:26 2023 +0800

    readme : update Android build instructions (#1922)

    Add steps for using termux on android devices to prevent common errors.

commit 8ab8ba62eb27cc340be2edf3418e051b1d967416
Author: Kawrakow <[email protected]>
Date:   Sun Jun 18 11:13:43 2023 +0300

    llama : prevent usage of k-quants when tensor size is not a multiple of 256 (#1921)

    * Fix examples/metal

    * k-quants: prevent usage when tensor size is not divisible by 256

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 90cc59d6ab1363a5c69c60c4b94db647d3a54a18
Author: Kawrakow <[email protected]>
Date:   Sun Jun 18 10:52:10 2023 +0300

    examples : fix examples/metal (#1920)

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit ce2c7d72e2d06988b5ddec6811ab923254542077
Author: Georgi Gerganov <[email protected]>
Date:   Sun Jun 18 09:09:47 2023 +0300

    metal : handle buffers larger than device's maxBufferLength (#1826)

    * metal : handle buffers larger than device's maxBufferLength

    * metal : print more verbose device info + handle errors

    * metal : fix prints for overlapping views

    * metal : minimize view overlap to try to utilize device memory better

commit 57cd69460f736031a3fc54af1e97c03f80128478
Author: Howard Su <[email protected]>
Date:   Sun Jun 18 12:29:47 2023 +0800

    cmake : add CUDA_ARCHITECTURES to new target ggml_static (#1917)

commit b2416493ab3ab21686d47c96669da6d6c6af08a4
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 20:55:03 2023 +0300

    make : do not print help for simple example

commit 4f9c43e3bd488b7561119785485e1155dba338d7
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 20:24:11 2023 +0300

    minor : warning fixes

commit 2c9380dd2f77e41149340f3ecb09764d793b16db
Author: Johannes Gäßler <[email protected]>
Date:   Sat Jun 17 19:15:02 2023 +0200

    Only one CUDA stream per device for async compute (#1898)

commit 051e1b0e6a6e3aee7d989b47760980e6fda5861c
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 17 19:30:22 2023 +0300

    llama : fix kv_cache `n` init (close #1903)

commit 86c7571864ff331f8cdb9e092f3abeb123729a56
Author: DaniAndTheWeb <[email protected]>
Date:   Sat Jun 17 18:17:22 2023 +0200

    make : update for latest Arch (#1701)

    With the upcoming change to the openblas package in arch the Makefile workaround is no longer needed.

commit 3d59ec5935ea1d33e9d51060a8dd737169b9b89b
Author: Howard Su <[email protected]>
Date:   Sat Jun 17 23:46:15 2023 +0800

    ggml : fix warnings under MSVC (#1908)

commit 0711a5f6dce7f04c2a791b14bc47f7d4cb545408
Author: Aaron Miller <[email protected]>
Date:   Sat Jun 17 07:37:49 2023 -0700

    metal : add norm, cpy f16->f16, alibi kernels (#1823)

commit fc45a81bc642b9ef33d9004f2b363d558438a6c9
Author: Faez Shakil <[email protected]>
Date:   Sat Jun 17 17:13:05 2023 +0500

    exposed modules so that they can be invoked by nix run github:ggerganov/llama.cpp#server etc (#1863)

commit 794db3e7b982fee37e3995db9c3a216a57ff65e3
Author: Randall Fitzgerald <[email protected]>
Date:   Sat Jun 17 07:53:04 2023 -0400

    Server Example Refactor and Improvements (#1570)

    A major rewrite for the server example.

    Note that if you have built something on the previous server API, it will probably be incompatible.
    Check out the examples for how a typical chat app could work.

    This took a lot of effort, there are 24 PR's closed in the submitter's repo alone, over 160 commits and a lot of comments and testing.

    Summary of the changes:

    - adds missing generation parameters: tfs_z, typical_p, repeat_last_n, repeat_penalty, presence_penalty, frequency_penalty, mirostat, penalize_nl, seed, ignore_eos
    - applies missing top k sampler
    - removes interactive mode/terminal-like behavior, removes exclude parameter
    - moves threads and batch size to server command-line parameters
    - adds LoRA loading and matches command line parameters with main example
    - fixes stopping on EOS token and with the specified token amount with n_predict
    - adds server timeouts, host, and port settings
    - adds expanded generation complete response; adds generation settings, stop reason, prompt truncated, model used, and final text
    - sets defaults for unspecified parameters between requests
    - removes /next-token endpoint and as_loop parameter, adds stream parameter and server-sent events for streaming
    - adds CORS headers to responses
    - adds request logging, exception printing and optional verbose logging
    - adds better stopping words handling when matching multiple tokens and while streaming, or when it finishes on a partial stop string
    - adds printing an error when it can't bind to the host/port specified
    - fixes multi-byte character handling and replaces invalid UTF-8 characters on responses
    - prints timing and build info on startup
    - adds logit bias to request parameters
    - removes embedding mode
    - updates documentation; adds streaming Node.js and Bash examples
    - fixes code formatting
    - sets server threads to 1 since the current global state doesn't work well with simultaneous requests
    - adds truncation of the input prompt and better context reset
    - removes token limit from the input prompt
    - significantly simplified the logic and removed a lot of variables

    ---------

    Co-authored-by: anon998 <[email protected]>
    Co-authored-by: Henri Vasserman <[email protected]>
    Co-authored-by: Felix Hellmann <[email protected]>
    Co-authored-by: Johannes Gäßler <[email protected]>
    Co-authored-by: Lesaun Harvey <[email protected]>

commit 5ddf7ea1fb42bac21026de2f77e0f9c069b92234
Author: Jiří Podivín <[email protected]>
Date:   Sat Jun 17 12:32:48 2023 +0200

    hooks : setting up flake8 and pre-commit hooks (#1681)

    Small, non-functional changes were made to non-compliant files.
    These include breaking up long lines, whitespace sanitation and
    unused import removal.

    Maximum line length in python files was set to a generous 125 chars,
    in order to minimize number of changes needed in scripts and general
    annoyance. The "txt" prompts directory is excluded from the checks
    as it may contain oddly formatted files and strings for a good reason.

    Signed-off-by: Jiri Podivin <[email protected]>

commit bac19927c302737465a1deb14ac0943a221863e8
Author: Gustavo Rocha Dias <[email protected]>
Date:   Sat Jun 17 06:01:06 2023 -0300

    readme :  alternative way to build for Android with CLBlast. (#1828)

commit b4c6f46f17b6e02f1cd55a81339e7e64f3aaa688
Author: Kerfuffle <[email protected]>
Date:   Sat Jun 17 01:49:42 2023 -0600

    Allow cmake to build ggml as a library (#1896)

    * Allow cmake to build ggml as a library

    * A ggml_static library will be created

    * When BUILD_SHARED_LIBS is enabled, ggml_shared will also be built

commit 92f20d9942c86daeb78637bdad7296a572f4da28
Author: David Yang <[email protected]>
Date:   Sat Jun 17 14:51:54 2023 +0800

    train : get raw text instead of page with html (#1905)

    We probably want to train using just the text of Shakespeare instead of the html of the page displaying his work.

commit d411968e990c37f51328849c96a743dd78f3c3dd
Author: 0cc4m <[email protected]>
Date:   Fri Jun 16 20:59:49 2023 +0200

    opencl : support k-quants (#1836)

    * Porting q2_k kernel to OpenCL

    * Set global and local sizes for kernel calls for dequantizing k-quants

    * Added q6_k kernel

    * Fix q4_k opencl struct order

    * Replace uchar with uint8_t

    * Finish dequant kernels

    * Added OpenCL DMMV kernels

    * Fix q2_k, improve code

    * Fix q3_k

    * Shorten switch statements

    * Improve code formatting

    ---------

    Co-authored-by: Concedo <[email protected]>

commit b41b4cad6f956b5f501db0711dd7007c32b5eee5
Author: SuperUserNameMan <[email protected]>
Date:   Fri Jun 16 20:58:09 2023 +0200

    examples : add "simple" (#1840)

    * Create `simple.cpp`

    * minimalist example `CMakeLists.txt`

    * Update Makefile for minimalist example

    * remove 273: Trailing whitespace

    * removed trailing white spaces simple.cpp

    * typo and comments simple.cpp

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 13fe9d2d84f30cab613c960bf66ac83916006694
Author: Zenix <[email protected]>
Date:   Sat Jun 17 03:53:04 2023 +0900

    cmake : add auto detection of BLAS_INCLUDE_DIRS (#1886)

commit ac3b8869538c7fbdb48ff141d78c4dea091789f0
Author: Johannes Gäßler <[email protected]>
Date:   Fri Jun 16 20:25:51 2023 +0200

    llama : fix embd when offloading non-repeating layers (#1891)

commit 5b9ccaf104cc1054d4f8f17bc8a4b8dc949e5527
Author: FrankHB <[email protected]>
Date:   Sat Jun 17 02:25:01 2023 +0800

    Fixed possible macro redefinition (#1892)

    MinGW libstdc++ may define `NOMINMAX` unconditionally. This fixes the case when it is already defined.

commit 9cbf50c041a525d781c7764f493a5443924e4e38
Author: Borislav Stanimirov <[email protected]>
Date:   Fri Jun 16 21:23:53 2023 +0300

    build : fix and ignore MSVC warnings (#1889)

commit 3d0112261042b356621e93db3fa4c6798a5d098f
Author: Kawrakow <[email protected]>
Date:   Fri Jun 16 20:08:44 2023 +0300

    CUDA : faster k-quant dot kernels (#1862)

    * cuda : faster k-quant dot kernels

    * Imrove Q2_K dot kernel on older GPUs

    We now have a K_QUANTS_PER_ITERATION macro, which should be
    set to 1 on older and to 2 on newer GPUs.
    With this, we preserve the performance of the original
    PR on RTX-4080, and are faster compared to master on
    GTX-1660.

    * Imrove Q6_K dot kernel on older GPUs

    Using the same K_QUANTS_PER_ITERATION macro as last commit,
    we preserve performance on RTX-4080 and speed up
    Q6_K on a GTX-1660.

    * Add LLAMA_CUDA_KQUANTS_ITER to CMakeLists.txt and Makefile

    Allowed values are 1 or 2. 2 gives the best performance on
    modern GPUs and is set as default. On older GPUs 1 may work
    better.

    * PR comments

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 602c748863e15270d80d74aa2c3bf86ab8139e07
Author: Borislav Stanimirov <[email protected]>
Date:   Fri Jun 16 09:58:11 2023 +0300

    gitignore : add several entries specific to Visual Studio (#1888)

commit a09f9195be39afb4b023b646c0a6ec8a86915174
Author: Johannes Gäßler <[email protected]>
Date:   Thu Jun 15 21:49:08 2023 +0200

    Fixed CUDA runtime version check (#1879)

commit bed92756172d4514b23aaf9744cf8e2dc892fc7b
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 15 21:56:50 2023 +0300

    cmake : remove whitespaces

commit c36e81da62ebfe09a768201cc44fa8d712dd00ed
Author: yangli2 <[email protected]>
Date:   Thu Jun 15 11:05:53 2023 -0700

    examples : add chat-vicuna.sh (#1854)

    Co-authored-by: Yang Li <[email protected]>

commit 3559433fecedf365e7aba2fe3d5f89d9abb817c1
Author: Igor Okulist <[email protected]>
Date:   Thu Jun 15 12:51:26 2023 -0500

    cmake : set include path for OpenBlas (#1830)

commit 69b34a0e80300bfb3e996983ac3ea075f5526675
Author: Frederik Vogel <[email protected]>
Date:   Fri Jun 16 02:47:04 2023 +0900

    swift : Package compile breaks due to ggml-metal.metal (#1831)

    * Ignore metal file in spm

    * Add ggml.h to spm public Headers

    ---------

    Co-authored-by: Vogel Frederik <[email protected]>

commit cf267d1c71a781700698f8518e903239c3bcc929
Author: daboe01 <[email protected]>
Date:   Thu Jun 15 19:42:48 2023 +0200

    make : add train-text-from-scratch (#1850)

    * make finetuning example accessible

    * fixed: targed was in wrong line

    * fixed: name of executable was wrong

    * fixed: naming of binary

    * fixed: model path was wrong

    * fixed clean target

    * Update examples/train-text-from-scratch/README.md

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 9dda13e5e1f70bdfc25fbc0f0378f27c8b67e983
Author: Srinivas Billa <[email protected]>
Date:   Thu Jun 15 18:36:38 2023 +0100

    readme : server compile flag (#1874)

    Explicitly include the server make instructions for C++ noobsl like me ;)

commit 37e257c48e350cf03c353c10d31e777f8d00123d
Author: sandyiscool <[email protected]>
Date:   Thu Jun 15 23:06:06 2023 +0530

    make : clean *.so files (#1857)

commit 64cc19b4fe3df03bc20e520aa111c30cff3a655e
Author: Howard Su <[email protected]>
Date:   Fri Jun 16 01:29:59 2023 +0800

    Fix the validation of main device (#1872)

commit 4bfcc855abdb2c9fcc3c5a84747974521909fa41
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 15 20:29:48 2023 +0300

    metal : parallel command buffer encoding (#1860)

    * metal : parallel command buffer encoding

    * metal : determine number of command buffers based on gf->n_threads

commit 6b8312e7979b852f6b6ac9d29cd51fda16c17948
Author: Johannes Gäßler <[email protected]>
Date:   Thu Jun 15 19:06:46 2023 +0200

    Better error when using both LoRA + GPU layers (#1861)

commit 254a7a7a5ff4c874ff8488f1f5cbdd7e9c89d682
Author: Johannes Gäßler <[email protected]>
Date:   Wed Jun 14 19:47:19 2023 +0200

    CUDA full GPU acceleration, KV cache in VRAM (#1827)

    * Fixed CUDA RoPE

    * ggml_cuda_mul_mat_vec_p021

    * ggml_cuda_scale

    * ggml_cuda_diag_mask_inf

    * ggml_is_permuted

    * ggml_cuda_cpy

    * flatten rows for ggml_cuda_op

    * Added a --low-vram option

    * Fixed Windows performance

    * Fixed LLAMA_CUDA_DMMV_Y > 1 for WizardLM

commit 92549202659fc23ba9fec5e688227d0da9b06b40
Author: 0xspringtime <[email protected]>
Date:   Tue Jun 13 15:37:54 2023 -0400

    baby-llama : fix operator!= (#1821)

    * Update baby-llama.cpp

    Seems to be an error in the implementation of the operator!= function. It attempts to compare the this pointer (a llama_hparams_lora object) with the other pointer (a llama_hparams object) using memcmp. This can lead to incorrect results because the sizes of the objects being compared (sizeof(llama_hparams) and sizeof(llama_hparams_lora)) are different, should now be able to compare two llama_hparams_lora objects for inequality.

    * Update baby-llama.cpp

    * Update baby-llama.cpp

commit e32089b2c20b1b87b22912f4a8b93fe01647d5b9
Author: xaedes <[email protected]>
Date:   Tue Jun 13 21:04:40 2023 +0200

    train : improved training-from-scratch example (#1652)

    * add python wrapper

    https://gist.github.com/abetlen/2b90e5f153f6efd00931d098de5c73ce

    * fix decoding error. adds errors=ignore parameter

    * add python bindings for functions to get and set the whole llama state
    (rng, logits, embedding and kv_cache)

    * update python bindings

    * add text generating baby-llama from scratch example

    * fix race condition bug in ggml_compute_forward_diag_mask_f32

    * implement ggml_soft_max_back for more performant backward pass of soft_max

    avoids creating big intermediate matrices of size n_embd x n_embd for llama layers and n_vocab x n_vocab for cross entropy loss

    * improve softmax backward pass

    go from quadratic runtime to linear runtime by simplifying the formulas

    * fix race condition bug in non-inplace ggml_compute_forward_diag_mask_f32

    memcpy needs to be synchronized across threads to avoid race conditions.
    => do it in INIT phase

    * fix bug in ggml_compute_forward_soft_max_back_f32 on DEBUG build

    * improve performance of mul_mat backward pass

    avoid transpose by using mul_mat with swapped arguments

    * avoid printing too much newlines in baby-llama-text

    * activate threading in baby-llama-text

    * add ggml_out_prod and use it for mul_mat backward pass for improved performance

    performance stats report improvement from 37 seconds to 16 seconds runtime during my training tests

    * better weight initialization improves training convergence at start

    * better weight initialization improves training convergence at start

    * improve ggml_out_prod performance

    - change iteration order (>15s -> 10s runtime)
    - parallelize over one more dimension: over dst matrix rows (10s -> <5s runtime)

    * add llama sampler, shuffle samples and constrain sampling to tokens occurring in train data

    * fix get_samples call, add model tensor names, increase model size, start training samples after newline

    * save train trained model to checkpoint and load model to be trained from checkpoint

    * use inplace functions where possible

    * initialize rng with srand

    * use different arguments for input and output checkpoint

    * ggml fixes to support backward pass on inplace operations

    * remove duplicate include

    * fix cross entropy loss

    - add target probabilities for each sample which is then used in cross entropy loss

    * print used memory before and after optimization

    * sample with non-greedy sampling parameters at the end of training

    * add cmake target for baby-llama-text

    * add ggml_add1_inplace to header

    * enable gradient propagation for inplace add1 and scale operations

    those functions backward passes don't need the original src0, so they also work when forward is inplace

    * implement AdamW in ggml_opt_adam by adding weight decay parameter (default 0.001f)

    also add a schedule parameter (default 1.0f) that can be used to scale alpha and decay according to learning schedule.
    setting the decay parameter to zero disables AdamW resulting in normal Adam optimizer.

    since the difference between Adam and AdamW is minimal it is not implemented as another optimizer, but integrated into the existing Adam optimizer.

    * use inplace operations in cross_entropy_loss

    * fix random weight initialization scale

    * add missing default parameters for adam optimizer

    * add ggml_opt_context, so that we can properly resume training

    otherwise the optimizer states, tracking statistics about the error function and its derivates,
    will reset to zero each time ggml_opt is called, hindering convergence on resumed training.

    now the optimizer context and all its memory is stored in a separate struct.

    * fix bug in llama_sample_token_mirostat_v2

    when all candidates are filtered out through mu threshold, the following soft_max operation will fail.
    so keep at least one.

    * add forward function without using cache, for more performant training

    during training on whole samples no cache is required.
    removing the cache and simplifying the remaining code results in performance and memory usage improvement.

    * print suppressed newline tokens as string "\n"

    printing too much actual newlines is suppressed to avoid flooding the console.

    * store optimizer state in training checkpoint and add learning schedule

    persistent optimizer state allows to resume training without resetting the optimizer
    learning schedule consists of linear warmup ramp followed by cosine decay with restarts

    * remove unused functions

    * fix bug in get_samples which corrupted training targets

    * save checkpoint only when it was trained

    * simplify code

    * remove trailing whitespace

    * simplify backward pass for SQRT

    * replace inefficient repeat backward pass with dedicated repeat_back operation

    * add ggml_cross_entropy_loss with backward pass for faster training

    cross entropy loss can also be implemented using softmax and log, but as dedicated operation it is faster and especially avoids unnecessary memory overhead.

    * add tests for cross_entropy_loss backward pass

    finite differences regularly results in estimated gradient of zero, despite the backward pass giving non zero gradient.
    _probably_ the finite differences fails due to numerical issues

    * use ggml_cross_entropy_loss in text training example

    * remove trailing whitespace

    * slightly improve how cross entropy loss is compute

    btw: directly implemented cross entropy loss seems to have way lower magnitudes than when implemented with softmax and log.
    probably the input to log gets closer to zero due to float numerics.
    maybe the multiplication by (1.0-eps)/sum is more accurate..

    * add llama_get_vocab to get the vocabulary as output parameters

    * set default model.type for unknown models with few layers

    * add export of training checkpoint to llama compatible model file

    * get vocabulary for exporting training checkpoint to llama compatible model file

    * implement backward pass of flash attention

    * bugfixes for backward pass of flash attention

    * test flash attention backward pass

    need to set loose error bounds to pass.
    the finitie differences are close to numeric limits and often return quite different values than the backward pass.
    reducing eps further lets the gradients vanish completely.
    likewise setting eps to big results in wronger values.
    the softmax in the middle of the function is probably the most responsible for the numeric issues using finite differences.

    * add option to train with flash attention and move options to the top of the main function

    training from scratch also works with flash attention
    training convergence and generation results after fix number of iterations are worse than when not using flash attention.
    maybe there still lingers a bug in the flash attention backward pass?
    but training works, just with slower convergence.

    flash attention is still worth to use, because it requires way less memory and is faster with high n_ctx

    * add train_params and command line option parser

    * remove unnecessary comments

    * add train params to specify memory size

    * remove python bindings

    * rename baby-llama-text to train-text-from-scratch

    * replace auto parameters in lambda function

    * add #include <climits>

    * add explicit cast to fix compile error

    "error: non-constant-expression cannot be narrowed from type 'int64_t' (aka 'long long') to 'uint32_t' (aka 'unsigned int') in initializer list [-Wc++11-narrowing]"

    * remove trailing whitespace

    * add ggml_opt_resume_g which accepts forward and backward cgraphs

    * fix formulas in comments

    * bug fix for ggml_compute_forward_get_rows_back_f32

    the result should be set to zero, not to whatever data is in opt0

    * improve training memory usage with scratch buffers

    instead of relying on the automatic backward pass, we manually create the graph for the backward pass.
    it turns out that all backward pass operations need only temporary memory which can be reused after each layer.

    will compute backward pass for ALL model parameters

    * add option to use scratch buffers in training or not

    make it configurable because currently training with scratch buffers implies flash attention and optimization over all parameters.

    * ci : disable temporary

    * store view offset and permute axes in opt[0] instead of storing it in padding

    use memcpy to store offset, because offset is of type size_t.
    when storing it as int32_t offset would have to be smaller than 2^31 which is not necessarily true.

    * minor : fix compile warnings + minor style changes

    * fix bug in threaded indices calculation of ggml_compute_forward_flash_attn_back_f32

    * store view offset like in master branch

    * bug fix in forward_batch_wo_cache_flash_attn_train

    * scratch buffer bug fixes in forward_batch_wo_cache_flash_attn_train

    data of permute and reshape is the same as their input.
    if we want to preserve the output of permute/reshape, we also need to preserve their inputs.

    replace reshape(src0, src1) with reshape_nd calls so that we don't need src1.

    replace (temporary) t03 with ggml_repeat(ctx0, layer.attention_norm, t02).
    in the future we could also use the new broadcasting ggml_mul to avoid these repeat calls.
    for this we need backward pass of broadcasting ggml_mul.

    * remove unnecessary scratch buffer 0

    buf 0 is persistent memory, so we can just disable scratch for this by using buf -1

    * avoid creating unnecessary grad tensors

    previously we need to create grads for model parameters, so that expand(..) correctly populates cgraph->leafs & cgraph->grads
    this wasted memory, because unnecessary grad for each op were automatically created:
    the automatically generated grad was unnecessary because we later manually set the grad (e.g. t35->grad = expand(gb, ...) ).
    this discarded the automatically generated grad resulting in wasted memory.

    improved this by changing expand(..) to not use ggml_build_forward_expand.
    expand set cgraph->nodes but not the leafs.
    cgraph->leafs & cgraph->grads are set in another pass after the last expand call.

    * print used training seed

    * zero initialize gfbuf and gbbuf

    * ci : re-enable workflows + add README for training

    ---------

    Co-authored-by: Georgi Gerganov <[email protected]>

commit 2347e45e7bdb09c9a7d74b2c0bc86c2b65f0c343
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 13 20:20:07 2023 +0300

    llama : do a warm-up eval at start for better timings (#1824)

commit 74d4cfa3438cb58bd177eed30014e6588694aaa8
Author: Kerfuffle <[email protected]>
Date:   Tue Jun 13 04:23:23 2023 -0600

    Allow "quantizing" to f16 and f32 (#1787)

    * Allow "quantizing" to f16 and f32

    Fix an issue where quantizing didn't respect LLAMA_NO_K_QUANTS

    Add brief help to the list of quantization types in the quantize tool

    Ignore case for quantization type arguments in the quantize tool

commit 74a6d922f12ccfe16b0c265f43be8978c6f25e98
Author: Kawrakow <[email protected]>
Date:   Mon Jun 12 22:39:21 2023 +0300

    Metal implementation for all k_quants (#1807)

    * metal : improve q4_K

    28.3 -> 26.0 ms/token by avoiding a branch in the
    calculation of the scales.

    * metal : small improvement for Q4_K

    * metal : still optimizing Q4_K

    This commit pushes it down to 25.3 ms / token.

    The crazy idea of using 6 bits for the scales is really costly on
    Metal: if I remove the bit fiddling necessary to make the block
    scales, time goes almost to the Q4_0 23 ms/token.

    Before pushing the k-quants upstream I had a Q4_K variant that
    had used 8-bit scales. It wasn't more accurate, used 0.125 bits more per weight,
    was running slightly slower on the CPU (due to the larger model size
    and being memory bound there), and the difference was entirely
    negligible under CUDA. So, I decided to publish the version with 6-bit
    scales. Perhaps I should re-consider and change to 8-bit scales?

    * metal : some more optimizations

    Q2_K: 25.4 ms/token
    Q6_K: 27.3 ms/token
    Q4_0: 22.8 ms/token
    Q4_1: 23.1 ms/token

    * metal : Q3_K support

    Something is not quite right yet.

    * metal : Q5_K support

    Initial version achieves 31.2 ms/token, 210 GB/s

    * metal : still not able to figure out why q3_K does not work

    * Minor

    * metal : yet another failed attempt to make q3_K work

    * metal : optimize Q5_K

    31.2 ms -> 27.8 ms.
    250 GB/s.

    * metal : q3_K still not working

    Adding a heavily commented q3_K metal kernel to explain
    my obviously faulty logic. Perhaps someone could spot the issue?

    * metal : q3_K finally working

    Not optimized at all.

    What was the issue? The scales are not 4-bytes aligned,
    and I was accessing them with a uint32_t pointer.
    When I tried that on CUDA, I got an error (illegal memory access)
    and added a memcpy to a local array of 3 uint32_t's.
    But on Metal it told me there is no memcpy, so I tried
    accessing directly. There is no error, just garbage results.
    At some point I did try accessing the scales with an uint16_t
    pointer (the scales are for sure 2-byte aligned), but was
    still getting garbage. I guess, there must have been another bug.

    No access to scales is via a uint16_t pointer and, after starting
    from scratch from the C dequantize function, it finally works.

    * metal : Q3_K 1st optimization pass

    * metal : Q3_K second optimization pass - 29.6 ms/token

    * metal : Q3_K cleanup

    * metal : fixed accidentally broken Q2_K

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit e4caa8da59c1c97dc23fa336f4d726984a20560f
Author: slaren <[email protected]>
Date:   Mon Jun 12 19:12:47 2023 +0200

    ci : run when changing only the CUDA sources (#1800)

commit 58970a4c39124a647ac2a640d9e178ea6c961e65
Author: Howard Su <[email protected]>
Date:   Mon Jun 12 20:44:16 2023 +0800

    Leverage mmap for offloading tensors to GPU (#1597)

    * Rebase to latest

    * Show progress

    * Add assert to make sure we only allocate temp buffer for non-CPU backend tensor

    Co-authored-by: Johannes Gäßler <[email protected]>

    ---------

    Co-authored-by: Johannes Gäßler <[email protected]>

commit 8c0a10e64dbf60fd9946c0cd5e6f59690800b123
Author: Kawrakow <[email protected]>
Date:   Mon Jun 12 14:31:36 2023 +0300

    metal : fix failure to load model (#1817)

    The number of buffers in the ggml context was left unitialized.
    This leads to sporadic failures to load the model on
    startup. It is actually strange that the failure occurred so
    infrequantly.

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit fa84c4b3e80199a5683438f062009c031a06c4fa
Author: Kerfuffle <[email protected]>
Date:   Sun Jun 11 08:19:17 2023 -0600

    Fix issue where interactive mode crashes when input exceeds ctx size (#1789)

    * Fix issue where interactive mode in the main example crashes when input exceeds ctx size

    * Ensure the context size is at least 8 tokens in the main example.

    Closes #1768

commit 12b063f0ecf280e98028e444fc492ee6222cdcdc
Author: Kyle Liang <[email protected]>
Date:   Sun Jun 11 21:20:52 2023 +0800

    Fixed WSL cuda's OOM error (#1594)

    * In the function , add the cuda error bypass.

    * remove excessive codes and prints

    ---------

    Co-authored-by: liang <[email protected]>

commit 31d2b5f4a4bae081e59b36ab37c6ff6f5b5940ad
Author: Ryan Landay <[email protected]>
Date:   Sun Jun 11 17:38:53 2023 +0800

    Update SHA256SUMS with current hashes for models quantized using q4_0 (#1798)

commit 4de0334f5cabf4696eced2e5d6e279fdfaa6c0f2
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 10 22:56:53 2023 +0300

    cmake : fix Metal build (close #1791)

commit 3f1223155a462477ac933474ebc4eab0ce3ca264
Author: Artyom Lebedev <[email protected]>
Date:   Sat Jun 10 22:51:36 2023 +0300

    k-quants : GCC12 compilation fix (#1792)

commit 303f5809f1b4ec49823dbe70cacd2124ec1d0df0
Author: Andrei <[email protected]>
Date:   Sat Jun 10 10:47:34 2023 -0400

    metal : fix issue with ggml-metal.metal path. Closes #1769 (#1782)

    * Fix issue with ggml-metal.metal path

    * Add ggml-metal.metal as a resource for llama target

    * Update flake.nix metal kernel substitution

commit 059e99066d95d73d1ca26c3375d47c0e35596229
Author: Aisuko <[email protected]>
Date:   Sun Jun 11 00:08:11 2023 +1000

    doc : fix wrong address of BLIS.md (#1772)

    Signed-off-by: Aisuko <[email protected]>

commit 17c10acfb44ecb7af25e37fb67b9501cbc0034d2
Author: Georgi Gerganov <[email protected]>
Date:   Sat Jun 10 12:06:45 2023 +0300

    ggml : force no_alloc == false when creating opt tensors (close #1699)

    This is needed to make operators like ggml_view() be able to store their
    parameters in the ggml context's memory and not get discarded when
    no_alloc is true

commit e9b66ee9829039d4ab54550d6222e42a0b31e52a
Author: Kawrakow <[email protected]>
Date:   Sat Jun 10 11:28:11 2023 +0300

    metal : add Q4_1 implementation (#1785)

    23.3 ms / token, so just ~1% slower than q4_0.
    Achieves 290 GB/s memory throughput.

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 4f0154b0bad775ac4651bf73b5c216eb43c45cdc
Author: Kerfuffle <[email protected]>
Date:   Sat Jun 10 01:59:17 2023 -0600

    llama : support requantizing models instead of only allowing quantization from 16/32bit (#1691)

    * Add support for quantizing already quantized models

    * Threaded dequantizing and f16 to f32 conversion

    * Clean up thread blocks with spares calculation a bit

    * Use std::runtime_error exceptions.

commit ef3171d16241c18581d4d08374f0b9e396ade6b7
Author: Xingchen Song(宋星辰) <[email protected]>
Date:   Sat Jun 10 15:49:40 2023 +0800

    ggml : workaround for missing _mm256_setr_m128i in GCC < 8 (#1638)

commit 555275a693843273759230547001f9ae07fb537e
Author: rankaiyx <[email protected]>
Date:   Sat Jun 10 14:41:59 2023 +0800

    make : add SSSE3 compilation use case (#1659)

commit 98ed16557432d7a5179c57eddcc3a08a7ae6d54d
Author: Robert Sung-wook Shin <[email protected]>
Date:   Sat Jun 10 01:24:40 2023 +0900

    OpenCL: Add release memory (#1741)

    * Add opencl release memory

    * Rename function name

commit ae9663f1887513e152839e91f61c513075a19422
Author: Johannes Gäßler <[email protected]>
Date:   Fri Jun 9 13:58:15 2023 +0200

    Windows nvcc workaround (#1753)

    Fix gibberish output on Windows when using CUDA

commit b33dee282f5d8032b5f780152732dc45cbf2d349
Author: Georgi Gerganov <[email protected]>
Date:   Fri Jun 9 11:11:04 2023 +0300

    metal : fix build "tanhf" -> "tanh"

commit 92f44ff7f778ef1b94028b2ba6d39943b5ca0ada
Author: AT <[email protected]>
Date:   Fri Jun 9 04:00:51 2023 -0400

    metal : add GELU implementation (#1770)

    Co-authored-by: Adam Treat <[email protected]>

commit 245fc3c37da5ac5963f9f11a9f4f2ac08d96afc6
Author: Kawrakow <[email protected]>
Date:   Fri Jun 9 10:39:59 2023 +0300

    metal : faster q4_0 (#1775)

    * metal : 8% faster q4_0

    Avoid copying into local uchar4 anf float4.

    * metal : 17% faster Q4_0

    Use 64 threads in a thread group.

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 72ff5282bf0388c60821f504c4c8cc2b1f491aa6
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 22:28:21 2023 +0300

    metal : add Q2_K implementation (#1762)

    * metal : add Q2_K implementation

    27.1 ms / token on M2 Max 30-core GPU, so about the
    same speed as Q4_0. Memory throughput is ~156 GB/s.

    The access pattern used in the Q2_K
    CUDA implementation resulted in significantly lower
    performance (~31 ms/token).

    * Fixing merge conflicts

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 0bf7cf1b296fc9fca05411b37afdf08a531487d2
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 8 20:48:14 2023 +0300

    Revert "ggml : load data into int8x16x4_t using vld4q_s8 on arm64 (#1738)"

    This reverts commit 8432d4d9f716b25133e3ed671d91e21f6f3be867.

commit 8432d4d9f716b25133e3ed671d91e21f6f3be867
Author: le.chang <[email protected]>
Date:   Fri Jun 9 00:47:56 2023 +0800

    ggml : load data into int8x16x4_t using vld4q_s8 on arm64 (#1738)

commit 0f291e1f65c1d68201e71ce99c89562a36686b6d
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 19:46:22 2023 +0300

    metal : Q6_K implementation (#1752)

    * Metal implementation for Q4_K

    Very slow for now:
    42 ms / token, Q4_0 runs in 28 ms/token on my
    30-core M2 Max GPU.

    * Optimizing Q4_K on metal

    The first token always takes longer, I guess because
    the metal kernel is being jit-compiled.
    So, using n = 128 to measure time.

    At this point Q4_K takes 29.5 ms / token
    compared to 27.2 ms / token for Q4_0.
    Quite a bit better than the initial attempt,
    but still not good enough.

    * Optimizing q4_K metal dot some more

    For n = 256 it is now 28.1 ms/token compared to
    27 ms/token for q4_0.

    * Fix after merge with master

    * Metal implementation for Q6_K

    Similar to the CUDA implementation.
    No idea if this is the optimum for Metal, but the few
    alternative variants I tried all had a lower performance.

    We get 36.5 ms / token on M2 Max with 30 GPU cores.
    This corresponds to ~200 GB/second throughput.

    * clang-tidy : add config back

    * Much better Q6_K implementation for metal

    28.3 ms / token for 7B. Subtracting ~9 ms that is spent in
    other compute graph operations, we are left with ~19 ms
    for the matrix multiplications. The model is ~5.5 GB,
    so we are getting 1000 / 19 * 5.5 = 290 GB/s!

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 8fc8179919a11738910db07a800f2b176f8adf09
Author: qingfengfenga <[email protected]>
Date:   Thu Jun 8 15:58:53 2023 +0800

    Add llama.cpp docker support for non-latin languages (#1673)

    * Modify Dockerfile default character set to improve compatibility (#1673)

commit b50b570ed9d699d3d126d72fc02de92926bcd937
Author: Steven Roussey <[email protected]>
Date:   Thu Jun 8 00:12:28 2023 -0700

    ggml : fix fprintf warnings (#1720)

commit 53aba3f393f2e02a78ddaba2e934893a8bbf3246
Author: Georgi Gerganov <[email protected]>
Date:   Thu Jun 8 10:09:08 2023 +0300

    clang-tidy : restore dot file from accidental deletion

commit 4161bdc04debb70bf5f275492b4d89fd9330087c
Author: Kawrakow <[email protected]>
Date:   Thu Jun 8 10:08:23 2023 +0300

    metal : add Q4_K implementation (#1733)

    * Metal implementation for Q4_K

    Very slow for now:
    42 ms / token, Q4_0 runs in 28 ms/token on my
    30-core M2 Max GPU.

    * Optimizing Q4_K on metal

    The first token always takes longer, I guess because
    the metal kernel is being jit-compiled.
    So, using n = 128 to measure time.

    At this point Q4_K takes 29.5 ms / token
    compared to 27.2 ms / token for Q4_0.
    Quite a bit better than the initial attempt,
    but still not good enough.

    * Optimizing q4_K metal dot some more

    For n = 256 it is now 28.1 ms/token compared to
    27 ms/token for q4_0.

    * Fix after merge with master

    ---------

    Co-authored-by: Iwan Kawrakow <[email protected]>

commit 0035858273ebe0694926bf4414d279f3e1cd109d
Author: johnson442 <[email protected]>
Date:   Thu Jun 8 08:02:48 2023 +0100

    k-quants : add missing compile definition to CMakeLists (#1748)

commit 5c64a0952ee58b2d742ee84e8e3d43cce5d366db
Author: Georgi Gerganov <[email protected]>
Date:   Wed Jun 7 10:59:52 2023 +0300

    k-quants : allow to optionally disable at compile time (#1734)

    * k-quants : put behind optional compile flag LLAMA_K_QUANTS

    * build : enable k-quants by default

commit 5b57a5b72676540b6a45a3f527126299969ad241
Author: jacobi petrucciani <[email protected]>
Date:   Wed Jun 7 00:15:31 2023 -0400

    flake : update to support metal on m1/m2 (#1724)

commit 4dc62c545df0af60635d579e9e4dd91bc5afff51
Author: Georgi Gerganov <[email protected]>
Date:   Wed Jun 7 07:15:08 2023 +0300

    readme : add June roadmap

commit 35a84916fb029905c44746127026079268216e7a
Author: Willy Tarreau <[email protected]>
Date:   Wed Jun 7 04:10:17 2023 +0200

    main: add the possibility to open the prompt cache read-only (#1640)

    The prompt cache constitutes a nice speed up when using the same prompt
    prefix across multiple evaluations, but when using it, it will also be
    updated, which is not always desirable. One use case is to have a large
    prompt containing some context and usage rules, and a second part
    containing variable data of the problem being studied. In this case it's
    desirable to be able to save the first part once, and to always reuse it
    as-is without updating it with the second part.

    The new argument --prompt-cache-ro enables this read-only mode on the
    prompt cache. The prompt's contents that match the cache are loaded
    from the cache but the rest is not modified. This allowed to reduce a
    total analysis time from 112s to 49.7s here, without having to backup
    and restore a copy of the prompt, which takes significant time at 500
    MB.

    Signed-off-by: Willy Tarreau <[email protected]>

commit 2d7bf110edd8c49209401a16132052cba706ffd0
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 22:54:39 2023 +0300

    llama : fix vram_scratch var

commit 2a4e41a086ce80da68c402457c75c77e52dcc698
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 22:41:53 2023 +0300

    llama : fix compile warnings

commit 17366df842e358768c0df7024484fffecfc7865b
Author: Johannes Gäßler <[email protected]>
Date:   Tue Jun 6 21:33:23 2023 +0200

    Multi GPU support, CUDA refactor, CUDA scratch buffer (#1703)

    * CUDA multi GPU + scratch

    ggml_cuda_compute_forward

    Tensor parallelism

    ggml_cuda_add

    ggml_cuda_rms_norm

    ggml_cuda_silu

    CUDA scratch buffer

    --main-gpu CLI option

commit 44f906e8537fcec965e312d621c80556d6aa9bec
Author: Georgi Gerganov <[email protected]>
Date:   Tue Jun 6 20:16:57 2023 +0300

    metal : add f16 support

commit d5b111f53d14972669eb52055f9df2567663ad8b
Author: LostRuins <[email protected]>
Date:   Wed Jun 7 01:00:01 2023 +0800

    Clblast fixes + enhancements to save VRAM and offload more layers (#1675)

    * Use events instead of clFinish, where possible

    * OpenCL: Don't load gpu layers into RAM, add mul_f32 kernel

    * Reduce queueing overhead for contiguous tensors by using single mul kernel call

    * Adapt to #1612 cl_mem malloc changes

    * Reduce code duplication between cuda and opencl branches

    * Improve implementation

    * Clblast fixes + enhancements to save VRAM:

    1. Change all Clblast buffers to CL_MEM_READ_WRITE, as the pool malloc currently doesn't properly handle them.
    2. When recycling buffers in pool malloc, always assign the SMALLEST available buffer that fits, instead of the FIRST available buffer
    3. When failing to recycle a buffer in pool malloc (all too small), instead recycle the largest available free buffer by resizing it.

    * change max value size_t to use limits

    * removed flags from the CL pool malloc, apply code tidying suggestions.

commit 2d43387dafe9c60f15f57aa23ee0b37864b98b32
Author: Georgi Gerganov <ggerga…
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

10 participants