2023-04-20 01:14:14 +00:00
# Define the default target now so that it is always the first target
2023-10-11 18:25:33 +00:00
BUILD_TARGETS = \
2024-01-12 05:59:57 +00:00
main quantize quantize-stats perplexity imatrix embedding vdot q8dot train-text-from-scratch convert-llama2c-to-ggml \
2023-11-06 21:36:23 +00:00
simple batched batched-bench save-load-state server gguf llama-bench libllava.a llava-cli baby-llama beam-search \
2024-03-07 17:55:27 +00:00
speculative infill tokenize benchmark-matmult parallel finetune export-lora lookahead lookup passkey gritlm tests/test-c.o
2023-05-27 17:04:14 +00:00
2023-07-21 10:09:16 +00:00
# Binaries only useful for tests
2023-10-11 18:25:33 +00:00
TEST_TARGETS = \
tests/test-llama-grammar tests/test-grammar-parser tests/test-double-float tests/test-grad0 tests/test-opt \
tests/test-quantize-fns tests/test-quantize-perf tests/test-sampling tests/test-tokenizer-0-llama \
2023-12-07 20:26:54 +00:00
tests/test-tokenizer-0-falcon tests/test-tokenizer-1-llama tests/test-tokenizer-1-bpe tests/test-rope \
2024-01-26 12:18:00 +00:00
tests/test-backend-ops tests/test-model-load-cancel tests/test-autorelease
2023-07-21 10:09:16 +00:00
2023-09-03 08:48:49 +00:00
# Code coverage output files
COV_TARGETS = *.gcno tests/*.gcno *.gcda tests/*.gcda *.gcov tests/*.gcov lcov-report gcovr-report
2023-09-04 19:26:24 +00:00
i f n d e f U N A M E _ S
UNAME_S := $( shell uname -s)
e n d i f
i f n d e f U N A M E _ P
UNAME_P := $( shell uname -p)
e n d i f
i f n d e f U N A M E _ M
UNAME_M := $( shell uname -m)
e n d i f
# Mac OS + Arm can report x86_64
# ref: https://github.com/ggerganov/whisper.cpp/issues/66#issuecomment-1282546789
i f e q ( $( UNAME_S ) , D a r w i n )
ifndef LLAMA_NO_METAL
LLAMA_METAL := 1
endif
ifneq ( $( UNAME_P) ,arm)
SYSCTL_M := $( shell sysctl -n hw.optional.arm64 2>/dev/null)
ifeq ( $( SYSCTL_M) ,1)
# UNAME_P := arm
# UNAME_M := arm64
warn := $( warning Your arch is announced as x86_64, but it seems to actually be ARM64. Not fixing that can lead to bad performance. For more info see: https://github.com/ggerganov/whisper.cpp/issues/66\# issuecomment-1282546789)
endif
endif
e n d i f
2023-05-27 17:04:14 +00:00
default : $( BUILD_TARGETS )
2023-04-20 01:14:14 +00:00
2023-09-07 14:15:01 +00:00
test : $( TEST_TARGETS )
@failures= 0; \
for test_target in $( TEST_TARGETS) ; do \
2023-08-30 09:42:51 +00:00
if [ " $$ test_target " = "tests/test-tokenizer-0-llama" ] ; then \
./$$ test_target $( CURDIR) /models/ggml-vocab-llama.gguf; \
elif [ " $$ test_target " = "tests/test-tokenizer-0-falcon" ] ; then \
2023-10-03 07:16:26 +00:00
./$$ test_target $( CURDIR) /models/ggml-vocab-falcon.gguf; \
2023-09-13 13:19:44 +00:00
elif [ " $$ test_target " = "tests/test-tokenizer-1-llama" ] ; then \
2023-08-30 09:42:51 +00:00
continue ; \
2023-10-03 07:16:26 +00:00
elif [ " $$ test_target " = "tests/test-tokenizer-1-bpe" ] ; then \
continue ; \
2023-08-30 09:42:51 +00:00
else \
2023-09-07 14:15:01 +00:00
echo " Running test $$ test_target... " ; \
2023-08-30 09:42:51 +00:00
./$$ test_target; \
fi ; \
2023-09-07 14:15:01 +00:00
if [ $$ ? -ne 0 ] ; then \
2023-12-21 20:07:46 +00:00
printf 'Test %s FAILED!\n\n' $$ test_target; \
2023-09-07 14:15:01 +00:00
failures = $$ ( ( failures + 1 ) ) ; \
else \
printf 'Test %s passed.\n\n' $$ test_target; \
fi ; \
done ; \
if [ $$ failures -gt 0 ] ; then \
printf '\n%s tests failed.\n' $$ failures; \
exit 1; \
fi
@echo 'All tests passed.'
2023-08-30 09:42:51 +00:00
all : $( BUILD_TARGETS ) $( TEST_TARGETS )
2023-09-03 08:48:49 +00:00
coverage : ## Run code coverage
gcov -pb tests/*.cpp
lcov-report : coverage ## Generate lcov report
mkdir -p lcov-report
lcov --capture --directory . --output-file lcov-report/coverage.info
genhtml lcov-report/coverage.info --output-directory lcov-report
gcovr-report : coverage ## Generate gcovr report
mkdir -p gcovr-report
gcovr --root . --html --html-details --output gcovr-report/coverage.html
2023-09-01 13:27:40 +00:00
i f d e f R I S C V _ C R O S S _ C O M P I L E
CC := riscv64-unknown-linux-gnu-gcc
CXX := riscv64-unknown-linux-gnu-g++
e n d i f
2023-03-10 18:40:58 +00:00
#
# Compile flags
#
2023-03-21 15:29:41 +00:00
# keep standard at C11 and C++11
2024-02-19 11:41:51 +00:00
MK_CPPFLAGS = -I. -Icommon
MK_CFLAGS = -std= c11 -fPIC
MK_CXXFLAGS = -std= c++11 -fPIC
MK_NVCCFLAGS = -std= c++11
2023-09-16 14:55:43 +00:00
ggml : add SOTA 2,3,4,5,6 bit k-quantizations (#1684)
* Starting to add k-quantization to ggml
I think it is better to have quantization separate from
ggml. For now just adding the k-quants there, but it would be
better to also factor out the existing ggml quantizations.
* Adding Q3_K and Q8_K (de)-quantization
* Q3_K now working on CUDA and AVX2/scalar
CUDA is not ideal - ~50% slower than Q4_0 for
single token prediction, about the same in batch
mode (perplexity). CPU single token is ~55 ms
(on Ryzen 7950X).
* Some improvement for Q3_K on CUDA
It is now ~22.5 ms/token on my GPU, so ~30% slower than Q4_0.
* Some more CUDA optimizations for Q3_K
Single token is now 20.5 ms/token (~20% slower than Q4_0).
Perplexity is on par with Q4_0.
* Adding Q4_K - scalar, AVX2, CUDA
Performance is the same or perhaps very slightly better than Q4_0 on the CPU.
On the GPU, single token prediction is ~10% better than Q4_0,
batch mode (perplexity is about the same).
* Adding Q6_K - scalar, AVX2, CUDA
Performance is ~40% lower compared to Q4_K on the CPU.
This is to be expected, considering that we are memory bound
on the CPU and the 6-bit model is ~44% larger than the 4-bit.
On the GPU, single token prediction is ~6% lower than Q4_0,
batch mode (perplexity) is even closer (but still slower).
* Adding Q5_K - scalar, AVX2, CUDA
Performance is ~20% lower compared to Q4_K on the CPU.
This is to be expected, considering that we are memory bound
on the CPU and the 5-bit model is ~22% larger than the 4-bit.
On the GPU, single token prediction is about the same as Q4_0
for both, single token and batch prediction.
* Per convention, all QX_K quantizations use Q5_K for output.weight
* Adding quantization mixes
* Quantization mixes: didn't quite get what I wanted in the last commit
* Q4_K dot product for ARM_NEON
* Q6_K dot product for ARM_NEON
* Q5_K dot product for ARM_NEON
* Adding Q3_K dot for ARM_NEON
It is 22% slower than Q4_K, despite the smaller model size.
On x86_64, where we are memory bound, the Q3_K model is
quite a bit faster than Q4_K.
* A very slightly faster ARM_NEON Q3_K dot
* Adding Q2_K - just CUDA for now
Token prediction is pretty good - about 15.5 ms on a RTX 4080.
Perplexity is about the same as Q4_K.
* Adding scalar and AVX2 Q2_K dot
* Adding ARM_NEON Q2_K dot
About the same performance as Q4_K.
* A slightly faster ARM_NEON Q2_K dot
Single token prediction is now ~36 ms on M2 Max.
The code is much simpler too.
* Fixed bug in Q2_K CUDA dot product kernel
Stranegly enough, for the few prompts I tried with the 7B model
the responses looked perfectly reasonable. Only realized something
is not quite right when I tried the larger models and started getting
nonse back.
In any case, Q2_K single token evaluation time on an RTX 4080 in a Ryzen7950X
box iusing CUDA and model fully loaded on the GPU are
~15.5 ms for 7B, ~25.4 ms for 13B, and ~55.8 ms for 30B.
The max number of layers that fit in VRAM for The 65B is 32.
With that, we get ~330 ms per token, which is not that much faster
than just running on the CPU (~470 ms per token).
* Don't print zeros/NaNs when no count histogram has been collected
* A 10% faster CUDA vector dot kernel for Q3_K
Q3_K is now running at ~18.5 ms / token on CUDA,
so the gap to Q4_0 is only 10%.
It seems memory acccess pattern is more important for
performance than the amount of computation the kernel
does.
* A slightly daster Q4_K AVX2 dot product
For perplexity, where we are less memory bound, time per
pass drops by ~5%. Barely measurable difference for single
token prediction.
* A slightly faster ARM_NEON A4_K dot product
* Minor
* Fix quantization error test
We cannot possibly be expecting rmse < 0.002 for 2- and 3-bit
quantization variants.
* Fix docker build
I have been sloppy with vector reinterpret casts on ARM_NEON.
It seems clang is very forgiving in that regard.
* Added forgotten ggml.o dependence on k_quants.h to the Makefile
* Had unintentionally committed the Makefile with -Ofast enabled
* ggml : rename k_quants -> ggml-quants-k, use lowercase in code
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-06-05 19:56:18 +00:00
# -Ofast tends to produce faster code, but may not be available for some compilers.
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 <iwan.kawrakow@gmail.com>
2023-06-26 16:43:07 +00:00
i f d e f L L A M A _ F A S T
2023-12-13 17:10:10 +00:00
MK_CFLAGS += -Ofast
HOST_CXXFLAGS += -Ofast
MK_NVCCFLAGS += -O3
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 <iwan.kawrakow@gmail.com>
2023-06-26 16:43:07 +00:00
e l s e
2023-12-13 17:10:10 +00:00
MK_CFLAGS += -O3
MK_CXXFLAGS += -O3
2024-02-03 19:14:59 +00:00
MK_NVCCFLAGS += -O3
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 <iwan.kawrakow@gmail.com>
2023-06-26 16:43:07 +00:00
e n d i f
2023-03-10 18:40:58 +00:00
2024-02-05 18:33:00 +00:00
i f n d e f L L A M A _ N O _ C C A C H E
CCACHE := $( shell which ccache)
i f d e f C C A C H E
export CCACHE_SLOPPINESS = time_macros
$( info I ccache found , compilation results will be cached . Disable with LLAMA_NO_CCACHE .)
CC := $( CCACHE) $( CC)
CXX := $( CCACHE) $( CXX)
e l s e
$( info I ccache not found . Consider installing it for faster compilation .)
e n d i f # CCACHE
e n d i f # LLAMA_NO_CCACHE
2023-09-08 12:09:21 +00:00
# clock_gettime came in POSIX.1b (1993)
# CLOCK_MONOTONIC came in POSIX.1-2001 / SUSv3 as optional
# posix_memalign came in POSIX.1-2001 / SUSv3
# M_PI is an XSI extension since POSIX.1-2001 / SUSv3, came in XPG1 (1985)
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D_XOPEN_SOURCE= 600
2023-09-08 12:09:21 +00:00
# Somehow in OpenBSD whenever POSIX conformance is specified
# some string functions rely on locale_t availability,
# which was introduced in POSIX.1-2008, forcing us to go higher
i f e q ( $( UNAME_S ) , O p e n B S D )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -U_XOPEN_SOURCE -D_XOPEN_SOURCE= 700
2023-09-08 12:09:21 +00:00
e n d i f
# Data types, macros and functions related to controlling CPU affinity and
# some memory allocation are available on Linux through GNU extensions in libc
i f e q ( $( UNAME_S ) , L i n u x )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D_GNU_SOURCE
2023-09-08 12:09:21 +00:00
e n d i f
# RLIMIT_MEMLOCK came in BSD, is not specified in POSIX.1,
# and on macOS its availability depends on enabling Darwin extensions
# similarly on DragonFly, enabling BSD extensions is necessary
i f e q ( $( UNAME_S ) , D a r w i n )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D_DARWIN_C_SOURCE
2023-09-08 12:09:21 +00:00
e n d i f
i f e q ( $( UNAME_S ) , D r a g o n F l y )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D__BSD_VISIBLE
2023-09-08 12:09:21 +00:00
e n d i f
# alloca is a non-standard interface that is not visible on BSDs when
# POSIX conformance is specified, but not all of them provide a clean way
# to enable it in such cases
i f e q ( $( UNAME_S ) , F r e e B S D )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D__BSD_VISIBLE
2023-09-08 12:09:21 +00:00
e n d i f
i f e q ( $( UNAME_S ) , N e t B S D )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D_NETBSD_SOURCE
2023-09-08 12:09:21 +00:00
e n d i f
i f e q ( $( UNAME_S ) , O p e n B S D )
2023-09-14 17:22:47 +00:00
MK_CPPFLAGS += -D_BSD_SOURCE
2023-09-08 12:09:21 +00:00
e n d i f
2023-05-28 19:01:02 +00:00
i f d e f L L A M A _ D E B U G
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -O0 -g
MK_CXXFLAGS += -O0 -g
MK_LDFLAGS += -g
2023-12-01 18:18:35 +00:00
ifeq ( $( UNAME_S) ,Linux)
2024-02-20 19:06:17 +00:00
MK_CPPFLAGS += -D_GLIBCXX_ASSERTIONS
2023-12-01 18:18:35 +00:00
endif
2023-05-28 19:01:02 +00:00
e l s e
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DNDEBUG
2023-04-29 15:43:28 +00:00
e n d i f
2023-10-11 19:35:46 +00:00
i f d e f L L A M A _ S A N I T I Z E _ T H R E A D
MK_CFLAGS += -fsanitize= thread -g
MK_CXXFLAGS += -fsanitize= thread -g
MK_LDFLAGS += -fsanitize= thread -g
e n d i f
i f d e f L L A M A _ S A N I T I Z E _ A D D R E S S
MK_CFLAGS += -fsanitize= address -fno-omit-frame-pointer -g
MK_CXXFLAGS += -fsanitize= address -fno-omit-frame-pointer -g
MK_LDFLAGS += -fsanitize= address -fno-omit-frame-pointer -g
e n d i f
i f d e f L L A M A _ S A N I T I Z E _ U N D E F I N E D
MK_CFLAGS += -fsanitize= undefined -g
MK_CXXFLAGS += -fsanitize= undefined -g
MK_LDFLAGS += -fsanitize= undefined -g
e n d i f
2023-07-04 12:38:04 +00:00
i f d e f L L A M A _ S E R V E R _ V E R B O S E
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DSERVER_VERBOSE= $( LLAMA_SERVER_VERBOSE)
2023-07-04 12:38:04 +00:00
e n d i f
2023-09-03 08:48:49 +00:00
i f d e f L L A M A _ C O D E _ C O V E R A G E
2023-09-05 19:12:00 +00:00
MK_CXXFLAGS += -fprofile-arcs -ftest-coverage -dumpbase ''
2023-09-03 08:48:49 +00:00
e n d i f
2023-09-01 09:07:06 +00:00
i f d e f L L A M A _ D I S A B L E _ L O G S
2023-09-05 19:12:00 +00:00
MK_CPPFLAGS += -DLOG_DISABLE_LOGS
2023-09-01 09:07:06 +00:00
e n d i f # LLAMA_DISABLE_LOGS
2023-03-28 16:48:20 +00:00
# warnings
2023-09-28 21:41:44 +00:00
WARN_FLAGS = -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function
MK_CFLAGS += $( WARN_FLAGS) -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror= implicit-int \
-Werror= implicit-function-declaration
MK_CXXFLAGS += $( WARN_FLAGS) -Wmissing-declarations -Wmissing-noreturn
2024-02-17 21:03:14 +00:00
i f e q ( $( LLAMA_FATAL_WARNINGS ) , 1 )
2024-02-19 12:45:41 +00:00
MK_CFLAGS += -Werror
2024-02-17 21:03:14 +00:00
MK_CXXFLAGS += -Werror
e n d i f
2023-11-14 17:34:41 +00:00
# this version of Apple ld64 is buggy
i f n e q '' '$(findstring dyld-1015.7,$(shell $(CC) $(LDFLAGS) -Wl,-v 2>&1))'
MK_CPPFLAGS += -DHAVE_BUGGY_APPLE_LINKER
e n d i f
2023-03-10 18:40:58 +00:00
# OS specific
# TODO: support Windows
2023-09-03 05:26:59 +00:00
i f n e q '' '$(filter $(UNAME_S),Linux Darwin FreeBSD NetBSD OpenBSD Haiku)'
MK_CFLAGS += -pthread
MK_CXXFLAGS += -pthread
2023-03-10 18:40:58 +00:00
e n d i f
2023-07-21 07:42:21 +00:00
# detect Windows
i f n e q ( $( findstring _NT ,$ ( UNAME_S ) ) , )
_WIN32 := 1
e n d i f
# library name prefix
i f n e q ( $( _WIN 32) , 1 )
LIB_PRE := lib
e n d i f
# Dynamic Shared Object extension
i f n e q ( $( _WIN 32) , 1 )
DSO_EXT := .so
e l s e
DSO_EXT := .dll
e n d i f
# Windows Sockets 2 (Winsock) for network-capable apps
i f e q ( $( _WIN 32) , 1 )
LWINSOCK2 := -lws2_32
e n d i f
2023-05-13 14:25:09 +00:00
i f d e f L L A M A _ G P R O F
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -pg
MK_CXXFLAGS += -pg
2023-05-13 14:25:09 +00:00
e n d i f
i f d e f L L A M A _ P E R F
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_PERF
2023-05-13 14:25:09 +00:00
e n d i f
2023-03-10 18:40:58 +00:00
# Architecture specific
# TODO: probably these flags need to be tweaked on some architectures
# feel free to update the Makefile for your architecture and send a pull request or issue
2023-09-01 13:27:40 +00:00
i f n d e f R I S C V
2023-07-21 10:53:27 +00:00
i f e q ( $( UNAME_M ) , $( filter $ ( UNAME_M ) ,x 86_ 64 i 686 amd 64) )
2023-04-02 07:17:05 +00:00
# Use all CPU extensions that are available:
2023-12-13 17:10:10 +00:00
MK_CFLAGS += -march= native -mtune= native
HOST_CXXFLAGS += -march= native -mtune= native
2023-04-22 08:08:12 +00:00
# Usage AVX-only
2023-09-03 05:26:59 +00:00
#MK_CFLAGS += -mfma -mf16c -mavx
#MK_CXXFLAGS += -mfma -mf16c -mavx
2023-06-10 06:41:59 +00:00
# Usage SSSE3-only (Not is SSE3!)
2023-09-03 05:26:59 +00:00
#MK_CFLAGS += -mssse3
#MK_CXXFLAGS += -mssse3
2023-03-10 18:40:58 +00:00
e n d i f
2023-06-04 20:34:30 +00:00
2023-09-01 13:53:14 +00:00
i f n e q '' '$(findstring mingw,$(shell $(CC) -dumpmachine))'
2023-12-12 09:27:26 +00:00
# The stack is only 16-byte aligned on Windows, so don't let gcc emit aligned moves.
# https://gcc.gnu.org/bugzilla/show_bug.cgi?id=54412
# https://github.com/ggerganov/llama.cpp/issues/2922
2023-09-05 19:12:00 +00:00
MK_CFLAGS += -Xassembler -muse-unaligned-vector-move
MK_CXXFLAGS += -Xassembler -muse-unaligned-vector-move
2023-12-12 09:27:26 +00:00
# Target Windows 8 for PrefetchVirtualMemory
MK_CPPFLAGS += -D_WIN32_WINNT= 0x602
2023-09-01 13:53:14 +00:00
e n d i f
2023-08-07 06:21:46 +00:00
i f n e q ( $( filter aarch 64%,$ ( UNAME_M ) ) , )
# Apple M1, M2, etc.
# Raspberry Pi 3, 4, Zero 2 (64-bit)
2023-12-22 15:11:12 +00:00
# Nvidia Jetson
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -mcpu= native
MK_CXXFLAGS += -mcpu= native
2023-12-22 15:11:12 +00:00
JETSON_RELEASE_INFO = $( shell jetson_release)
ifdef JETSON_RELEASE_INFO
ifneq ( $( filter TX2%,$( JETSON_RELEASE_INFO) ) ,)
JETSON_EOL_MODULE_DETECT = 1
CC = aarch64-unknown-linux-gnu-gcc
cxx = aarch64-unknown-linux-gnu-g++
endif
endif
2023-08-07 06:21:46 +00:00
e n d i f
i f n e q ( $( filter armv 6%,$ ( UNAME_M ) ) , )
# Raspberry Pi 1, Zero
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -mfpu= neon-fp-armv8 -mfp16-format= ieee -mno-unaligned-access
MK_CXXFLAGS += -mfpu= neon-fp-armv8 -mfp16-format= ieee -mno-unaligned-access
2023-08-07 06:21:46 +00:00
e n d i f
i f n e q ( $( filter armv 7%,$ ( UNAME_M ) ) , )
# Raspberry Pi 2
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -mfpu= neon-fp-armv8 -mfp16-format= ieee -mno-unaligned-access -funsafe-math-optimizations
MK_CXXFLAGS += -mfpu= neon-fp-armv8 -mfp16-format= ieee -mno-unaligned-access -funsafe-math-optimizations
2023-08-07 06:21:46 +00:00
e n d i f
i f n e q ( $( filter armv 8%,$ ( UNAME_M ) ) , )
# Raspberry Pi 3, 4, Zero 2 (32-bit)
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -mfp16-format= ieee -mno-unaligned-access
MK_CXXFLAGS += -mfp16-format= ieee -mno-unaligned-access
2023-08-07 06:21:46 +00:00
e n d i f
2023-03-10 18:40:58 +00:00
i f n e q ( $( filter ppc 64%,$ ( UNAME_M ) ) , )
POWER9_M := $( shell grep "POWER9" /proc/cpuinfo)
ifneq ( ,$( findstring POWER9,$( POWER9_M) ) )
2023-09-03 05:26:59 +00:00
MK_CFLAGS += -mcpu= power9
MK_CXXFLAGS += -mcpu= power9
2023-03-10 18:40:58 +00:00
endif
e n d i f
2023-06-04 20:34:30 +00:00
2023-11-17 16:11:23 +00:00
i f n e q ( $( filter ppc 64le %,$ ( UNAME_M ) ) , )
MK_CFLAGS += -mcpu= powerpc64le
MK_CXXFLAGS += -mcpu= powerpc64le
CUDA_POWER_ARCH = 1
e n d i f
2023-09-01 13:27:40 +00:00
e l s e
2023-09-05 19:12:00 +00:00
MK_CFLAGS += -march= rv64gcv -mabi= lp64d
MK_CXXFLAGS += -march= rv64gcv -mabi= lp64d
2023-09-01 13:27:40 +00:00
e n d i f
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 <iwan.kawrakow@gmail.com>
2023-06-26 16:43:07 +00:00
i f d e f L L A M A _ Q K K _ 6 4
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_QKK_64
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 <iwan.kawrakow@gmail.com>
2023-06-26 16:43:07 +00:00
e n d i f
2023-06-07 07:59:52 +00:00
2023-03-11 10:26:16 +00:00
i f n d e f L L A M A _ N O _ A C C E L E R A T E
2023-09-04 19:26:24 +00:00
# Mac OS - include Accelerate framework.
# `-framework Accelerate` works both with Apple Silicon and Mac Intel
2023-03-10 18:40:58 +00:00
ifeq ( $( UNAME_S) ,Darwin)
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_ACCELERATE
2023-09-27 15:34:32 +00:00
MK_CPPFLAGS += -DACCELERATE_NEW_LAPACK
MK_CPPFLAGS += -DACCELERATE_LAPACK_ILP64
2023-09-03 05:26:59 +00:00
MK_LDFLAGS += -framework Accelerate
2023-03-10 18:40:58 +00:00
endif
2023-06-04 20:34:30 +00:00
e n d i f # LLAMA_NO_ACCELERATE
2023-07-10 15:49:56 +00:00
i f d e f L L A M A _ M P I
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_MPI
MK_CFLAGS += -Wno-cast-qual
MK_CXXFLAGS += -Wno-cast-qual
2023-10-29 16:32:28 +00:00
OBJS += ggml-mpi.o
2023-07-10 15:49:56 +00:00
e n d i f # LLAMA_MPI
2023-03-11 10:26:16 +00:00
i f d e f L L A M A _ O P E N B L A S
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_OPENBLAS $( shell pkg-config --cflags-only-I openblas)
MK_CFLAGS += $( shell pkg-config --cflags-only-other openblas)
MK_LDFLAGS += $( shell pkg-config --libs openblas)
2023-06-04 20:34:30 +00:00
e n d i f # LLAMA_OPENBLAS
2023-05-20 14:58:31 +00:00
i f d e f L L A M A _ B L I S
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_OPENBLAS -I/usr/local/include/blis -I/usr/include/blis
MK_LDFLAGS += -lblis -L/usr/local/lib
2023-06-04 20:34:30 +00:00
e n d i f # LLAMA_BLIS
2023-04-19 09:22:45 +00:00
i f d e f L L A M A _ C U B L A S
2024-02-27 02:03:06 +00:00
ifneq ( '' , '$(wildcard /opt/cuda)' )
CUDA_PATH ?= /opt/cuda
else
CUDA_PATH ?= /usr/local/cuda
endif
MK_CPPFLAGS += -DGGML_USE_CUBLAS -I$( CUDA_PATH) /include -I$( CUDA_PATH) /targets/$( UNAME_M) -linux/include
MK_LDFLAGS += -lcuda -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L$( CUDA_PATH) /lib64 -L/usr/lib64 -L$( CUDA_PATH) /targets/$( UNAME_M) -linux/lib -L/usr/lib/wsl/lib
2023-10-29 16:32:28 +00:00
OBJS += ggml-cuda.o
2024-02-03 19:14:59 +00:00
MK_NVCCFLAGS += -use_fast_math
2024-02-19 12:45:41 +00:00
i f d e f L L A M A _ F A T A L _ W A R N I N G S
MK_NVCCFLAGS += -Werror all-warnings
e n d i f # LLAMA_FATAL_WARNINGS
2023-12-22 15:11:12 +00:00
i f n d e f J E T S O N _ E O L _ M O D U L E _ D E T E C T
MK_NVCCFLAGS += --forward-unknown-to-host-compiler
e n d i f # JETSON_EOL_MODULE_DETECT
2023-12-13 12:04:25 +00:00
i f d e f L L A M A _ D E B U G
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -lineinfo
2023-12-24 13:34:22 +00:00
e n d i f # LLAMA_DEBUG
2023-07-21 10:38:57 +00:00
i f d e f L L A M A _ C U D A _ N V C C
2024-02-05 18:33:00 +00:00
NVCC = $( CCACHE) $( LLAMA_CUDA_NVCC)
2023-07-21 10:38:57 +00:00
e l s e
2024-02-05 18:33:00 +00:00
NVCC = $( CCACHE) nvcc
2023-07-21 10:38:57 +00:00
e n d i f #LLAMA_CUDA_NVCC
2023-07-07 18:25:25 +00:00
i f d e f C U D A _ D O C K E R _ A R C H
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -Wno-deprecated-gpu-targets -arch= $( CUDA_DOCKER_ARCH)
e l s e i f n d e f C U D A _ P O W E R _ A R C H
MK_NVCCFLAGS += -arch= native
2023-10-23 20:46:05 +00:00
e n d i f # CUDA_DOCKER_ARCH
2023-07-05 12:19:42 +00:00
i f d e f L L A M A _ C U D A _ F O R C E _ D M M V
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_FORCE_DMMV
2023-07-05 12:19:42 +00:00
e n d i f # LLAMA_CUDA_FORCE_DMMV
2023-10-27 14:01:23 +00:00
i f d e f L L A M A _ C U D A _ F O R C E _ M M Q
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_FORCE_MMQ
2023-10-27 14:01:23 +00:00
e n d i f # LLAMA_CUDA_FORCE_MMQ
2023-05-25 21:07:29 +00:00
i f d e f L L A M A _ C U D A _ D M M V _ X
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_DMMV_X= $( LLAMA_CUDA_DMMV_X)
2023-05-25 21:07:29 +00:00
e l s e
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_DMMV_X= 32
2023-05-25 21:07:29 +00:00
e n d i f # LLAMA_CUDA_DMMV_X
2023-07-05 12:19:42 +00:00
i f d e f L L A M A _ C U D A _ M M V _ Y
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_MMV_Y= $( LLAMA_CUDA_MMV_Y)
2023-07-05 12:19:42 +00:00
e l s e i f d e f L L A M A _ C U D A _ D M M V _ Y
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_MMV_Y= $( LLAMA_CUDA_DMMV_Y) # for backwards compatibility
2023-05-25 21:07:29 +00:00
e l s e
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_MMV_Y= 1
2023-07-05 12:19:42 +00:00
e n d i f # LLAMA_CUDA_MMV_Y
2023-07-29 21:04:44 +00:00
i f d e f L L A M A _ C U D A _ F 1 6
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_F16
2023-07-29 21:04:44 +00:00
e n d i f # LLAMA_CUDA_F16
2023-06-19 08:23:56 +00:00
i f d e f L L A M A _ C U D A _ D M M V _ F 1 6
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_F16
2023-06-19 08:23:56 +00:00
e n d i f # LLAMA_CUDA_DMMV_F16
2023-06-16 17:08:44 +00:00
i f d e f L L A M A _ C U D A _ K Q U A N T S _ I T E R
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DK_QUANTS_PER_ITERATION= $( LLAMA_CUDA_KQUANTS_ITER)
2023-06-16 17:08:44 +00:00
e l s e
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DK_QUANTS_PER_ITERATION= 2
2023-06-16 17:08:44 +00:00
e n d i f
2023-09-17 14:37:53 +00:00
i f d e f L L A M A _ C U D A _ P E E R _ M A X _ B A T C H _ S I Z E
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_PEER_MAX_BATCH_SIZE= $( LLAMA_CUDA_PEER_MAX_BATCH_SIZE)
2023-09-17 14:37:53 +00:00
e l s e
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -DGGML_CUDA_PEER_MAX_BATCH_SIZE= 128
2023-09-17 14:37:53 +00:00
e n d i f # LLAMA_CUDA_PEER_MAX_BATCH_SIZE
2023-07-31 13:44:35 +00:00
#ifdef LLAMA_CUDA_CUBLAS
2023-12-13 17:10:10 +00:00
# MK_NVCCFLAGS += -DGGML_CUDA_CUBLAS
2023-07-31 13:44:35 +00:00
#endif # LLAMA_CUDA_CUBLAS
2023-07-21 10:38:57 +00:00
i f d e f L L A M A _ C U D A _ C C B I N
2023-12-13 17:10:10 +00:00
MK_NVCCFLAGS += -ccbin $( LLAMA_CUDA_CCBIN)
2023-07-21 10:38:57 +00:00
e n d i f
2023-04-20 01:14:14 +00:00
ggml-cuda.o : ggml -cuda .cu ggml -cuda .h
2023-12-22 15:11:12 +00:00
i f d e f J E T S O N _ E O L _ M O D U L E _ D E T E C T
2024-02-19 20:54:12 +00:00
$( NVCC) -I. -Icommon -D_XOPEN_SOURCE= 600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/usr/local/cuda/targets/aarch64-linux/include -std= c++11 -O3 $( NVCCFLAGS) $( CPPFLAGS) -Xcompiler " $( CUDA_CXXFLAGS) " -c $< -o $@
2023-12-22 15:11:12 +00:00
e l s e
2024-02-19 20:54:12 +00:00
$( NVCC) $( NVCCFLAGS) $( CPPFLAGS) -Xcompiler " $( CUDA_CXXFLAGS) " -c $< -o $@
2023-12-22 15:11:12 +00:00
e n d i f # JETSON_EOL_MODULE_DETECT
2023-05-25 21:07:29 +00:00
e n d i f # LLAMA_CUBLAS
2023-06-04 20:34:30 +00:00
ggml : add CLBlast support (#1164)
* Allow use of OpenCL GPU-based BLAS using ClBlast instead of OpenBLAS for context processing
* Improve ClBlast implementation, avoid recreating buffers, remove redundant transfers
* Finish merge of ClBlast support
* Move CLBlast implementation to separate file
Add buffer reuse code (adapted from slaren's cuda implementation)
* Add q4_2 and q4_3 CLBlast support, improve code
* Double CLBlast speed by disabling OpenBLAS thread workaround
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <2141330+slaren@users.noreply.github.com>
* Fix device selection env variable names
* Fix cast in opencl kernels
* Add CLBlast to CMakeLists.txt
* Replace buffer pool with static buffers a, b, qb, c
Fix compile warnings
* Fix typos, use GGML_TYPE defines, improve code
* Improve btype dequant kernel selection code, add error if type is unsupported
* Improve code quality
* Move internal stuff out of header
* Use internal enums instead of CLBlast enums
* Remove leftover C++ includes and defines
* Make event use easier to read
Co-authored-by: Henri Vasserman <henv@hot.ee>
* Use c compiler for opencl files
* Simplify code, fix include
* First check error, then release event
* Make globals static, fix indentation
* Rename dequant kernels file to conform with other file names
* Fix import cl file name
---------
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <2141330+slaren@users.noreply.github.com>
Co-authored-by: Henri Vasserman <henv@hot.ee>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-04-28 14:57:16 +00:00
i f d e f L L A M A _ C L B L A S T
2023-07-23 11:52:08 +00:00
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_CLBLAST $( shell pkg-config --cflags-only-I clblast OpenCL)
MK_CFLAGS += $( shell pkg-config --cflags-only-other clblast OpenCL)
MK_CXXFLAGS += $( shell pkg-config --cflags-only-other clblast OpenCL)
2023-07-23 11:52:08 +00:00
2023-05-05 12:18:21 +00:00
# Mac provides OpenCL as a framework
ifeq ( $( UNAME_S) ,Darwin)
2023-09-03 05:26:59 +00:00
MK_LDFLAGS += -lclblast -framework OpenCL
2023-05-05 12:18:21 +00:00
else
2023-09-03 05:26:59 +00:00
MK_LDFLAGS += $( shell pkg-config --libs clblast OpenCL)
2023-05-05 12:18:21 +00:00
endif
ggml : add CLBlast support (#1164)
* Allow use of OpenCL GPU-based BLAS using ClBlast instead of OpenBLAS for context processing
* Improve ClBlast implementation, avoid recreating buffers, remove redundant transfers
* Finish merge of ClBlast support
* Move CLBlast implementation to separate file
Add buffer reuse code (adapted from slaren's cuda implementation)
* Add q4_2 and q4_3 CLBlast support, improve code
* Double CLBlast speed by disabling OpenBLAS thread workaround
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <2141330+slaren@users.noreply.github.com>
* Fix device selection env variable names
* Fix cast in opencl kernels
* Add CLBlast to CMakeLists.txt
* Replace buffer pool with static buffers a, b, qb, c
Fix compile warnings
* Fix typos, use GGML_TYPE defines, improve code
* Improve btype dequant kernel selection code, add error if type is unsupported
* Improve code quality
* Move internal stuff out of header
* Use internal enums instead of CLBlast enums
* Remove leftover C++ includes and defines
* Make event use easier to read
Co-authored-by: Henri Vasserman <henv@hot.ee>
* Use c compiler for opencl files
* Simplify code, fix include
* First check error, then release event
* Make globals static, fix indentation
* Rename dequant kernels file to conform with other file names
* Fix import cl file name
---------
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <2141330+slaren@users.noreply.github.com>
Co-authored-by: Henri Vasserman <henv@hot.ee>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-04-28 14:57:16 +00:00
OBJS += ggml-opencl.o
2023-06-04 20:34:30 +00:00
2023-05-22 21:33:24 +00:00
ggml-opencl.o : ggml -opencl .cpp ggml -opencl .h
$( CXX) $( CXXFLAGS) -c $< -o $@
2023-06-04 20:34:30 +00:00
e n d i f # LLAMA_CLBLAST
ggml : add Vulkan backend (#2059)
* Vulkan loader code
* Fix matmul kernel, continue implementation
* Continue implementation
* Vulkan memory management
* Vulkan development
* Matmul call
* Add aligned malloc and free for VMA
* Continue implementation
* First matmul success
* GEMM Kernel optimization
* 1D Blocktiling
* 2D Blocktiling
* Write coalescing
* Continue vulkan implementation and optimization
* First FP16 attempt, disabled for now
* Code abstraction, FP16 implementation, fix kernel, add FP16 to FP32 kernel
* Enable device extensions properly, restore fp16 matmul op
* Fix mulmat_f16
* Output FP32 in fp16 matmul shader
* Fix f16_to_f32 kernel
* dequant_q4_0 kernel
* Add VMA library
* Avoid requesting dedicated memory, VMA can decide that by itself
* Add bounds checking to matmul kernels, improve implementation, fix command buffers not freed properly
* add cmake commands
* Add 2d write operation, profiling code
* Fix 2d write
* Fix queue selection for AMD RADV
* Fix trailing whitespace in vk_mem_alloc.h
* Add WIP warp tile mat mul shaders
* Disable glslc optimization
* Disable glslc optimization for CMake
* Optimize warptile matmul shader, replace blocktile with it
* Add split-k optimization for small matrix multiplication
Use semaphores for synchronization instead of fences or waitidle
Rework async write/read for synchronization
* Fix validation errors, improve compatibility with AMD GPUs
* Rework command buffer handling
* Variable matmul kernel using specialization constants
* Fix synchronization on AMD, add barriers for buffer ownership transfer, add debug flag and prints
* Reuse semaphores
* Handle stage flags during command buffer submission properly
* Increase matmul test runs for consistent results
* Fix F32 matmul
* Add vectorized loading and zeropadding for matrix multiplication
* Use pinned memory for f16 preprocessing
* Don't force aligned matmul
* Don't free before queue done
* Replace VMA library with native Vulkan buffer management
* Basic offloading support with mul_f32 and dmmv for q4_0
* Run glslc commands in parallel
* Unroll loops in dmmv shader
* Reduce usage of waitIdle
* Reuse pinned allocation for f16 conversion
* Handle devices with only a single queue
* Fix trailing whitespace in CMakeLists.txt
* Allow parallel execution of kernels, parallelize third and fourth dimension calls
* Add fallback for devices only supporting one DescriptorSet per DescriptorPool
* Move to graph function similar to CUDA implementation
* Use F16 kernel for most things, replace q_f32 with mul_mat_q_f16 function
* Add F32 dmmv shaders
* Batch submissions
* Add .spv to gitignore
* Split off matrix vector multiplication for separate optimization
* Use single command buffer for matrix vector multiplication ops
* Reduce overhead of mul_f32 calls by using a single command buffer
* Add submission batching to mul_f32
* Fix tests
* Add missing barrier
* Add further missing barrier
* Add further ops
* Replace vk::QueueFamilyIgnored with VK_QUEUE_FAMILY_IGNORED to support more Vulkan header versions
* Remove unnecessary cblas link
* Fix descriptor set pre-allocation assert
* Add runtime shader compilation, start transferring shaders to this approach
* Transfer remaining shaders to header and compile on runtime
* Fix fp32 fallback if device doesn't support fp16, add force disable env var GGML_VULKAN_DISABLE_F16
* Add support for q4_1, q5_0, q5_1 and q8_0
* Remove unnecessary scalar layout extension
* Parse graph early to pre-record command buffers
* Add q6_k support
* Add multi-submit for command buffers
* Fix q6_k dequant shader for AMD
* Fix q6_k for GPUs without fp16 support
* Simplify q6_k fp16 fix
* Minor fixes
* Fix wg_denom of m-mulmat shaders
* Add Python-based Vulkan shader generator
* Replace shaderc dependency with precompiled shaders
Fix python script to generate shaders
* Clean up code
* Fix shader generator script Windows compatibility
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
* Close file before deletion
* Fix vulkan shader fp32 name
* Add q2_k and q3_k support
Add validation check to compare shader results to cpu results
* Add q4_k support
* Add q5_k support
* Bake SPIR-V bytecode into the library instead of loading shaders from file
* Switch to signal semaphores for flexibility
Prepare broadcasting support for mul mat
* Finish broadcasting mul mat support for GQA
* Clean up unused functions
Add repeat op
* Add further ops, not yet enabled. Improve semaphore code
* Reduce number of used semaphores by utilizing timelines more properly
* Remove queue information
* Reuse timeline semaphores, allow parallel operation with binary semaphores to work around nvidia driver limitations
* Add Vulkan to llama-bench
* Remove cblas dependency
* Fix matmul k-split bug
* Fix q4_k dmmv K_QUANTS_PER_ITERATION 1 shader
* Add RMS Norm shader, rework op_f32 shader setup, fix matmul bug
* Fix issues with float16 overflows in shaders
* Fix issues with older Vulkan headers on Ubuntu 22.04
* Allow multi-op partial offloading by parsing the graph to preallocate enough between-op buffers
* Implement further ops, rework op_f32 calls, fix bugs
* Finish full offloading support, add last remaining ops, fix bugs, remove redundant code
* Upload generated file ggml-vulkan-shaders.hpp, remove redundant shaders
* Merge upstream changes, fix conflicts, adapt soft_max op
* Fix Python and shader header format
* Free model gpu buffers on exit
* Use single queue per device to simplify code
* Add matmul shader support for running multiple calculations in parallel
* Switch from semaphore-synchronized multiple command buffers per op to single command buffer for multiple ops, whole graph if possible
* Fix missing event cast
* Replace uint64_t(-1) with UINT64_MAX, rename function for clarity
* Fix warning about empty C function parameters
* Fix compiler warnings
* Properly implement Vulkan backend buffer handling
* Fix oversized host staging buffers
* Simplify barrier synchronization calls
* Fix gcc warnings
* Implement max_size for backend buffer types to limit the size of a single allocation
* Use min of maxMemoryAllocationSize and maxBufferSize for device max allocation size
* refactor multi buf
* Disable unsupported ops to fix tests
* Check for maintenance4 support before using it
* Handle devices with only a single queue
* Fix single queue logic
* propagate buffer usage in multi buffers
* Implement rope_neox op
* Cleanup header and other files
* Simplify gpu_extras by removing events and putting staging memcpys into contexts
* Move queue into context
Add not-yet-enabled async backend ops
* Simplify context use, optimize matmul shader for warp size 64 (AMD GCN), fix split_k matmul shader optimization
* Add get_max_size to SYCL backend.
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* llama : fix trailing whitespace
---------
Co-authored-by: Henri Vasserman <henv@hot.ee>
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <slarengh@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2024-01-28 17:03:59 +00:00
i f d e f L L A M A _ V U L K A N
MK_CPPFLAGS += -DGGML_USE_VULKAN
MK_LDFLAGS += -lvulkan
OBJS += ggml-vulkan.o
i f d e f L L A M A _ V U L K A N _ C H E C K _ R E S U L T S
MK_CPPFLAGS += -DGGML_VULKAN_CHECK_RESULTS
e n d i f
2024-02-03 17:15:00 +00:00
i f d e f L L A M A _ V U L K A N _ D E B U G
MK_CPPFLAGS += -DGGML_VULKAN_DEBUG
e n d i f
i f d e f L L A M A _ V U L K A N _ V A L I D A T E
MK_CPPFLAGS += -DGGML_VULKAN_VALIDATE
e n d i f
i f d e f L L A M A _ V U L K A N _ R U N _ T E S T S
MK_CPPFLAGS += -DGGML_VULKAN_RUN_TESTS
e n d i f
ggml : add Vulkan backend (#2059)
* Vulkan loader code
* Fix matmul kernel, continue implementation
* Continue implementation
* Vulkan memory management
* Vulkan development
* Matmul call
* Add aligned malloc and free for VMA
* Continue implementation
* First matmul success
* GEMM Kernel optimization
* 1D Blocktiling
* 2D Blocktiling
* Write coalescing
* Continue vulkan implementation and optimization
* First FP16 attempt, disabled for now
* Code abstraction, FP16 implementation, fix kernel, add FP16 to FP32 kernel
* Enable device extensions properly, restore fp16 matmul op
* Fix mulmat_f16
* Output FP32 in fp16 matmul shader
* Fix f16_to_f32 kernel
* dequant_q4_0 kernel
* Add VMA library
* Avoid requesting dedicated memory, VMA can decide that by itself
* Add bounds checking to matmul kernels, improve implementation, fix command buffers not freed properly
* add cmake commands
* Add 2d write operation, profiling code
* Fix 2d write
* Fix queue selection for AMD RADV
* Fix trailing whitespace in vk_mem_alloc.h
* Add WIP warp tile mat mul shaders
* Disable glslc optimization
* Disable glslc optimization for CMake
* Optimize warptile matmul shader, replace blocktile with it
* Add split-k optimization for small matrix multiplication
Use semaphores for synchronization instead of fences or waitidle
Rework async write/read for synchronization
* Fix validation errors, improve compatibility with AMD GPUs
* Rework command buffer handling
* Variable matmul kernel using specialization constants
* Fix synchronization on AMD, add barriers for buffer ownership transfer, add debug flag and prints
* Reuse semaphores
* Handle stage flags during command buffer submission properly
* Increase matmul test runs for consistent results
* Fix F32 matmul
* Add vectorized loading and zeropadding for matrix multiplication
* Use pinned memory for f16 preprocessing
* Don't force aligned matmul
* Don't free before queue done
* Replace VMA library with native Vulkan buffer management
* Basic offloading support with mul_f32 and dmmv for q4_0
* Run glslc commands in parallel
* Unroll loops in dmmv shader
* Reduce usage of waitIdle
* Reuse pinned allocation for f16 conversion
* Handle devices with only a single queue
* Fix trailing whitespace in CMakeLists.txt
* Allow parallel execution of kernels, parallelize third and fourth dimension calls
* Add fallback for devices only supporting one DescriptorSet per DescriptorPool
* Move to graph function similar to CUDA implementation
* Use F16 kernel for most things, replace q_f32 with mul_mat_q_f16 function
* Add F32 dmmv shaders
* Batch submissions
* Add .spv to gitignore
* Split off matrix vector multiplication for separate optimization
* Use single command buffer for matrix vector multiplication ops
* Reduce overhead of mul_f32 calls by using a single command buffer
* Add submission batching to mul_f32
* Fix tests
* Add missing barrier
* Add further missing barrier
* Add further ops
* Replace vk::QueueFamilyIgnored with VK_QUEUE_FAMILY_IGNORED to support more Vulkan header versions
* Remove unnecessary cblas link
* Fix descriptor set pre-allocation assert
* Add runtime shader compilation, start transferring shaders to this approach
* Transfer remaining shaders to header and compile on runtime
* Fix fp32 fallback if device doesn't support fp16, add force disable env var GGML_VULKAN_DISABLE_F16
* Add support for q4_1, q5_0, q5_1 and q8_0
* Remove unnecessary scalar layout extension
* Parse graph early to pre-record command buffers
* Add q6_k support
* Add multi-submit for command buffers
* Fix q6_k dequant shader for AMD
* Fix q6_k for GPUs without fp16 support
* Simplify q6_k fp16 fix
* Minor fixes
* Fix wg_denom of m-mulmat shaders
* Add Python-based Vulkan shader generator
* Replace shaderc dependency with precompiled shaders
Fix python script to generate shaders
* Clean up code
* Fix shader generator script Windows compatibility
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
* Close file before deletion
* Fix vulkan shader fp32 name
* Add q2_k and q3_k support
Add validation check to compare shader results to cpu results
* Add q4_k support
* Add q5_k support
* Bake SPIR-V bytecode into the library instead of loading shaders from file
* Switch to signal semaphores for flexibility
Prepare broadcasting support for mul mat
* Finish broadcasting mul mat support for GQA
* Clean up unused functions
Add repeat op
* Add further ops, not yet enabled. Improve semaphore code
* Reduce number of used semaphores by utilizing timelines more properly
* Remove queue information
* Reuse timeline semaphores, allow parallel operation with binary semaphores to work around nvidia driver limitations
* Add Vulkan to llama-bench
* Remove cblas dependency
* Fix matmul k-split bug
* Fix q4_k dmmv K_QUANTS_PER_ITERATION 1 shader
* Add RMS Norm shader, rework op_f32 shader setup, fix matmul bug
* Fix issues with float16 overflows in shaders
* Fix issues with older Vulkan headers on Ubuntu 22.04
* Allow multi-op partial offloading by parsing the graph to preallocate enough between-op buffers
* Implement further ops, rework op_f32 calls, fix bugs
* Finish full offloading support, add last remaining ops, fix bugs, remove redundant code
* Upload generated file ggml-vulkan-shaders.hpp, remove redundant shaders
* Merge upstream changes, fix conflicts, adapt soft_max op
* Fix Python and shader header format
* Free model gpu buffers on exit
* Use single queue per device to simplify code
* Add matmul shader support for running multiple calculations in parallel
* Switch from semaphore-synchronized multiple command buffers per op to single command buffer for multiple ops, whole graph if possible
* Fix missing event cast
* Replace uint64_t(-1) with UINT64_MAX, rename function for clarity
* Fix warning about empty C function parameters
* Fix compiler warnings
* Properly implement Vulkan backend buffer handling
* Fix oversized host staging buffers
* Simplify barrier synchronization calls
* Fix gcc warnings
* Implement max_size for backend buffer types to limit the size of a single allocation
* Use min of maxMemoryAllocationSize and maxBufferSize for device max allocation size
* refactor multi buf
* Disable unsupported ops to fix tests
* Check for maintenance4 support before using it
* Handle devices with only a single queue
* Fix single queue logic
* propagate buffer usage in multi buffers
* Implement rope_neox op
* Cleanup header and other files
* Simplify gpu_extras by removing events and putting staging memcpys into contexts
* Move queue into context
Add not-yet-enabled async backend ops
* Simplify context use, optimize matmul shader for warp size 64 (AMD GCN), fix split_k matmul shader optimization
* Add get_max_size to SYCL backend.
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* llama : fix trailing whitespace
---------
Co-authored-by: Henri Vasserman <henv@hot.ee>
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com>
Co-authored-by: slaren <slarengh@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2024-01-28 17:03:59 +00:00
ggml-vulkan.o : ggml -vulkan .cpp ggml -vulkan .h
$( CXX) $( CXXFLAGS) -c $< -o $@
e n d i f # LLAMA_VULKAN
2023-08-25 09:09:42 +00:00
i f d e f L L A M A _ H I P B L A S
2023-12-17 15:23:33 +00:00
ifeq ( $( wildcard /opt/rocm) ,)
ROCM_PATH ?= /usr
GPU_TARGETS ?= $( shell $( shell which amdgpu-arch) )
else
ROCM_PATH ?= /opt/rocm
GPU_TARGETS ?= $( shell $( ROCM_PATH) /llvm/bin/amdgpu-arch)
endif
2024-02-05 18:33:00 +00:00
HIPCC ?= $( CCACHE) $( ROCM_PATH) /bin/hipcc
2023-08-25 09:09:42 +00:00
LLAMA_CUDA_DMMV_X ?= 32
LLAMA_CUDA_MMV_Y ?= 1
LLAMA_CUDA_KQUANTS_ITER ?= 2
2023-09-03 05:26:59 +00:00
MK_CPPFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS
2023-12-22 08:03:25 +00:00
i f d e f L L A M A _ H I P _ U M A
MK_CPPFLAGS += -DGGML_HIP_UMA
e n d i f # LLAMA_HIP_UMA
2023-09-03 05:26:59 +00:00
MK_LDFLAGS += -L$( ROCM_PATH) /lib -Wl,-rpath= $( ROCM_PATH) /lib
MK_LDFLAGS += -lhipblas -lamdhip64 -lrocblas
2023-08-25 09:09:42 +00:00
HIPFLAGS += $( addprefix --offload-arch= ,$( GPU_TARGETS) )
HIPFLAGS += -DGGML_CUDA_DMMV_X= $( LLAMA_CUDA_DMMV_X)
HIPFLAGS += -DGGML_CUDA_MMV_Y= $( LLAMA_CUDA_MMV_Y)
HIPFLAGS += -DK_QUANTS_PER_ITERATION= $( LLAMA_CUDA_KQUANTS_ITER)
i f d e f L L A M A _ C U D A _ F O R C E _ D M M V
HIPFLAGS += -DGGML_CUDA_FORCE_DMMV
e n d i f # LLAMA_CUDA_FORCE_DMMV
OBJS += ggml-cuda.o
ggml-cuda.o : ggml -cuda .cu ggml -cuda .h
$( HIPCC) $( CXXFLAGS) $( HIPFLAGS) -x hip -c -o $@ $<
e n d i f # LLAMA_HIPBLAS
2023-06-04 20:34:30 +00:00
i f d e f L L A M A _ M E T A L
2023-09-05 19:12:00 +00:00
MK_CPPFLAGS += -DGGML_USE_METAL
2023-09-03 05:26:59 +00:00
MK_LDFLAGS += -framework Foundation -framework Metal -framework MetalKit
OBJS += ggml-metal.o
2023-09-05 22:21:10 +00:00
i f d e f L L A M A _ M E T A L _ N D E B U G
MK_CPPFLAGS += -DGGML_METAL_NDEBUG
e n d i f
2024-02-20 09:58:36 +00:00
i f d e f L L A M A _ M E T A L _ E M B E D _ L I B R A R Y
MK_CPPFLAGS += -DGGML_METAL_EMBED_LIBRARY
OBJS += ggml-metal-embed.o
e n d i f
2023-06-04 20:34:30 +00:00
e n d i f # LLAMA_METAL
2023-07-14 17:34:40 +00:00
i f d e f L L A M A _ M E T A L
ggml-metal.o : ggml -metal .m ggml -metal .h
$( CC) $( CFLAGS) -c $< -o $@
2024-02-20 09:58:36 +00:00
i f d e f L L A M A _ M E T A L _ E M B E D _ L I B R A R Y
ggml-metal-embed.o : ggml -metal .metal
@echo "Embedding Metal library"
$( eval TEMP_ASSEMBLY = $( shell mktemp) )
@echo ".section __DATA, __ggml_metallib" > $( TEMP_ASSEMBLY)
@echo ".globl _ggml_metallib_start" >> $( TEMP_ASSEMBLY)
@echo "_ggml_metallib_start:" >> $( TEMP_ASSEMBLY)
@echo " .incbin \" $<\" " >> $( TEMP_ASSEMBLY)
@echo ".globl _ggml_metallib_end" >> $( TEMP_ASSEMBLY)
@echo "_ggml_metallib_end:" >> $( TEMP_ASSEMBLY)
@$( AS) $( TEMP_ASSEMBLY) -o $@
@rm -f ${ TEMP_ASSEMBLY }
e n d i f
2023-07-14 17:34:40 +00:00
e n d i f # LLAMA_METAL
i f d e f L L A M A _ M P I
ggml-mpi.o : ggml -mpi .c ggml -mpi .h
$( CC) $( CFLAGS) -c $< -o $@
e n d i f # LLAMA_MPI
2023-12-13 17:10:10 +00:00
GF_CC := $( CC)
i n c l u d e s c r i p t s / g e t - f l a g s . m k
2023-09-03 05:26:59 +00:00
# combine build flags with cmdline overrides
2024-02-19 20:54:12 +00:00
override CPPFLAGS : = $( MK_CPPFLAGS ) $( CPPFLAGS )
override CFLAGS : = $( CPPFLAGS ) $( MK_CFLAGS ) $( GF_CFLAGS ) $( CFLAGS )
BASE_CXXFLAGS := $( MK_CXXFLAGS) $( CXXFLAGS)
override CXXFLAGS : = $( BASE_CXXFLAGS ) $( HOST_CXXFLAGS ) $( GF_CXXFLAGS ) $( CPPFLAGS )
2023-12-13 17:10:10 +00:00
override NVCCFLAGS : = $( MK_NVCCFLAGS ) $( NVCCFLAGS )
override LDFLAGS : = $( MK_LDFLAGS ) $( LDFLAGS )
# identify CUDA host compiler
i f d e f L L A M A _ C U B L A S
GF_CC := $( NVCC) $( NVCCFLAGS) 2>/dev/null .c -Xcompiler
i n c l u d e s c r i p t s / g e t - f l a g s . m k
2024-02-18 21:21:52 +00:00
CUDA_CXXFLAGS := $( BASE_CXXFLAGS) $( GF_CXXFLAGS) -Wno-pedantic
2023-12-13 17:10:10 +00:00
e n d i f
2023-09-03 05:26:59 +00:00
2023-03-10 18:40:58 +00:00
#
# Print build information
#
$(info I llama.cpp build info : )
2023-09-16 14:55:43 +00:00
$(info I UNAME_S : $( UNAME_S ) )
$(info I UNAME_P : $( UNAME_P ) )
$(info I UNAME_M : $( UNAME_M ) )
$(info I CFLAGS : $( CFLAGS ) )
$(info I CXXFLAGS : $( CXXFLAGS ) )
$(info I NVCCFLAGS : $( NVCCFLAGS ) )
$(info I LDFLAGS : $( LDFLAGS ) )
2024-02-03 19:15:13 +00:00
$(info I CC : $( shell $ ( CC ) --version | head -n 1) )
$(info I CXX : $( shell $ ( CXX ) --version | head -n 1) )
i f d e f L L A M A _ C U B L A S
$(info I NVCC : $( shell $ ( NVCC ) --version | tail -n 1) )
2024-02-25 16:46:49 +00:00
CUDA_VERSION := $( shell $( NVCC) --version | grep -oP 'release (\K[0-9]+\.[0-9])' )
2024-02-13 11:38:37 +00:00
i f e q ( $( shell awk -v "v =$ ( CUDA_VERSION ) " 'BEGIN { print ( v < 11.7) }') , 1 )
i f n d e f C U D A _ D O C K E R _ A R C H
i f n d e f C U D A _ P O W E R _ A R C H
$(error I ERROR : For CUDA versions < 11.7 a target CUDA architecture must be explicitly provided via CUDA_DOCKER_ARCH )
e n d i f # CUDA_POWER_ARCH
e n d i f # CUDA_DOCKER_ARCH
e n d i f # eq ($(shell echo "$(CUDA_VERSION) < 11.7" | bc),1)
2024-02-03 19:15:13 +00:00
e n d i f # LLAMA_CUBLAS
2023-03-10 18:40:58 +00:00
$( info )
#
# Build library
#
2023-06-07 07:59:52 +00:00
ggml.o : ggml .c ggml .h ggml -cuda .h
2023-04-14 19:39:48 +00:00
$( CC) $( CFLAGS) -c $< -o $@
2023-03-10 18:40:58 +00:00
2023-07-30 13:58:01 +00:00
ggml-alloc.o : ggml -alloc .c ggml .h ggml -alloc .h
$( CC) $( CFLAGS) -c $< -o $@
2023-10-08 17:19:14 +00:00
ggml-backend.o : ggml -backend .c ggml .h ggml -backend .h
$( CC) $( CFLAGS) -c $< -o $@
2023-10-29 16:32:28 +00:00
ggml-quants.o : ggml -quants .c ggml .h ggml -quants .h
$( CC) $( CFLAGS) -c $< -o $@
OBJS += ggml-alloc.o ggml-backend.o ggml-quants.o
2023-07-30 13:58:01 +00:00
2023-10-08 17:19:14 +00:00
llama.o : llama .cpp ggml .h ggml -alloc .h ggml -backend .h ggml -cuda .h ggml -metal .h llama .h
2023-04-14 19:39:48 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $@
2023-03-22 05:32:36 +00:00
2023-10-29 16:33:47 +00:00
COMMON_H_DEPS = common/common.h common/sampling.h common/log.h
2023-11-02 06:50:16 +00:00
COMMON_DEPS = common.o sampling.o grammar-parser.o build-info.o
2023-10-11 19:35:46 +00:00
2023-11-02 06:50:16 +00:00
common.o : common /common .cpp $( COMMON_H_DEPS )
2023-10-11 19:35:46 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $@
sampling.o : common /sampling .cpp $( COMMON_H_DEPS )
2023-04-14 19:39:48 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $@
2023-03-10 18:40:58 +00:00
2023-08-21 20:07:43 +00:00
console.o : common /console .cpp common /console .h
2023-08-04 15:20:12 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $@
2023-08-21 20:07:43 +00:00
grammar-parser.o : common /grammar -parser .cpp common /grammar -parser .h
2023-07-24 03:58:10 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $@
train : finetune LORA (#2632)
* fix track_max_mem in forward_batch_wo_cache_flash_attn_train
* remove unnecessary Adam(W) optimizer tensors.
reduces optimizer memory overhead from 7*modelsize to 2*modelsize.
additionally allows to optimize models with more than 2^31 parameters by replacing int with int64_t.
bumps training checkpoint file version, but old checkpoints can still be read.
new version with less tensors is saved.
* add gradient clipping to AdamW
* Fix reset of unused g->nodes and g->grads to NULL
* implement gradient checkpointing for training
reduces memory overhead from O(n_layer) to O(sqrt(n_layer))
as explained in readme of https://github.com/cybertronai/gradient-checkpointing
* remove unused compute buffer 3
* add and use function ggml_build_backward_expand to avoid stack overflows with large maximum number of nodes
GGML_API void ggml_build_backward_expand(struct ggml_context * ctx, struct ggml_cgraph * gf, struct ggml_cgraph * gb, bool keep);
* change AdamW decay parameter to work like the torch AdamW decay parameter
It is now relative to Adam learning rate `alpha*sched`.
Before that it was relative to `sched` only.
`alpha` being the maximum learning rate and `sched` being a scaling parameter in [0..1]
* change default AdamW weight decay parameter used in training to 0.1 as used in nanoGPT
* change default AdamW weight decay parameter defined in ggml to 0.0, making Adam default instead of AdamW
btw: the default weight decay parameter for torch.optim.AdamW is 0.01
* bug fixes for cross entropy loss
ggml_cross_entropy_loss: sums where not correctly added in workload of each thread
ggml_cross_entropy_loss_back: simplify backward process, reducing numerical issues
guard usage of exp f16 lookup in cross entropy by #define GGML_CROSS_ENTROPY_EXP_FP16
cross entropy loss is only used once during training, but it is quite sensitive to numerical errors introduced by exp-f16-lookup.
so exp-f16-lookup for cross entropy loss is disabled by default, trading better gradients for very slightly worse runtime performance.
* fix test-grad0 for cross_entropy_loss
the second argument to cross_entropy_loss must sum up to 1 for each row
* fix test-grad0 for soft_max
dont use only sum as aggregation, because sum of softmax is always 1 -> finite differences should not work
instead use sum(log(soft_max()*(1-eps)+eps)); use eps to avoid log(0)
* improve finite differences of test-grad0 by using double instead of float
* change cross_entropy_loss to output average over all rows
this helps keeping the loss and gradients in a sane range
* improve gradient checkpointing
sqrt(n_layers) is only the best checkpoint step when mem size of checkpoints and mem size of layers are equal.
since layers require more memory than the single-tensor-checkpoint we use, the optimal values are compute different:
```
given: n, u, v
objective: minimize(a*u+b*v) where a*b=n, a>0, b>0
b=n/a
minimize(a*u+v*n/a)
diff(a*u+v*n/a, a) = u - (v*n/a)/a
diff(a*u+v*n/a, a) == 0
u - (v*n/a)/a == 0
u == v*n/(a*a)
u*a*a = v*n
a*a = v*n/u
a = sqrt(n*v/u)
```
this change results in more checkpoints, requiring less layers to store between checkpoints, overall improving memory usage.
* disable gradient checkpointing debug output
* llama : fix rope usage in train-text-from-scratch after ChatGLM change
* add more training parameters:
--enable-restart N Only for Adam optimizer. Enable restarts of cos-decay
--disable-restart N Only for Adam optimizer. Disable restarts of cos-decay
--opt-past N Number of optimization iterations to track for delta convergence test. Disabled when zero.
--opt-delta N Maximum delta for delta convergence test. Disabled when <= zero.
--opt-max-no-improvement N Maximum number of optimization iterations with no improvement. Disabled when <= zero.
--adam-epsf N AdamW epsilon for convergence test. Disabled when <= zero.
--adam-min-alpha N Adam minimum learning rate alpha, usually 0.1 * alpha
* replace memcpy with reshape operation so that the graph is not cut at the input
this makes it possible to store other values into the input tensor and then simply recompute the graph without rebuilding it
* remove unused function argument from get_example_targets_batch
* measure and print total training time
* add optimization callback to ggml_opt_resume_g
this callback is called before each iteration with custom data and pointer to learning schedule parameter (only used in Adam(W)).
can be used for dynamic learning schedule and setting input data for batches before each iteration
* use optimization callback in training
allows dynamic learning schedule and different batch data for each iteration without relying on low n_iter and high n_examples parameters
reduces runtime by avoiding restart of optimization function and improves training convergence by providing a different batch for each iteration
* add minimum number of tensor dimensions to apply weight decay (default 2)
this allows to not apply weight decay to bias parameters
* rename training parameter cos-decay-alpha to cos-decay-min and clarify that adam-min-alpha also applies to warmup
* fix increase of model.train_samples and model.train_tokens
now that each optimizer iteration gets its own batch we need to multiply by number of opt iterations
* change sampling parameters for prediction after training to defaults of common.h
and clarify what is context for prediction and what are generated tokens
* tighten abs error bounds for cross_entropy_loss in test-grad0
* add conditional compilation of using F16 exp in flash attention
uncomment `// #define GGML_FLASH_ATTN_EXP_FP16` to enable usage of f16 exp in flash attention
* tighten abs error bounds for flash_attn in test-grad0
* tighten abs error bounds for sqrt in test-grad0
* remove out-commented vectorized code of opt_adam
the vectorized code might be bit faster for low number of parameters, but it had a big memory usage overhead
* ggml : update ggml_rms_norm_back with configurable eps
* llama training : fix ggml_rms_norm_back calls to pass configurable eps
* remove trailing whitespace
* add train function using automatic gradient checkpointing backward pass and allocator
* in train function replace add_inplace by regular add
because using add_inplace seems to result in different gradients
* don't use allocate hash_map on context
because the context has no_alloc=True when using memory allocator resulting in NULL data pointers
* correctly clone reshape and permute operations by also cloning tensor->nb values
* fix variable name and add missing type cast
* terminate recursive tensor cloning when reaching tensor without src tensors
* correctly clone view tensors by setting data pointers
without this the checkpointing would only work when being used together with memory allocator
* fix variable names
* swap arguments to commutative ops to be the same as in `forward_batch_wo_cache_flash_attn`
* add input tensors as checkpoints
so that recursive tensor cloning of gradient checkpointing terminates on input tensors
* fix variable name and add missing boolean negation
* make sure some tensors are not reallocated by inserting new temporary nodes depending on them:
output and parameter gradient tensors need to be available at the end of the graph execution
parameter gradient tensors also need to be available before the graph execution because they are set to zero before each optimizer iteration
checkpoint tensors are allocated all together to reduce memory allocator fragmentation
afterwards, in addition to the temporary nodes, we also need to reset the temporary leafs
* fix ASSERT to work with zero layers
* add training options whether to use allocator and/or unified training function
* integrate unified training function which may use memory allocator
the unified training function also supports arguments whether to use flash attention and/or gradient checkpointing
* format name of cloned tensors with " (clone)" suffix
* set names for tensors in unified train function for easier debugging
* allocate graph on context using ggml_new_graph
* remove handwritten training functions
* remove unused training parameters "use_scratch" and "use_unified"
* remove trailing whitespace
* remove unused train params: mem_compute1_gb & mem_compute2_gb
mem_compute_gb is used for compute when automatic memory allocator is not enabled, otherwise it can be very small to only hold the tensor definitions
mem_compute0_gb is used for automatic memory allocator (as long as measurement of max required size is not implemented)
* remove unused forward_batch function
* add debug asserts in ggml_allocr_alloc to some common pitfalls when using this function directly
* only use ggml_allocr_alloc when tensor has NULL data and is no view
* fix test when to create temporary backward graph
temporary backward graph is only necessary when using checkpointing
* fix memory "leak" in optimizers
each iteration a new cplan with new memory for work data was allocated.
now cplan creation only happens at the start of optimization, with each iteration reusing the cplan and its work data.
* reverse order of for loop in ggml_build_backward_expand to save memory when using gradient checkpointing and allocator
with this loop order gradient checkpointing with allocator on 16 layer model saves 13% memory; 2 layer memory it saves 2% memory.
the computation results are the same
* add API functions to access llama model tensors
* add stub example for finetuning, based on train-text-from-scratch
* move and remove code
* add API functions to access remaining model parameters:
mult, head and rot
* first draft for LORA finetune training
* remove const model and layer arguments in API functions for accessing model tensors
* bug fixes to make finetune compile
automatic allocator does not work yet
* add debug prints for training memory improvements
* fix names of lora tensors
* avoid stack overflow resulting from big ggml_cgraph
replace stack allocation and ggml_build_forward by ggml_new_graph in combination with ggml_build_forward_expand
* replace llama API functions to get model tensors by one function to get model tensor by name
LLAMA_API struct ggml_tensor * llama_get_model_tensor(struct llama_model * model, const char * name);
* remove unused call to not existing llama_get_layer_from_model
* implement ggml_compute_forward_out_prod_q_f32
* remove trailing whitespace
* add lora finetune support on quantized base model tensors
* add ggml_add_cast API function
this function works like ggml_add, but accepts a data type for the resulting tensor.
only supported for quantized src0 input.
* use ggml_add_cast in finetuning
lora-applied weights will now have data type F32, which improves gradients when finetuning quantized base models
* bug fix: actually use result type passed to ggml_add_cast
* make sure base model tensors data cannot be used in viewable operations
memory allocator would try to make lora application inplace on base model tensors.
since those are memory mapped this will result in memory access violations
* fix bug in ggml_out_prod which resulted in wrong n_dims of result tensors
* avoid keeping in memory ALL of the gradients
The problem here stems from ggml_graph_reset. This function is called in the optimization function, before each graph computation, to reset the gradients to zero. This required a unique memory slot for each gradient: allocating memory from a previosly freed memory location might lead to non-zero input gradients.
During ggml_compute_backward the gradients are build stepwise by adding or substracting new values, starting from a OP_NONE tensor which needs to contain zero-values. This requires the graph reset.
To avoid this I now remember in ggml_build_backward_expand the original OP_NONE gradient tensors in a hash table, which is passed to ggml_compute_backward. There instead of using add (or sub or similar) I test whether the existing gradient to be changed is a zero-valued-tensor by looking up its existence in the hash table. When it is such a zero-tensor it will not be modified, but replaced by the value to be added, otherwise the regular add (not inplace, allocator will take care of this) will be used. This way none of those zero-tensor values will be necessary in the final backward graph and more importantly they won't need a unique memory slot, just to make them zero.
* remove trailing whitespace
* remove debug prints and function to compute tensor data hash
* improve optimization iteration prints
* adjust maximal values to support finetuning 3B models
* change default finetune params lora_r and lora_alpha to match the n_rank parameters of 4
* bug fix: make sure finetune input gradient is allocated at begin and kept until end
* remove unnecessary src tensor from ggml_get_rows_back
we don't need data of src[2] for computation, only to setup the correct output shape.
remove dependency on src[2], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included.
this is similar to how ggml_reshape does it.
* remove unnecessary src tensor from ggml_repeat & ggml_repeat_back
we don't need data of src[1] for computation, only to setup the correct output shape.
remove dependency on src[1], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included
* resolve todo
allocator will only make it inplace when they are of the same type
* mixing multiple LORA adapters is now possible
pass more than one '--lora FNAME' argument to apply more than one LORA.
use '--lora-scaled FNAME S' when you want to specify a user-defined scale for an adapter.
* add option to save finetune output every N iterations
* also save latest finetune output with ITERATION="LATEST" and print where files are saved
saving with LATEST makes it easier to resume training from the latest checkpoint
the string "LATEST" can be configured with command line option "--fn-latest STR"
* update checkpoint train stats before saving via "--save-every"
* add command line option `--rank-wo N` for rank of wo tensor
* update finetune README
* fix dump_non_result_info_yaml to output multiple lora adapters
* bug fix: replace GGML_TYPE_SIZE[t] by ggml_type_size(t)
* replace llama_n_mult by llama_n_ff
* finetune bug fixes to compile with merged in code from master
* remove prediction related code to reduce duplicated code with main
use main instead
* reduce large memory overhead in train-text-from-scratch
all gradients had to be pinned so that graph_reset works correctly.
this is no longer necessary with the changes to ggml_compute_backward introduced in this PR.
* add comment explaining why finetune checkpoints are allocated in one block
* make default value of float member a float literal
* handle rms_norm and rope parameters the same as in train-text-from-scratch
* remove unused code
* remove vocab related code as it is unnecessary
* add LLM_KV_TRAINING_TYPE to train-text-from-scratch checkpoints
so that they can be differentiated from lora finetune checkpoints
* add gguf constants and load/save functions from train-text-from-scratch
* add load & save lora finetune checkpoints via gguf
* add python script to convert old finetune checkpoint files to gguf
* remove old checkpoint save & load code
* remove code to print data checksums which was used to verify correctness of new gguf code
* omit tokenization when training is disabled, only save llama lora adapter
training can be disabled by passing '-n 0' to finetune
* remove trailing whitespace
* update README.md
* implement ggml_compute_forward_repeat_f16
* avoid stack overflow of large cgraphs in test-grad0
* add ggml API functions ggml_unravel_index, ggml_get_i32_nd and its analogs for set and for f32
ggml_get_i32_1d, ggml_set_i32_1d, ggml_get_f32_1d, ggml_set_f32_1d now support non-contiguous tensors.
in case of non-contiguous tensor, the 1d index is unraveled into a multi index using ggml_unravel_index to be passed to '_nd' function equivalent.
this fixes a bug in test-grad0 which happens due to ggml_build_backward not building purely contiguous tensors anymore
* increase test-grad0 context mem size to accommodate for bigger cgraph
* add sanity check to ggml_compute_backward, asserting the correct shape of gradients
* fix ggml_acc_or_set to return tensor of correct shape
* remove unused 'inplace' argument from ggml_compute_backward function
inplace operations to add gradients are no longer created by ggml_compute_backward
use allocator to automatically make inplace operations
* add missing argument 'int i0' to ggml_get_i32_nd & ggml_set_i32_nd header declarations
* fix error message in ggml_allocr_alloc to display actual max_avail
* fix check_gradient
ggml_build_backward_expand was previously replaced by ggml_build_backward, but the assignment of forward graph to backward graph missing
* use tensor->view_src instead of ggml_is_view and get_view_source
* move gradient checkpointing code into ggml, new API function:
// build gradient checkpointing backward graph gb for gf using provided checkpoints
// gb_tmp will contain original backward graph with rewritten backward process nodes,
// but without the second forward pass nodes.
GGML_API void ggml_build_backward_gradient_checkpointing(
struct ggml_context * ctx,
struct ggml_cgraph * gf,
struct ggml_cgraph * gb,
struct ggml_cgraph * gb_tmp,
struct ggml_tensor * * checkpoints,
int n_checkpoints);
* replace custom data getters and setters by ggml functions
* train-text-from-scratch can train (full finetune) gguf models
just pass the gguf model via `--checkpoint-in FN`.
after this, to continue training, pass the generated checkpoint instead of the original gguf model.
tested with smaller models, bigger models may exceed available memory.
use (LORA) finetune for those.
* remove trailing whitespace
* add option to save train-text-from-scratch output every N iterations
* update README.md
* fix warnings
* fix warnings
* remove finetune option to disable allocator
the allocator should always be used.
by making sure that it is always used it gets easier to implement automatic memory requirements computation
* add tensor checkpoints only when gradient checkpointing is enabled
* initialize opt ggml context if none was provided
* add ggml-alloc API function 'ggml_allocr_max_size' to get max size of alloc
GGML_API size_t ggml_allocr_max_size(struct ggml_allocr * alloc);
* finetune: automatically allocate all memory and changes to command line options
remove '--n_examples N' parameter, as it no longer makes sense to call optimization process multiple times in a loop.
add '--only_write_lora' command line option: will skip tokenization and training, to only write a llama.cpp comptabile LORA adapter.
remove memory buffer related command line options.
improve iteration console output.
* add finetune to Makefile
* update README.md
* print time per iteration and estimate remaining time
* increase measured alloc size by tensor_alignment
ggml_allocr_reset will reduce the given size by up to tensor_alignment-1
* fix README.md
* add some more allocator debug prints
* bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue
* revert last commit
"bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue"
"alloc was freeing an externally allocated tensor, because it calculated the end of allocator memory as alloc->data + alloc->max_size instead of alloc->data + alloc->size."
This is intentional to reduce the risk of freeing external tensors when measuring. Unless max_size is not properly calculated, I don't see why this is an issue.
* remove unnecessary "0x" before "%p" output
* move measurement memory segment to upper region of the address space
* update README.md
* fix printf format warnings
* add missing gguf_free in load_checkpoint_lora_file
* load default rms_norm and rope parameters from base model
* add gradient accumulation
specify number accumulation steps with '--grad-acc N'.
this will simulate a bigger batch size of grad_acc*batch.
* fix tracking of train_samples and train_tokens
* build : fix compile warnings
* ggml : fix L-BFGS linesearch loop
* improve finetune time measurement
fix printf warnings on system where int64_t is (long int).
change time datatypes to double because values get big with long training times.
exclude file saving from time measurement.
converge faster to actual time per iteration by removing very small first duration before first iteration was performed.
fix bug in output of total training time, the reported value was 1000 times to small.
* specify default lora rank with '--lora-r N'
'--lora-r N' will specify default rank for all tensors
'--rank-wq N', etc. will override this default rank for specific tensor types.
* fix gradient accumulation bug where the same batch was used for each microstep
* fix gradient accumulation bug where the same batch was used for each microstep
* support grouped-query-attention in ggml_flash_attn and ggml_flash_attn_back
k and v can now be repeated in q along ne[2]
in forward pass just use modulo to compute k and v indices, like ik2 = iq2 % nek2.
in backard pass this won't work as easy, because multiple threads will compete to accumulate to the same k->grad[:,ik1,ik2,ik3] and v->grad[:,iv1,iv2,iv3].
so we change the parallelization over q rows to be over k rows. this ensures non-overlapping (ik2,ik3) across threads.
in each thread we then iterate over the number of repetitions of k/v in q to compute iq2 as iq2 = ik2 + irep*nek2.
since ne2 is not the same for q,k and v we also change how the gradients are concatenated into the result tensor.
additionally the offsets of gradq, gradk and gradv in the result tensor are now memory aligned.
we also simplify the compute_backward part of flash_attn to use ggml_reshape instead of switching over the number of dimensions.
this needs a small change to ggml_reshape, removing the assertion of second argument to be contiguous.
since only the shape (ne) of the second reshape argument is of relevance, its memory layout (nb) is irrelevant -> it can very well be non-contiguous.
change test-grad0 to also test for repeated k/v in q.
this changes the rng and now results in small gradient differences in softmax. these solely come from using f16 exp table lookup in forward softmax: when temporarily changing softmax to use actual exp function, the reported gradient differences go away. gradient differences coming solely from f16 table lookup are acceptable.
added a note to explain this.
* add llama API functions to get grouped-query-attention n_head parameter 'n_head_kv'.
* fix finetune to support grouped-query-attention (using flash-attention)
note: ggml changes to ggml_out_prod are necessary to support grouped-query-attention without flash-attention.
* support broadcastable a in out_prod(a, b) and backward pass of broadcasting mul_mat(a, b)
* test broadcasting mul_mat backward pass
* decouple random number generator of each operation test
when changing one test the rng of others tests is not influenced anymore
* add comment briefly describing what ggml_repeat_back does
* simplify broadcasting mul_mat backward using ggml_repeat_back
* add cgraph evaluation order member and corresponding enum type
this controls in which order ggml_build_forward visits source nodes.
by default the nodes are visited left to right, i.e. src[0] first.
in some cases it is beneficial for ggml-alloc to visit in a different order.
two possible orders are supported: left-to-right (src[0] first) and right-to-left (src[0] last).
* measure max compute size for each cgraph eval order and use best order
this can bring huge memory savings:
e.g. codellama-34b with n_ctx=64, n_batch=1 goes from 92927.8mb down to 4627.6 MB
* remove unused command line options
* add sample start patterns and options to force new or by default resume last shuffling
* update shuffle rng state on reshuffle
* exclude known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* remove probably unnecessary exception type flags from stringstream
* pass correct max number of tokens to llama_tokenize
* account for possible leading whitespace that will be added by tokenizer
e.g. '\t' will be tokenized by llama spm tokenizer to [29871, 12]
* use unrolled vec_mad in out_prod
y is vec_mad result vec.
x is vec_mad input vec.
v is vec_mad input scalar.
ggml_vec_mad_f32_unroll will internally loop over x and v with same y.
GGML_VEC_MAD_UNROLL is by default defined to 32.
This value is empirical optimized using performance test runs of out-prod in openllama-3b finetune with 256 context length and batch size 1. It gives 23% performance boost for out_prod.
Full measurements of out-prod runtime in ms:
unroll_xv unroll_yv
1 67014.643 87826.469
2 77117.552 89077.656
4 72091.311 109121.657
8 61077.543 88678.334
16 56914.67 79514.947
24 59024.595 84350.254
28 55952.446 83368.73
32 51476.658 85177.745
36 55973.792 84659.92
40 55139.616 93844.738
48 60736.392 93330.267
64 99856.878 116994.99
Second column is when unrollying yv instead of xv
* set lora_alpha to value of lora_r if it is not set via command line
otherwise only changing lora_r will change scaling of lora adapter used in prediction
* reshuffle original sample order instead of the previous shuffled order
otherwise resumed reshuffle will not result in same sample order
* block tiling for out-prod inspired by mul-mat
block sizes are empirically optimized
roughly doubles the flops of out-prod
* exclude some more known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* add static keywords
* remove outcommented old code
* update train-text-from-scratch with tokenization, sample selection and shuffling from finetune
* remove lbfgs related train parameters
* move common train functions into common/train.[h|cpp]
* move train state into struct train_state
* move train data saving code into callback to unify code of opt_callback
train_params are still different in finetune and train-text-from-scratch, so it can't yet be moved to train.h|cpp
* move common train params into common/train
* move common opt_callback into common/train
* fix consume_common_train_arg
* save and load head_count_kv in lora checkpoints
* increase train_samples by used_samples instead of number of batches
on batch can contain more than one sample when option "fill_with_next_samples" is used
* fix usage of llama_tokenize
* remove static from process_escape since we need it exposed in header
* fix code formating of long function declarations
* fix condition in load_train_state_gguf
* use die("msg") instead of replace GGML_ASSERT(!"msg") or throw std::runtime_error("msg")
* fix saving and loading of training type
* remove terminating '\0' from tokenization
(llama_tokenize is now passed the string length instead of relying on terminating '\0')
* fix compile warnings
* fix compile warnings
* use new/delete for train_state instead of malloc/free
using malloc may result in seg faults when trying to assign string fields
* assert that sample_count > 0, avoiding division by zero
* fix frand to return value in interval [0,1)
* add train option "--sample-random-offsets"
Use samples beginning at random offsets.
The offset is only applied to the first sample in each batch context window.
Together with "--fill-with-next-samples" this may help for training endless text generation.
For example given a dataset containing samples "abcd", "ABCD", "0123".
With context size of 8 and options "--fill-with-next-samples", "--no-separate-with-eos", "--no-separate-with-bos",
the context windows of batches could only be filled with "abcdABCD", "ABCDabcd", "0123abcd", etc.
With "--sample-random-offsets" it can also be filled with "23abcdAB", "bcd0123A", etc.
* deduplicate code into function
* remove n_rot hparam, as it must always be hparam.n_embd_head()
* align code
* assert correct base model tensor shapes
* move some params from lora hparams into model hparams and load model params from gguf
this equalizes the model definition in finetune and text-from-scratch and removes the need for additional llama api functions to get model parameters
* remove now unnecessary llama API functions to get model params that where added by this PR
* train-text-from-scratch: automatically allocate model tensors, remove option '--mem-model N'
* train-text-from-scratch: automatically allocate opt context
* train-text-from-scratch: automatically allocate input tensors
* train-text-from-scratch: automatically allocate compute memory
* remove unused options and equalize train-text-from-scratch with finetune
* initialize opt->loss_after with zero
* add export-lora program
* remove trailing whitespace
* add export-lora build in Makefile
* remove unused struct tensor_info from export-lora
* add export-lora build dependency to llama
because it depends on common, which depends on llama
* update finetune README.md
* cancel optimization when specified number of epochs is completed
* improve handling of export-lora arguments
print errors and warnings when files could not be read or created
* Fix export-lora.cpp "not enough space in the context's memory pool" (#1)
* Fix export-lora.cpp "not enough space in the context's memory pool"
Without this patch, export-lora would sometimes error with "not enough space in the context's memory pool (needed 656784, available 656800)".
* increase required context size by 5*GGML_MEM_ALIGN instead of plain 16
---------
Co-authored-by: xaedes <xaedes@gmail.com>
* improve handling of not yet supported tensor types
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: meatbag-18a <145869052+meatbag-18a@users.noreply.github.com>
2023-09-28 18:40:11 +00:00
train.o : common /train .cpp common /train .h
$( CXX) $( CXXFLAGS) -c $< -o $@
2023-06-07 07:59:52 +00:00
libllama.so : llama .o ggml .o $( OBJS )
2023-05-01 16:23:47 +00:00
$( CXX) $( CXXFLAGS) -shared -fPIC -o $@ $^ $( LDFLAGS)
2024-02-01 15:18:53 +00:00
libllama.a : llama .o ggml .o $( OBJS ) $( COMMON_DEPS )
ar rcs libllama.a llama.o ggml.o $( OBJS) $( COMMON_DEPS)
2023-03-10 18:40:58 +00:00
clean :
2024-02-01 15:18:53 +00:00
rm -vrf *.o tests/*.o *.so *.a *.dll benchmark-matmult common/build-info.cpp *.dot $( COV_TARGETS) $( BUILD_TARGETS) $( TEST_TARGETS)
2024-02-05 18:33:00 +00:00
find examples pocs -type f -name "*.o" -delete
2023-03-10 18:40:58 +00:00
2023-05-01 16:23:47 +00:00
#
# Examples
#
2024-02-05 18:33:00 +00:00
# $< is the first prerequisite, i.e. the source file.
# Explicitly compile this to an object file so that it can be cached with ccache.
# The source file is then filtered out from $^ (the list of all prerequisites) and the object file is added instead.
# Helper function that replaces .c, .cpp, and .cu file endings with .o:
GET_OBJ_FILE = $( patsubst %.c,%.o,$( patsubst %.cpp,%.o,$( patsubst %.cu,%.o,$( 1) ) ) )
2023-11-02 06:50:16 +00:00
main : examples /main /main .cpp ggml .o llama .o $( COMMON_DEPS ) console .o grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-03-23 11:41:32 +00:00
@echo
@echo '==== Run ./main -h for help. ===='
@echo
2023-03-10 18:40:58 +00:00
2023-11-02 06:50:16 +00:00
infill : examples /infill /infill .cpp ggml .o llama .o $( COMMON_DEPS ) console .o grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-10-02 07:42:02 +00:00
2023-11-02 06:50:16 +00:00
simple : examples /simple /simple .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-06-16 18:58:09 +00:00
2023-11-18 21:48:17 +00:00
tokenize : examples /tokenize /tokenize .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-11-18 21:48:17 +00:00
2023-11-02 06:50:16 +00:00
batched : examples /batched /batched .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-09-28 16:04:36 +00:00
2023-11-02 06:50:16 +00:00
batched-bench : examples /batched -bench /batched -bench .cpp build -info .o ggml .o llama .o common .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-10-11 18:25:33 +00:00
2023-11-02 06:50:16 +00:00
quantize : examples /quantize /quantize .cpp build -info .o ggml .o llama .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-03-25 18:26:40 +00:00
2023-11-02 06:50:16 +00:00
quantize-stats : examples /quantize -stats /quantize -stats .cpp build -info .o ggml .o llama .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-04-07 22:09:18 +00:00
2023-11-02 06:50:16 +00:00
perplexity : examples /perplexity /perplexity .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2024-01-12 05:59:57 +00:00
imatrix : examples /imatrix /imatrix .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-03-10 18:40:58 +00:00
2023-11-02 06:50:16 +00:00
embedding : examples /embedding /embedding .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-03-28 06:11:09 +00:00
2024-02-29 15:50:41 +00:00
gritlm : examples /gritlm /gritlm .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-11-02 06:50:16 +00:00
save-load-state : examples /save -load -state /save -load -state .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-04-18 19:00:14 +00:00
2024-02-20 19:07:22 +00:00
server : examples /server /server .cpp examples /server /oai .hpp examples /server /utils .hpp examples /server /httplib .h examples /server /json .hpp examples /server /index .html .hpp examples /server /index .js .hpp examples /server /completion .js .hpp examples /llava /clip .cpp examples /llava /clip .h examples /llava /llava .h examples /llava /llava .cpp common /stb_image .h ggml .o llama .o $( COMMON_DEPS ) grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) -c examples/llava/clip.cpp -o $( call GET_OBJ_FILE, examples/llava/clip.cpp) -Wno-cast-qual
$( CXX) $( CXXFLAGS) -Iexamples/server $( filter-out %.h %.hpp $< examples/llava/clip.cpp,$^) $( call GET_OBJ_FILE, $<) $( call GET_OBJ_FILE, examples/llava/clip.cpp) -o $@ $( LDFLAGS) $( LWINSOCK2)
2023-05-27 17:04:14 +00:00
2023-12-21 21:07:58 +00:00
gguf : examples /gguf /gguf .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-21 20:07:43 +00:00
2023-10-11 19:35:46 +00:00
train-text-from-scratch : examples /train -text -from -scratch /train -text -from -scratch .cpp ggml .o llama .o $( COMMON_DEPS ) train .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-06-15 17:42:48 +00:00
2023-08-29 08:42:41 +00:00
convert-llama2c-to-ggml : examples /convert -llama 2c -to -ggml /convert -llama 2c -to -ggml .cpp ggml .o llama .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-11 23:17:25 +00:00
2023-11-02 06:50:16 +00:00
llama-bench : examples /llama -bench /llama -bench .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-18 10:44:58 +00:00
2023-11-06 21:36:23 +00:00
libllava.a : examples /llava /llava .cpp examples /llava /llava .h examples /llava /clip .cpp examples /llava /clip .h common /stb_image .h common /base 64.hpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2023-11-07 17:25:32 +00:00
$( CXX) $( CXXFLAGS) -static -fPIC -c $< -o $@ -Wno-cast-qual
2023-11-06 21:36:23 +00:00
llava-cli : examples /llava /llava -cli .cpp examples /llava /clip .h examples /llava /clip .cpp examples /llava /llava .h examples /llava /llava .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) -c examples/llava/clip.cpp -o $( call GET_OBJ_FILE, examples/llava/clip.cpp) -Wno-cast-qual
$( CXX) $( CXXFLAGS) -c examples/llava/llava.cpp -o $( call GET_OBJ_FILE, examples/llava/llava.cpp)
$( CXX) $( CXXFLAGS) $( filter-out %.h $< examples/llava/clip.cpp examples/llava/llava.cpp,$^) $( call GET_OBJ_FILE, $<) $( call GET_OBJ_FILE, examples/llava/clip.cpp) $( call GET_OBJ_FILE, examples/llava/llava.cpp) -o $@ $( LDFLAGS)
2023-10-12 15:23:18 +00:00
2023-10-11 19:35:46 +00:00
baby-llama : examples /baby -llama /baby -llama .cpp ggml .o llama .o $( COMMON_DEPS ) train .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-29 08:42:41 +00:00
2023-11-02 06:50:16 +00:00
beam-search : examples /beam -search /beam -search .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-29 08:42:41 +00:00
2023-11-02 06:50:16 +00:00
finetune : examples /finetune /finetune .cpp ggml .o llama .o $( COMMON_DEPS ) train .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
train : finetune LORA (#2632)
* fix track_max_mem in forward_batch_wo_cache_flash_attn_train
* remove unnecessary Adam(W) optimizer tensors.
reduces optimizer memory overhead from 7*modelsize to 2*modelsize.
additionally allows to optimize models with more than 2^31 parameters by replacing int with int64_t.
bumps training checkpoint file version, but old checkpoints can still be read.
new version with less tensors is saved.
* add gradient clipping to AdamW
* Fix reset of unused g->nodes and g->grads to NULL
* implement gradient checkpointing for training
reduces memory overhead from O(n_layer) to O(sqrt(n_layer))
as explained in readme of https://github.com/cybertronai/gradient-checkpointing
* remove unused compute buffer 3
* add and use function ggml_build_backward_expand to avoid stack overflows with large maximum number of nodes
GGML_API void ggml_build_backward_expand(struct ggml_context * ctx, struct ggml_cgraph * gf, struct ggml_cgraph * gb, bool keep);
* change AdamW decay parameter to work like the torch AdamW decay parameter
It is now relative to Adam learning rate `alpha*sched`.
Before that it was relative to `sched` only.
`alpha` being the maximum learning rate and `sched` being a scaling parameter in [0..1]
* change default AdamW weight decay parameter used in training to 0.1 as used in nanoGPT
* change default AdamW weight decay parameter defined in ggml to 0.0, making Adam default instead of AdamW
btw: the default weight decay parameter for torch.optim.AdamW is 0.01
* bug fixes for cross entropy loss
ggml_cross_entropy_loss: sums where not correctly added in workload of each thread
ggml_cross_entropy_loss_back: simplify backward process, reducing numerical issues
guard usage of exp f16 lookup in cross entropy by #define GGML_CROSS_ENTROPY_EXP_FP16
cross entropy loss is only used once during training, but it is quite sensitive to numerical errors introduced by exp-f16-lookup.
so exp-f16-lookup for cross entropy loss is disabled by default, trading better gradients for very slightly worse runtime performance.
* fix test-grad0 for cross_entropy_loss
the second argument to cross_entropy_loss must sum up to 1 for each row
* fix test-grad0 for soft_max
dont use only sum as aggregation, because sum of softmax is always 1 -> finite differences should not work
instead use sum(log(soft_max()*(1-eps)+eps)); use eps to avoid log(0)
* improve finite differences of test-grad0 by using double instead of float
* change cross_entropy_loss to output average over all rows
this helps keeping the loss and gradients in a sane range
* improve gradient checkpointing
sqrt(n_layers) is only the best checkpoint step when mem size of checkpoints and mem size of layers are equal.
since layers require more memory than the single-tensor-checkpoint we use, the optimal values are compute different:
```
given: n, u, v
objective: minimize(a*u+b*v) where a*b=n, a>0, b>0
b=n/a
minimize(a*u+v*n/a)
diff(a*u+v*n/a, a) = u - (v*n/a)/a
diff(a*u+v*n/a, a) == 0
u - (v*n/a)/a == 0
u == v*n/(a*a)
u*a*a = v*n
a*a = v*n/u
a = sqrt(n*v/u)
```
this change results in more checkpoints, requiring less layers to store between checkpoints, overall improving memory usage.
* disable gradient checkpointing debug output
* llama : fix rope usage in train-text-from-scratch after ChatGLM change
* add more training parameters:
--enable-restart N Only for Adam optimizer. Enable restarts of cos-decay
--disable-restart N Only for Adam optimizer. Disable restarts of cos-decay
--opt-past N Number of optimization iterations to track for delta convergence test. Disabled when zero.
--opt-delta N Maximum delta for delta convergence test. Disabled when <= zero.
--opt-max-no-improvement N Maximum number of optimization iterations with no improvement. Disabled when <= zero.
--adam-epsf N AdamW epsilon for convergence test. Disabled when <= zero.
--adam-min-alpha N Adam minimum learning rate alpha, usually 0.1 * alpha
* replace memcpy with reshape operation so that the graph is not cut at the input
this makes it possible to store other values into the input tensor and then simply recompute the graph without rebuilding it
* remove unused function argument from get_example_targets_batch
* measure and print total training time
* add optimization callback to ggml_opt_resume_g
this callback is called before each iteration with custom data and pointer to learning schedule parameter (only used in Adam(W)).
can be used for dynamic learning schedule and setting input data for batches before each iteration
* use optimization callback in training
allows dynamic learning schedule and different batch data for each iteration without relying on low n_iter and high n_examples parameters
reduces runtime by avoiding restart of optimization function and improves training convergence by providing a different batch for each iteration
* add minimum number of tensor dimensions to apply weight decay (default 2)
this allows to not apply weight decay to bias parameters
* rename training parameter cos-decay-alpha to cos-decay-min and clarify that adam-min-alpha also applies to warmup
* fix increase of model.train_samples and model.train_tokens
now that each optimizer iteration gets its own batch we need to multiply by number of opt iterations
* change sampling parameters for prediction after training to defaults of common.h
and clarify what is context for prediction and what are generated tokens
* tighten abs error bounds for cross_entropy_loss in test-grad0
* add conditional compilation of using F16 exp in flash attention
uncomment `// #define GGML_FLASH_ATTN_EXP_FP16` to enable usage of f16 exp in flash attention
* tighten abs error bounds for flash_attn in test-grad0
* tighten abs error bounds for sqrt in test-grad0
* remove out-commented vectorized code of opt_adam
the vectorized code might be bit faster for low number of parameters, but it had a big memory usage overhead
* ggml : update ggml_rms_norm_back with configurable eps
* llama training : fix ggml_rms_norm_back calls to pass configurable eps
* remove trailing whitespace
* add train function using automatic gradient checkpointing backward pass and allocator
* in train function replace add_inplace by regular add
because using add_inplace seems to result in different gradients
* don't use allocate hash_map on context
because the context has no_alloc=True when using memory allocator resulting in NULL data pointers
* correctly clone reshape and permute operations by also cloning tensor->nb values
* fix variable name and add missing type cast
* terminate recursive tensor cloning when reaching tensor without src tensors
* correctly clone view tensors by setting data pointers
without this the checkpointing would only work when being used together with memory allocator
* fix variable names
* swap arguments to commutative ops to be the same as in `forward_batch_wo_cache_flash_attn`
* add input tensors as checkpoints
so that recursive tensor cloning of gradient checkpointing terminates on input tensors
* fix variable name and add missing boolean negation
* make sure some tensors are not reallocated by inserting new temporary nodes depending on them:
output and parameter gradient tensors need to be available at the end of the graph execution
parameter gradient tensors also need to be available before the graph execution because they are set to zero before each optimizer iteration
checkpoint tensors are allocated all together to reduce memory allocator fragmentation
afterwards, in addition to the temporary nodes, we also need to reset the temporary leafs
* fix ASSERT to work with zero layers
* add training options whether to use allocator and/or unified training function
* integrate unified training function which may use memory allocator
the unified training function also supports arguments whether to use flash attention and/or gradient checkpointing
* format name of cloned tensors with " (clone)" suffix
* set names for tensors in unified train function for easier debugging
* allocate graph on context using ggml_new_graph
* remove handwritten training functions
* remove unused training parameters "use_scratch" and "use_unified"
* remove trailing whitespace
* remove unused train params: mem_compute1_gb & mem_compute2_gb
mem_compute_gb is used for compute when automatic memory allocator is not enabled, otherwise it can be very small to only hold the tensor definitions
mem_compute0_gb is used for automatic memory allocator (as long as measurement of max required size is not implemented)
* remove unused forward_batch function
* add debug asserts in ggml_allocr_alloc to some common pitfalls when using this function directly
* only use ggml_allocr_alloc when tensor has NULL data and is no view
* fix test when to create temporary backward graph
temporary backward graph is only necessary when using checkpointing
* fix memory "leak" in optimizers
each iteration a new cplan with new memory for work data was allocated.
now cplan creation only happens at the start of optimization, with each iteration reusing the cplan and its work data.
* reverse order of for loop in ggml_build_backward_expand to save memory when using gradient checkpointing and allocator
with this loop order gradient checkpointing with allocator on 16 layer model saves 13% memory; 2 layer memory it saves 2% memory.
the computation results are the same
* add API functions to access llama model tensors
* add stub example for finetuning, based on train-text-from-scratch
* move and remove code
* add API functions to access remaining model parameters:
mult, head and rot
* first draft for LORA finetune training
* remove const model and layer arguments in API functions for accessing model tensors
* bug fixes to make finetune compile
automatic allocator does not work yet
* add debug prints for training memory improvements
* fix names of lora tensors
* avoid stack overflow resulting from big ggml_cgraph
replace stack allocation and ggml_build_forward by ggml_new_graph in combination with ggml_build_forward_expand
* replace llama API functions to get model tensors by one function to get model tensor by name
LLAMA_API struct ggml_tensor * llama_get_model_tensor(struct llama_model * model, const char * name);
* remove unused call to not existing llama_get_layer_from_model
* implement ggml_compute_forward_out_prod_q_f32
* remove trailing whitespace
* add lora finetune support on quantized base model tensors
* add ggml_add_cast API function
this function works like ggml_add, but accepts a data type for the resulting tensor.
only supported for quantized src0 input.
* use ggml_add_cast in finetuning
lora-applied weights will now have data type F32, which improves gradients when finetuning quantized base models
* bug fix: actually use result type passed to ggml_add_cast
* make sure base model tensors data cannot be used in viewable operations
memory allocator would try to make lora application inplace on base model tensors.
since those are memory mapped this will result in memory access violations
* fix bug in ggml_out_prod which resulted in wrong n_dims of result tensors
* avoid keeping in memory ALL of the gradients
The problem here stems from ggml_graph_reset. This function is called in the optimization function, before each graph computation, to reset the gradients to zero. This required a unique memory slot for each gradient: allocating memory from a previosly freed memory location might lead to non-zero input gradients.
During ggml_compute_backward the gradients are build stepwise by adding or substracting new values, starting from a OP_NONE tensor which needs to contain zero-values. This requires the graph reset.
To avoid this I now remember in ggml_build_backward_expand the original OP_NONE gradient tensors in a hash table, which is passed to ggml_compute_backward. There instead of using add (or sub or similar) I test whether the existing gradient to be changed is a zero-valued-tensor by looking up its existence in the hash table. When it is such a zero-tensor it will not be modified, but replaced by the value to be added, otherwise the regular add (not inplace, allocator will take care of this) will be used. This way none of those zero-tensor values will be necessary in the final backward graph and more importantly they won't need a unique memory slot, just to make them zero.
* remove trailing whitespace
* remove debug prints and function to compute tensor data hash
* improve optimization iteration prints
* adjust maximal values to support finetuning 3B models
* change default finetune params lora_r and lora_alpha to match the n_rank parameters of 4
* bug fix: make sure finetune input gradient is allocated at begin and kept until end
* remove unnecessary src tensor from ggml_get_rows_back
we don't need data of src[2] for computation, only to setup the correct output shape.
remove dependency on src[2], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included.
this is similar to how ggml_reshape does it.
* remove unnecessary src tensor from ggml_repeat & ggml_repeat_back
we don't need data of src[1] for computation, only to setup the correct output shape.
remove dependency on src[1], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included
* resolve todo
allocator will only make it inplace when they are of the same type
* mixing multiple LORA adapters is now possible
pass more than one '--lora FNAME' argument to apply more than one LORA.
use '--lora-scaled FNAME S' when you want to specify a user-defined scale for an adapter.
* add option to save finetune output every N iterations
* also save latest finetune output with ITERATION="LATEST" and print where files are saved
saving with LATEST makes it easier to resume training from the latest checkpoint
the string "LATEST" can be configured with command line option "--fn-latest STR"
* update checkpoint train stats before saving via "--save-every"
* add command line option `--rank-wo N` for rank of wo tensor
* update finetune README
* fix dump_non_result_info_yaml to output multiple lora adapters
* bug fix: replace GGML_TYPE_SIZE[t] by ggml_type_size(t)
* replace llama_n_mult by llama_n_ff
* finetune bug fixes to compile with merged in code from master
* remove prediction related code to reduce duplicated code with main
use main instead
* reduce large memory overhead in train-text-from-scratch
all gradients had to be pinned so that graph_reset works correctly.
this is no longer necessary with the changes to ggml_compute_backward introduced in this PR.
* add comment explaining why finetune checkpoints are allocated in one block
* make default value of float member a float literal
* handle rms_norm and rope parameters the same as in train-text-from-scratch
* remove unused code
* remove vocab related code as it is unnecessary
* add LLM_KV_TRAINING_TYPE to train-text-from-scratch checkpoints
so that they can be differentiated from lora finetune checkpoints
* add gguf constants and load/save functions from train-text-from-scratch
* add load & save lora finetune checkpoints via gguf
* add python script to convert old finetune checkpoint files to gguf
* remove old checkpoint save & load code
* remove code to print data checksums which was used to verify correctness of new gguf code
* omit tokenization when training is disabled, only save llama lora adapter
training can be disabled by passing '-n 0' to finetune
* remove trailing whitespace
* update README.md
* implement ggml_compute_forward_repeat_f16
* avoid stack overflow of large cgraphs in test-grad0
* add ggml API functions ggml_unravel_index, ggml_get_i32_nd and its analogs for set and for f32
ggml_get_i32_1d, ggml_set_i32_1d, ggml_get_f32_1d, ggml_set_f32_1d now support non-contiguous tensors.
in case of non-contiguous tensor, the 1d index is unraveled into a multi index using ggml_unravel_index to be passed to '_nd' function equivalent.
this fixes a bug in test-grad0 which happens due to ggml_build_backward not building purely contiguous tensors anymore
* increase test-grad0 context mem size to accommodate for bigger cgraph
* add sanity check to ggml_compute_backward, asserting the correct shape of gradients
* fix ggml_acc_or_set to return tensor of correct shape
* remove unused 'inplace' argument from ggml_compute_backward function
inplace operations to add gradients are no longer created by ggml_compute_backward
use allocator to automatically make inplace operations
* add missing argument 'int i0' to ggml_get_i32_nd & ggml_set_i32_nd header declarations
* fix error message in ggml_allocr_alloc to display actual max_avail
* fix check_gradient
ggml_build_backward_expand was previously replaced by ggml_build_backward, but the assignment of forward graph to backward graph missing
* use tensor->view_src instead of ggml_is_view and get_view_source
* move gradient checkpointing code into ggml, new API function:
// build gradient checkpointing backward graph gb for gf using provided checkpoints
// gb_tmp will contain original backward graph with rewritten backward process nodes,
// but without the second forward pass nodes.
GGML_API void ggml_build_backward_gradient_checkpointing(
struct ggml_context * ctx,
struct ggml_cgraph * gf,
struct ggml_cgraph * gb,
struct ggml_cgraph * gb_tmp,
struct ggml_tensor * * checkpoints,
int n_checkpoints);
* replace custom data getters and setters by ggml functions
* train-text-from-scratch can train (full finetune) gguf models
just pass the gguf model via `--checkpoint-in FN`.
after this, to continue training, pass the generated checkpoint instead of the original gguf model.
tested with smaller models, bigger models may exceed available memory.
use (LORA) finetune for those.
* remove trailing whitespace
* add option to save train-text-from-scratch output every N iterations
* update README.md
* fix warnings
* fix warnings
* remove finetune option to disable allocator
the allocator should always be used.
by making sure that it is always used it gets easier to implement automatic memory requirements computation
* add tensor checkpoints only when gradient checkpointing is enabled
* initialize opt ggml context if none was provided
* add ggml-alloc API function 'ggml_allocr_max_size' to get max size of alloc
GGML_API size_t ggml_allocr_max_size(struct ggml_allocr * alloc);
* finetune: automatically allocate all memory and changes to command line options
remove '--n_examples N' parameter, as it no longer makes sense to call optimization process multiple times in a loop.
add '--only_write_lora' command line option: will skip tokenization and training, to only write a llama.cpp comptabile LORA adapter.
remove memory buffer related command line options.
improve iteration console output.
* add finetune to Makefile
* update README.md
* print time per iteration and estimate remaining time
* increase measured alloc size by tensor_alignment
ggml_allocr_reset will reduce the given size by up to tensor_alignment-1
* fix README.md
* add some more allocator debug prints
* bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue
* revert last commit
"bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue"
"alloc was freeing an externally allocated tensor, because it calculated the end of allocator memory as alloc->data + alloc->max_size instead of alloc->data + alloc->size."
This is intentional to reduce the risk of freeing external tensors when measuring. Unless max_size is not properly calculated, I don't see why this is an issue.
* remove unnecessary "0x" before "%p" output
* move measurement memory segment to upper region of the address space
* update README.md
* fix printf format warnings
* add missing gguf_free in load_checkpoint_lora_file
* load default rms_norm and rope parameters from base model
* add gradient accumulation
specify number accumulation steps with '--grad-acc N'.
this will simulate a bigger batch size of grad_acc*batch.
* fix tracking of train_samples and train_tokens
* build : fix compile warnings
* ggml : fix L-BFGS linesearch loop
* improve finetune time measurement
fix printf warnings on system where int64_t is (long int).
change time datatypes to double because values get big with long training times.
exclude file saving from time measurement.
converge faster to actual time per iteration by removing very small first duration before first iteration was performed.
fix bug in output of total training time, the reported value was 1000 times to small.
* specify default lora rank with '--lora-r N'
'--lora-r N' will specify default rank for all tensors
'--rank-wq N', etc. will override this default rank for specific tensor types.
* fix gradient accumulation bug where the same batch was used for each microstep
* fix gradient accumulation bug where the same batch was used for each microstep
* support grouped-query-attention in ggml_flash_attn and ggml_flash_attn_back
k and v can now be repeated in q along ne[2]
in forward pass just use modulo to compute k and v indices, like ik2 = iq2 % nek2.
in backard pass this won't work as easy, because multiple threads will compete to accumulate to the same k->grad[:,ik1,ik2,ik3] and v->grad[:,iv1,iv2,iv3].
so we change the parallelization over q rows to be over k rows. this ensures non-overlapping (ik2,ik3) across threads.
in each thread we then iterate over the number of repetitions of k/v in q to compute iq2 as iq2 = ik2 + irep*nek2.
since ne2 is not the same for q,k and v we also change how the gradients are concatenated into the result tensor.
additionally the offsets of gradq, gradk and gradv in the result tensor are now memory aligned.
we also simplify the compute_backward part of flash_attn to use ggml_reshape instead of switching over the number of dimensions.
this needs a small change to ggml_reshape, removing the assertion of second argument to be contiguous.
since only the shape (ne) of the second reshape argument is of relevance, its memory layout (nb) is irrelevant -> it can very well be non-contiguous.
change test-grad0 to also test for repeated k/v in q.
this changes the rng and now results in small gradient differences in softmax. these solely come from using f16 exp table lookup in forward softmax: when temporarily changing softmax to use actual exp function, the reported gradient differences go away. gradient differences coming solely from f16 table lookup are acceptable.
added a note to explain this.
* add llama API functions to get grouped-query-attention n_head parameter 'n_head_kv'.
* fix finetune to support grouped-query-attention (using flash-attention)
note: ggml changes to ggml_out_prod are necessary to support grouped-query-attention without flash-attention.
* support broadcastable a in out_prod(a, b) and backward pass of broadcasting mul_mat(a, b)
* test broadcasting mul_mat backward pass
* decouple random number generator of each operation test
when changing one test the rng of others tests is not influenced anymore
* add comment briefly describing what ggml_repeat_back does
* simplify broadcasting mul_mat backward using ggml_repeat_back
* add cgraph evaluation order member and corresponding enum type
this controls in which order ggml_build_forward visits source nodes.
by default the nodes are visited left to right, i.e. src[0] first.
in some cases it is beneficial for ggml-alloc to visit in a different order.
two possible orders are supported: left-to-right (src[0] first) and right-to-left (src[0] last).
* measure max compute size for each cgraph eval order and use best order
this can bring huge memory savings:
e.g. codellama-34b with n_ctx=64, n_batch=1 goes from 92927.8mb down to 4627.6 MB
* remove unused command line options
* add sample start patterns and options to force new or by default resume last shuffling
* update shuffle rng state on reshuffle
* exclude known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* remove probably unnecessary exception type flags from stringstream
* pass correct max number of tokens to llama_tokenize
* account for possible leading whitespace that will be added by tokenizer
e.g. '\t' will be tokenized by llama spm tokenizer to [29871, 12]
* use unrolled vec_mad in out_prod
y is vec_mad result vec.
x is vec_mad input vec.
v is vec_mad input scalar.
ggml_vec_mad_f32_unroll will internally loop over x and v with same y.
GGML_VEC_MAD_UNROLL is by default defined to 32.
This value is empirical optimized using performance test runs of out-prod in openllama-3b finetune with 256 context length and batch size 1. It gives 23% performance boost for out_prod.
Full measurements of out-prod runtime in ms:
unroll_xv unroll_yv
1 67014.643 87826.469
2 77117.552 89077.656
4 72091.311 109121.657
8 61077.543 88678.334
16 56914.67 79514.947
24 59024.595 84350.254
28 55952.446 83368.73
32 51476.658 85177.745
36 55973.792 84659.92
40 55139.616 93844.738
48 60736.392 93330.267
64 99856.878 116994.99
Second column is when unrollying yv instead of xv
* set lora_alpha to value of lora_r if it is not set via command line
otherwise only changing lora_r will change scaling of lora adapter used in prediction
* reshuffle original sample order instead of the previous shuffled order
otherwise resumed reshuffle will not result in same sample order
* block tiling for out-prod inspired by mul-mat
block sizes are empirically optimized
roughly doubles the flops of out-prod
* exclude some more known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* add static keywords
* remove outcommented old code
* update train-text-from-scratch with tokenization, sample selection and shuffling from finetune
* remove lbfgs related train parameters
* move common train functions into common/train.[h|cpp]
* move train state into struct train_state
* move train data saving code into callback to unify code of opt_callback
train_params are still different in finetune and train-text-from-scratch, so it can't yet be moved to train.h|cpp
* move common train params into common/train
* move common opt_callback into common/train
* fix consume_common_train_arg
* save and load head_count_kv in lora checkpoints
* increase train_samples by used_samples instead of number of batches
on batch can contain more than one sample when option "fill_with_next_samples" is used
* fix usage of llama_tokenize
* remove static from process_escape since we need it exposed in header
* fix code formating of long function declarations
* fix condition in load_train_state_gguf
* use die("msg") instead of replace GGML_ASSERT(!"msg") or throw std::runtime_error("msg")
* fix saving and loading of training type
* remove terminating '\0' from tokenization
(llama_tokenize is now passed the string length instead of relying on terminating '\0')
* fix compile warnings
* fix compile warnings
* use new/delete for train_state instead of malloc/free
using malloc may result in seg faults when trying to assign string fields
* assert that sample_count > 0, avoiding division by zero
* fix frand to return value in interval [0,1)
* add train option "--sample-random-offsets"
Use samples beginning at random offsets.
The offset is only applied to the first sample in each batch context window.
Together with "--fill-with-next-samples" this may help for training endless text generation.
For example given a dataset containing samples "abcd", "ABCD", "0123".
With context size of 8 and options "--fill-with-next-samples", "--no-separate-with-eos", "--no-separate-with-bos",
the context windows of batches could only be filled with "abcdABCD", "ABCDabcd", "0123abcd", etc.
With "--sample-random-offsets" it can also be filled with "23abcdAB", "bcd0123A", etc.
* deduplicate code into function
* remove n_rot hparam, as it must always be hparam.n_embd_head()
* align code
* assert correct base model tensor shapes
* move some params from lora hparams into model hparams and load model params from gguf
this equalizes the model definition in finetune and text-from-scratch and removes the need for additional llama api functions to get model parameters
* remove now unnecessary llama API functions to get model params that where added by this PR
* train-text-from-scratch: automatically allocate model tensors, remove option '--mem-model N'
* train-text-from-scratch: automatically allocate opt context
* train-text-from-scratch: automatically allocate input tensors
* train-text-from-scratch: automatically allocate compute memory
* remove unused options and equalize train-text-from-scratch with finetune
* initialize opt->loss_after with zero
* add export-lora program
* remove trailing whitespace
* add export-lora build in Makefile
* remove unused struct tensor_info from export-lora
* add export-lora build dependency to llama
because it depends on common, which depends on llama
* update finetune README.md
* cancel optimization when specified number of epochs is completed
* improve handling of export-lora arguments
print errors and warnings when files could not be read or created
* Fix export-lora.cpp "not enough space in the context's memory pool" (#1)
* Fix export-lora.cpp "not enough space in the context's memory pool"
Without this patch, export-lora would sometimes error with "not enough space in the context's memory pool (needed 656784, available 656800)".
* increase required context size by 5*GGML_MEM_ALIGN instead of plain 16
---------
Co-authored-by: xaedes <xaedes@gmail.com>
* improve handling of not yet supported tensor types
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: meatbag-18a <145869052+meatbag-18a@users.noreply.github.com>
2023-09-28 18:40:11 +00:00
2023-11-30 22:23:08 +00:00
export-lora : examples /export -lora /export -lora .cpp ggml .o common /common .h $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
train : finetune LORA (#2632)
* fix track_max_mem in forward_batch_wo_cache_flash_attn_train
* remove unnecessary Adam(W) optimizer tensors.
reduces optimizer memory overhead from 7*modelsize to 2*modelsize.
additionally allows to optimize models with more than 2^31 parameters by replacing int with int64_t.
bumps training checkpoint file version, but old checkpoints can still be read.
new version with less tensors is saved.
* add gradient clipping to AdamW
* Fix reset of unused g->nodes and g->grads to NULL
* implement gradient checkpointing for training
reduces memory overhead from O(n_layer) to O(sqrt(n_layer))
as explained in readme of https://github.com/cybertronai/gradient-checkpointing
* remove unused compute buffer 3
* add and use function ggml_build_backward_expand to avoid stack overflows with large maximum number of nodes
GGML_API void ggml_build_backward_expand(struct ggml_context * ctx, struct ggml_cgraph * gf, struct ggml_cgraph * gb, bool keep);
* change AdamW decay parameter to work like the torch AdamW decay parameter
It is now relative to Adam learning rate `alpha*sched`.
Before that it was relative to `sched` only.
`alpha` being the maximum learning rate and `sched` being a scaling parameter in [0..1]
* change default AdamW weight decay parameter used in training to 0.1 as used in nanoGPT
* change default AdamW weight decay parameter defined in ggml to 0.0, making Adam default instead of AdamW
btw: the default weight decay parameter for torch.optim.AdamW is 0.01
* bug fixes for cross entropy loss
ggml_cross_entropy_loss: sums where not correctly added in workload of each thread
ggml_cross_entropy_loss_back: simplify backward process, reducing numerical issues
guard usage of exp f16 lookup in cross entropy by #define GGML_CROSS_ENTROPY_EXP_FP16
cross entropy loss is only used once during training, but it is quite sensitive to numerical errors introduced by exp-f16-lookup.
so exp-f16-lookup for cross entropy loss is disabled by default, trading better gradients for very slightly worse runtime performance.
* fix test-grad0 for cross_entropy_loss
the second argument to cross_entropy_loss must sum up to 1 for each row
* fix test-grad0 for soft_max
dont use only sum as aggregation, because sum of softmax is always 1 -> finite differences should not work
instead use sum(log(soft_max()*(1-eps)+eps)); use eps to avoid log(0)
* improve finite differences of test-grad0 by using double instead of float
* change cross_entropy_loss to output average over all rows
this helps keeping the loss and gradients in a sane range
* improve gradient checkpointing
sqrt(n_layers) is only the best checkpoint step when mem size of checkpoints and mem size of layers are equal.
since layers require more memory than the single-tensor-checkpoint we use, the optimal values are compute different:
```
given: n, u, v
objective: minimize(a*u+b*v) where a*b=n, a>0, b>0
b=n/a
minimize(a*u+v*n/a)
diff(a*u+v*n/a, a) = u - (v*n/a)/a
diff(a*u+v*n/a, a) == 0
u - (v*n/a)/a == 0
u == v*n/(a*a)
u*a*a = v*n
a*a = v*n/u
a = sqrt(n*v/u)
```
this change results in more checkpoints, requiring less layers to store between checkpoints, overall improving memory usage.
* disable gradient checkpointing debug output
* llama : fix rope usage in train-text-from-scratch after ChatGLM change
* add more training parameters:
--enable-restart N Only for Adam optimizer. Enable restarts of cos-decay
--disable-restart N Only for Adam optimizer. Disable restarts of cos-decay
--opt-past N Number of optimization iterations to track for delta convergence test. Disabled when zero.
--opt-delta N Maximum delta for delta convergence test. Disabled when <= zero.
--opt-max-no-improvement N Maximum number of optimization iterations with no improvement. Disabled when <= zero.
--adam-epsf N AdamW epsilon for convergence test. Disabled when <= zero.
--adam-min-alpha N Adam minimum learning rate alpha, usually 0.1 * alpha
* replace memcpy with reshape operation so that the graph is not cut at the input
this makes it possible to store other values into the input tensor and then simply recompute the graph without rebuilding it
* remove unused function argument from get_example_targets_batch
* measure and print total training time
* add optimization callback to ggml_opt_resume_g
this callback is called before each iteration with custom data and pointer to learning schedule parameter (only used in Adam(W)).
can be used for dynamic learning schedule and setting input data for batches before each iteration
* use optimization callback in training
allows dynamic learning schedule and different batch data for each iteration without relying on low n_iter and high n_examples parameters
reduces runtime by avoiding restart of optimization function and improves training convergence by providing a different batch for each iteration
* add minimum number of tensor dimensions to apply weight decay (default 2)
this allows to not apply weight decay to bias parameters
* rename training parameter cos-decay-alpha to cos-decay-min and clarify that adam-min-alpha also applies to warmup
* fix increase of model.train_samples and model.train_tokens
now that each optimizer iteration gets its own batch we need to multiply by number of opt iterations
* change sampling parameters for prediction after training to defaults of common.h
and clarify what is context for prediction and what are generated tokens
* tighten abs error bounds for cross_entropy_loss in test-grad0
* add conditional compilation of using F16 exp in flash attention
uncomment `// #define GGML_FLASH_ATTN_EXP_FP16` to enable usage of f16 exp in flash attention
* tighten abs error bounds for flash_attn in test-grad0
* tighten abs error bounds for sqrt in test-grad0
* remove out-commented vectorized code of opt_adam
the vectorized code might be bit faster for low number of parameters, but it had a big memory usage overhead
* ggml : update ggml_rms_norm_back with configurable eps
* llama training : fix ggml_rms_norm_back calls to pass configurable eps
* remove trailing whitespace
* add train function using automatic gradient checkpointing backward pass and allocator
* in train function replace add_inplace by regular add
because using add_inplace seems to result in different gradients
* don't use allocate hash_map on context
because the context has no_alloc=True when using memory allocator resulting in NULL data pointers
* correctly clone reshape and permute operations by also cloning tensor->nb values
* fix variable name and add missing type cast
* terminate recursive tensor cloning when reaching tensor without src tensors
* correctly clone view tensors by setting data pointers
without this the checkpointing would only work when being used together with memory allocator
* fix variable names
* swap arguments to commutative ops to be the same as in `forward_batch_wo_cache_flash_attn`
* add input tensors as checkpoints
so that recursive tensor cloning of gradient checkpointing terminates on input tensors
* fix variable name and add missing boolean negation
* make sure some tensors are not reallocated by inserting new temporary nodes depending on them:
output and parameter gradient tensors need to be available at the end of the graph execution
parameter gradient tensors also need to be available before the graph execution because they are set to zero before each optimizer iteration
checkpoint tensors are allocated all together to reduce memory allocator fragmentation
afterwards, in addition to the temporary nodes, we also need to reset the temporary leafs
* fix ASSERT to work with zero layers
* add training options whether to use allocator and/or unified training function
* integrate unified training function which may use memory allocator
the unified training function also supports arguments whether to use flash attention and/or gradient checkpointing
* format name of cloned tensors with " (clone)" suffix
* set names for tensors in unified train function for easier debugging
* allocate graph on context using ggml_new_graph
* remove handwritten training functions
* remove unused training parameters "use_scratch" and "use_unified"
* remove trailing whitespace
* remove unused train params: mem_compute1_gb & mem_compute2_gb
mem_compute_gb is used for compute when automatic memory allocator is not enabled, otherwise it can be very small to only hold the tensor definitions
mem_compute0_gb is used for automatic memory allocator (as long as measurement of max required size is not implemented)
* remove unused forward_batch function
* add debug asserts in ggml_allocr_alloc to some common pitfalls when using this function directly
* only use ggml_allocr_alloc when tensor has NULL data and is no view
* fix test when to create temporary backward graph
temporary backward graph is only necessary when using checkpointing
* fix memory "leak" in optimizers
each iteration a new cplan with new memory for work data was allocated.
now cplan creation only happens at the start of optimization, with each iteration reusing the cplan and its work data.
* reverse order of for loop in ggml_build_backward_expand to save memory when using gradient checkpointing and allocator
with this loop order gradient checkpointing with allocator on 16 layer model saves 13% memory; 2 layer memory it saves 2% memory.
the computation results are the same
* add API functions to access llama model tensors
* add stub example for finetuning, based on train-text-from-scratch
* move and remove code
* add API functions to access remaining model parameters:
mult, head and rot
* first draft for LORA finetune training
* remove const model and layer arguments in API functions for accessing model tensors
* bug fixes to make finetune compile
automatic allocator does not work yet
* add debug prints for training memory improvements
* fix names of lora tensors
* avoid stack overflow resulting from big ggml_cgraph
replace stack allocation and ggml_build_forward by ggml_new_graph in combination with ggml_build_forward_expand
* replace llama API functions to get model tensors by one function to get model tensor by name
LLAMA_API struct ggml_tensor * llama_get_model_tensor(struct llama_model * model, const char * name);
* remove unused call to not existing llama_get_layer_from_model
* implement ggml_compute_forward_out_prod_q_f32
* remove trailing whitespace
* add lora finetune support on quantized base model tensors
* add ggml_add_cast API function
this function works like ggml_add, but accepts a data type for the resulting tensor.
only supported for quantized src0 input.
* use ggml_add_cast in finetuning
lora-applied weights will now have data type F32, which improves gradients when finetuning quantized base models
* bug fix: actually use result type passed to ggml_add_cast
* make sure base model tensors data cannot be used in viewable operations
memory allocator would try to make lora application inplace on base model tensors.
since those are memory mapped this will result in memory access violations
* fix bug in ggml_out_prod which resulted in wrong n_dims of result tensors
* avoid keeping in memory ALL of the gradients
The problem here stems from ggml_graph_reset. This function is called in the optimization function, before each graph computation, to reset the gradients to zero. This required a unique memory slot for each gradient: allocating memory from a previosly freed memory location might lead to non-zero input gradients.
During ggml_compute_backward the gradients are build stepwise by adding or substracting new values, starting from a OP_NONE tensor which needs to contain zero-values. This requires the graph reset.
To avoid this I now remember in ggml_build_backward_expand the original OP_NONE gradient tensors in a hash table, which is passed to ggml_compute_backward. There instead of using add (or sub or similar) I test whether the existing gradient to be changed is a zero-valued-tensor by looking up its existence in the hash table. When it is such a zero-tensor it will not be modified, but replaced by the value to be added, otherwise the regular add (not inplace, allocator will take care of this) will be used. This way none of those zero-tensor values will be necessary in the final backward graph and more importantly they won't need a unique memory slot, just to make them zero.
* remove trailing whitespace
* remove debug prints and function to compute tensor data hash
* improve optimization iteration prints
* adjust maximal values to support finetuning 3B models
* change default finetune params lora_r and lora_alpha to match the n_rank parameters of 4
* bug fix: make sure finetune input gradient is allocated at begin and kept until end
* remove unnecessary src tensor from ggml_get_rows_back
we don't need data of src[2] for computation, only to setup the correct output shape.
remove dependency on src[2], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included.
this is similar to how ggml_reshape does it.
* remove unnecessary src tensor from ggml_repeat & ggml_repeat_back
we don't need data of src[1] for computation, only to setup the correct output shape.
remove dependency on src[1], so that allocator can work more freely.
the computational graph is still completely determined, because the output shape is naturally included
* resolve todo
allocator will only make it inplace when they are of the same type
* mixing multiple LORA adapters is now possible
pass more than one '--lora FNAME' argument to apply more than one LORA.
use '--lora-scaled FNAME S' when you want to specify a user-defined scale for an adapter.
* add option to save finetune output every N iterations
* also save latest finetune output with ITERATION="LATEST" and print where files are saved
saving with LATEST makes it easier to resume training from the latest checkpoint
the string "LATEST" can be configured with command line option "--fn-latest STR"
* update checkpoint train stats before saving via "--save-every"
* add command line option `--rank-wo N` for rank of wo tensor
* update finetune README
* fix dump_non_result_info_yaml to output multiple lora adapters
* bug fix: replace GGML_TYPE_SIZE[t] by ggml_type_size(t)
* replace llama_n_mult by llama_n_ff
* finetune bug fixes to compile with merged in code from master
* remove prediction related code to reduce duplicated code with main
use main instead
* reduce large memory overhead in train-text-from-scratch
all gradients had to be pinned so that graph_reset works correctly.
this is no longer necessary with the changes to ggml_compute_backward introduced in this PR.
* add comment explaining why finetune checkpoints are allocated in one block
* make default value of float member a float literal
* handle rms_norm and rope parameters the same as in train-text-from-scratch
* remove unused code
* remove vocab related code as it is unnecessary
* add LLM_KV_TRAINING_TYPE to train-text-from-scratch checkpoints
so that they can be differentiated from lora finetune checkpoints
* add gguf constants and load/save functions from train-text-from-scratch
* add load & save lora finetune checkpoints via gguf
* add python script to convert old finetune checkpoint files to gguf
* remove old checkpoint save & load code
* remove code to print data checksums which was used to verify correctness of new gguf code
* omit tokenization when training is disabled, only save llama lora adapter
training can be disabled by passing '-n 0' to finetune
* remove trailing whitespace
* update README.md
* implement ggml_compute_forward_repeat_f16
* avoid stack overflow of large cgraphs in test-grad0
* add ggml API functions ggml_unravel_index, ggml_get_i32_nd and its analogs for set and for f32
ggml_get_i32_1d, ggml_set_i32_1d, ggml_get_f32_1d, ggml_set_f32_1d now support non-contiguous tensors.
in case of non-contiguous tensor, the 1d index is unraveled into a multi index using ggml_unravel_index to be passed to '_nd' function equivalent.
this fixes a bug in test-grad0 which happens due to ggml_build_backward not building purely contiguous tensors anymore
* increase test-grad0 context mem size to accommodate for bigger cgraph
* add sanity check to ggml_compute_backward, asserting the correct shape of gradients
* fix ggml_acc_or_set to return tensor of correct shape
* remove unused 'inplace' argument from ggml_compute_backward function
inplace operations to add gradients are no longer created by ggml_compute_backward
use allocator to automatically make inplace operations
* add missing argument 'int i0' to ggml_get_i32_nd & ggml_set_i32_nd header declarations
* fix error message in ggml_allocr_alloc to display actual max_avail
* fix check_gradient
ggml_build_backward_expand was previously replaced by ggml_build_backward, but the assignment of forward graph to backward graph missing
* use tensor->view_src instead of ggml_is_view and get_view_source
* move gradient checkpointing code into ggml, new API function:
// build gradient checkpointing backward graph gb for gf using provided checkpoints
// gb_tmp will contain original backward graph with rewritten backward process nodes,
// but without the second forward pass nodes.
GGML_API void ggml_build_backward_gradient_checkpointing(
struct ggml_context * ctx,
struct ggml_cgraph * gf,
struct ggml_cgraph * gb,
struct ggml_cgraph * gb_tmp,
struct ggml_tensor * * checkpoints,
int n_checkpoints);
* replace custom data getters and setters by ggml functions
* train-text-from-scratch can train (full finetune) gguf models
just pass the gguf model via `--checkpoint-in FN`.
after this, to continue training, pass the generated checkpoint instead of the original gguf model.
tested with smaller models, bigger models may exceed available memory.
use (LORA) finetune for those.
* remove trailing whitespace
* add option to save train-text-from-scratch output every N iterations
* update README.md
* fix warnings
* fix warnings
* remove finetune option to disable allocator
the allocator should always be used.
by making sure that it is always used it gets easier to implement automatic memory requirements computation
* add tensor checkpoints only when gradient checkpointing is enabled
* initialize opt ggml context if none was provided
* add ggml-alloc API function 'ggml_allocr_max_size' to get max size of alloc
GGML_API size_t ggml_allocr_max_size(struct ggml_allocr * alloc);
* finetune: automatically allocate all memory and changes to command line options
remove '--n_examples N' parameter, as it no longer makes sense to call optimization process multiple times in a loop.
add '--only_write_lora' command line option: will skip tokenization and training, to only write a llama.cpp comptabile LORA adapter.
remove memory buffer related command line options.
improve iteration console output.
* add finetune to Makefile
* update README.md
* print time per iteration and estimate remaining time
* increase measured alloc size by tensor_alignment
ggml_allocr_reset will reduce the given size by up to tensor_alignment-1
* fix README.md
* add some more allocator debug prints
* bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue
* revert last commit
"bug fix, probably solves the 'ggml_allocr_alloc: not enough space in the buffer' issue"
"alloc was freeing an externally allocated tensor, because it calculated the end of allocator memory as alloc->data + alloc->max_size instead of alloc->data + alloc->size."
This is intentional to reduce the risk of freeing external tensors when measuring. Unless max_size is not properly calculated, I don't see why this is an issue.
* remove unnecessary "0x" before "%p" output
* move measurement memory segment to upper region of the address space
* update README.md
* fix printf format warnings
* add missing gguf_free in load_checkpoint_lora_file
* load default rms_norm and rope parameters from base model
* add gradient accumulation
specify number accumulation steps with '--grad-acc N'.
this will simulate a bigger batch size of grad_acc*batch.
* fix tracking of train_samples and train_tokens
* build : fix compile warnings
* ggml : fix L-BFGS linesearch loop
* improve finetune time measurement
fix printf warnings on system where int64_t is (long int).
change time datatypes to double because values get big with long training times.
exclude file saving from time measurement.
converge faster to actual time per iteration by removing very small first duration before first iteration was performed.
fix bug in output of total training time, the reported value was 1000 times to small.
* specify default lora rank with '--lora-r N'
'--lora-r N' will specify default rank for all tensors
'--rank-wq N', etc. will override this default rank for specific tensor types.
* fix gradient accumulation bug where the same batch was used for each microstep
* fix gradient accumulation bug where the same batch was used for each microstep
* support grouped-query-attention in ggml_flash_attn and ggml_flash_attn_back
k and v can now be repeated in q along ne[2]
in forward pass just use modulo to compute k and v indices, like ik2 = iq2 % nek2.
in backard pass this won't work as easy, because multiple threads will compete to accumulate to the same k->grad[:,ik1,ik2,ik3] and v->grad[:,iv1,iv2,iv3].
so we change the parallelization over q rows to be over k rows. this ensures non-overlapping (ik2,ik3) across threads.
in each thread we then iterate over the number of repetitions of k/v in q to compute iq2 as iq2 = ik2 + irep*nek2.
since ne2 is not the same for q,k and v we also change how the gradients are concatenated into the result tensor.
additionally the offsets of gradq, gradk and gradv in the result tensor are now memory aligned.
we also simplify the compute_backward part of flash_attn to use ggml_reshape instead of switching over the number of dimensions.
this needs a small change to ggml_reshape, removing the assertion of second argument to be contiguous.
since only the shape (ne) of the second reshape argument is of relevance, its memory layout (nb) is irrelevant -> it can very well be non-contiguous.
change test-grad0 to also test for repeated k/v in q.
this changes the rng and now results in small gradient differences in softmax. these solely come from using f16 exp table lookup in forward softmax: when temporarily changing softmax to use actual exp function, the reported gradient differences go away. gradient differences coming solely from f16 table lookup are acceptable.
added a note to explain this.
* add llama API functions to get grouped-query-attention n_head parameter 'n_head_kv'.
* fix finetune to support grouped-query-attention (using flash-attention)
note: ggml changes to ggml_out_prod are necessary to support grouped-query-attention without flash-attention.
* support broadcastable a in out_prod(a, b) and backward pass of broadcasting mul_mat(a, b)
* test broadcasting mul_mat backward pass
* decouple random number generator of each operation test
when changing one test the rng of others tests is not influenced anymore
* add comment briefly describing what ggml_repeat_back does
* simplify broadcasting mul_mat backward using ggml_repeat_back
* add cgraph evaluation order member and corresponding enum type
this controls in which order ggml_build_forward visits source nodes.
by default the nodes are visited left to right, i.e. src[0] first.
in some cases it is beneficial for ggml-alloc to visit in a different order.
two possible orders are supported: left-to-right (src[0] first) and right-to-left (src[0] last).
* measure max compute size for each cgraph eval order and use best order
this can bring huge memory savings:
e.g. codellama-34b with n_ctx=64, n_batch=1 goes from 92927.8mb down to 4627.6 MB
* remove unused command line options
* add sample start patterns and options to force new or by default resume last shuffling
* update shuffle rng state on reshuffle
* exclude known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* remove probably unnecessary exception type flags from stringstream
* pass correct max number of tokens to llama_tokenize
* account for possible leading whitespace that will be added by tokenizer
e.g. '\t' will be tokenized by llama spm tokenizer to [29871, 12]
* use unrolled vec_mad in out_prod
y is vec_mad result vec.
x is vec_mad input vec.
v is vec_mad input scalar.
ggml_vec_mad_f32_unroll will internally loop over x and v with same y.
GGML_VEC_MAD_UNROLL is by default defined to 32.
This value is empirical optimized using performance test runs of out-prod in openllama-3b finetune with 256 context length and batch size 1. It gives 23% performance boost for out_prod.
Full measurements of out-prod runtime in ms:
unroll_xv unroll_yv
1 67014.643 87826.469
2 77117.552 89077.656
4 72091.311 109121.657
8 61077.543 88678.334
16 56914.67 79514.947
24 59024.595 84350.254
28 55952.446 83368.73
32 51476.658 85177.745
36 55973.792 84659.92
40 55139.616 93844.738
48 60736.392 93330.267
64 99856.878 116994.99
Second column is when unrollying yv instead of xv
* set lora_alpha to value of lora_r if it is not set via command line
otherwise only changing lora_r will change scaling of lora adapter used in prediction
* reshuffle original sample order instead of the previous shuffled order
otherwise resumed reshuffle will not result in same sample order
* block tiling for out-prod inspired by mul-mat
block sizes are empirically optimized
roughly doubles the flops of out-prod
* exclude some more known zero values from computations in flash_attn_f32 & flash_attn_back_f32
* add static keywords
* remove outcommented old code
* update train-text-from-scratch with tokenization, sample selection and shuffling from finetune
* remove lbfgs related train parameters
* move common train functions into common/train.[h|cpp]
* move train state into struct train_state
* move train data saving code into callback to unify code of opt_callback
train_params are still different in finetune and train-text-from-scratch, so it can't yet be moved to train.h|cpp
* move common train params into common/train
* move common opt_callback into common/train
* fix consume_common_train_arg
* save and load head_count_kv in lora checkpoints
* increase train_samples by used_samples instead of number of batches
on batch can contain more than one sample when option "fill_with_next_samples" is used
* fix usage of llama_tokenize
* remove static from process_escape since we need it exposed in header
* fix code formating of long function declarations
* fix condition in load_train_state_gguf
* use die("msg") instead of replace GGML_ASSERT(!"msg") or throw std::runtime_error("msg")
* fix saving and loading of training type
* remove terminating '\0' from tokenization
(llama_tokenize is now passed the string length instead of relying on terminating '\0')
* fix compile warnings
* fix compile warnings
* use new/delete for train_state instead of malloc/free
using malloc may result in seg faults when trying to assign string fields
* assert that sample_count > 0, avoiding division by zero
* fix frand to return value in interval [0,1)
* add train option "--sample-random-offsets"
Use samples beginning at random offsets.
The offset is only applied to the first sample in each batch context window.
Together with "--fill-with-next-samples" this may help for training endless text generation.
For example given a dataset containing samples "abcd", "ABCD", "0123".
With context size of 8 and options "--fill-with-next-samples", "--no-separate-with-eos", "--no-separate-with-bos",
the context windows of batches could only be filled with "abcdABCD", "ABCDabcd", "0123abcd", etc.
With "--sample-random-offsets" it can also be filled with "23abcdAB", "bcd0123A", etc.
* deduplicate code into function
* remove n_rot hparam, as it must always be hparam.n_embd_head()
* align code
* assert correct base model tensor shapes
* move some params from lora hparams into model hparams and load model params from gguf
this equalizes the model definition in finetune and text-from-scratch and removes the need for additional llama api functions to get model parameters
* remove now unnecessary llama API functions to get model params that where added by this PR
* train-text-from-scratch: automatically allocate model tensors, remove option '--mem-model N'
* train-text-from-scratch: automatically allocate opt context
* train-text-from-scratch: automatically allocate input tensors
* train-text-from-scratch: automatically allocate compute memory
* remove unused options and equalize train-text-from-scratch with finetune
* initialize opt->loss_after with zero
* add export-lora program
* remove trailing whitespace
* add export-lora build in Makefile
* remove unused struct tensor_info from export-lora
* add export-lora build dependency to llama
because it depends on common, which depends on llama
* update finetune README.md
* cancel optimization when specified number of epochs is completed
* improve handling of export-lora arguments
print errors and warnings when files could not be read or created
* Fix export-lora.cpp "not enough space in the context's memory pool" (#1)
* Fix export-lora.cpp "not enough space in the context's memory pool"
Without this patch, export-lora would sometimes error with "not enough space in the context's memory pool (needed 656784, available 656800)".
* increase required context size by 5*GGML_MEM_ALIGN instead of plain 16
---------
Co-authored-by: xaedes <xaedes@gmail.com>
* improve handling of not yet supported tensor types
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: meatbag-18a <145869052+meatbag-18a@users.noreply.github.com>
2023-09-28 18:40:11 +00:00
2023-11-02 06:50:16 +00:00
speculative : examples /speculative /speculative .cpp ggml .o llama .o $( COMMON_DEPS ) grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-09-04 10:39:57 +00:00
2023-11-02 06:50:16 +00:00
parallel : examples /parallel /parallel .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-09-28 16:04:36 +00:00
2023-11-26 18:33:07 +00:00
lookahead : examples /lookahead /lookahead .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-11-26 18:33:07 +00:00
2023-12-22 16:05:56 +00:00
lookup : examples /lookup /lookup .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-12-22 16:05:56 +00:00
2024-01-08 09:14:04 +00:00
passkey : examples /passkey /passkey .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2024-01-08 09:14:04 +00:00
2023-10-11 11:14:05 +00:00
i f e q ( $( UNAME_S ) , D a r w i n )
swift : examples /batched .swift
( cd examples/batched.swift; make build)
e n d i f
2023-11-02 06:50:16 +00:00
common/build-info.cpp : $( wildcard .git /index ) scripts /build -info .sh
2024-02-05 18:33:00 +00:00
@sh scripts/build-info.sh " $( CC) " > $@ .tmp
2023-05-01 16:23:47 +00:00
@if ! cmp -s $@ .tmp $@ ; then \
mv $@ .tmp $@ ; \
else \
rm $@ .tmp; \
fi
2023-04-13 14:03:57 +00:00
2023-11-02 06:50:16 +00:00
build-info.o : common /build -info .cpp
$( CXX) $( CXXFLAGS) -c $( filter-out %.h,$^) -o $@
2023-03-10 18:40:58 +00:00
#
# Tests
#
2023-07-21 10:09:16 +00:00
tests : $( TEST_TARGETS )
2023-11-02 06:50:16 +00:00
benchmark-matmult : examples /benchmark /benchmark -matmult .cpp build -info .o ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-09-28 21:41:44 +00:00
run-benchmark-matmult : benchmark -matmult
2023-04-30 12:32:37 +00:00
./$@
2023-04-13 14:03:57 +00:00
2023-10-11 11:14:05 +00:00
.PHONY : run -benchmark -matmult swift
2023-09-28 21:41:44 +00:00
2023-06-07 07:59:52 +00:00
vdot : pocs /vdot /vdot .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-05-01 16:23:47 +00:00
2023-09-28 21:41:44 +00:00
q8dot : pocs /vdot /q 8dot .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-09-28 21:41:44 +00:00
2023-11-30 22:23:08 +00:00
tests/test-llama-grammar : tests /test -llama -grammar .cpp ggml .o grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-17 07:41:01 +00:00
2023-11-30 22:23:08 +00:00
tests/test-grammar-parser : tests /test -grammar -parser .cpp ggml .o llama .o grammar -parser .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-13 14:00:48 +00:00
2023-11-30 22:23:08 +00:00
tests/test-double-float : tests /test -double -float .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-11-30 22:23:08 +00:00
tests/test-grad0 : tests /test -grad 0.cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-11-30 22:23:08 +00:00
tests/test-opt : tests /test -opt .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-11-30 22:23:08 +00:00
tests/test-quantize-fns : tests /test -quantize -fns .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-11-30 22:23:08 +00:00
tests/test-quantize-perf : tests /test -quantize -perf .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-11-30 22:23:08 +00:00
tests/test-sampling : tests /test -sampling .cpp ggml .o llama .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-07-21 10:09:16 +00:00
2023-12-12 09:27:26 +00:00
tests/test-tokenizer-0-falcon : tests /test -tokenizer -0-falcon .cpp ggml .o llama .o $( COMMON_DEPS ) console .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-28 15:38:35 +00:00
2023-12-12 09:27:26 +00:00
tests/test-tokenizer-0-llama : tests /test -tokenizer -0-llama .cpp ggml .o llama .o $( COMMON_DEPS ) console .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-28 15:38:35 +00:00
2023-12-12 09:27:26 +00:00
tests/test-tokenizer-1-bpe : tests /test -tokenizer -1-bpe .cpp ggml .o llama .o $( COMMON_DEPS ) console .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-10-03 07:16:26 +00:00
2023-12-12 09:27:26 +00:00
tests/test-tokenizer-1-llama : tests /test -tokenizer -1-llama .cpp ggml .o llama .o $( COMMON_DEPS ) console .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-08-30 06:20:26 +00:00
2023-11-30 22:23:08 +00:00
tests/test-rope : tests /test -rope .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2023-11-30 22:23:08 +00:00
2023-08-30 06:20:26 +00:00
tests/test-c.o : tests /test -c .c llama .h
$( CC) $( CFLAGS) -c $( filter-out %.h,$^) -o $@
2023-12-07 20:26:54 +00:00
tests/test-backend-ops : tests /test -backend -ops .cpp ggml .o $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2024-01-17 16:38:39 +00:00
2024-01-26 12:18:00 +00:00
tests/test-model-load-cancel : tests /test -model -load -cancel .cpp ggml .o llama .o tests /get -model .cpp $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2024-01-26 12:18:00 +00:00
tests/test-autorelease : tests /test -autorelease .cpp ggml .o llama .o tests /get -model .cpp $( COMMON_DEPS ) $( OBJS )
2024-02-05 18:33:00 +00:00
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
2024-02-19 08:23:37 +00:00
tests/test-chat-template : tests /test -chat -template .cpp ggml .o llama .o $( COMMON_DEPS ) $( OBJS )
$( CXX) $( CXXFLAGS) -c $< -o $( call GET_OBJ_FILE, $<)
$( CXX) $( CXXFLAGS) $( filter-out %.h $<,$^) $( call GET_OBJ_FILE, $<) -o $@ $( LDFLAGS)
