ggml : add ggml_soft_max_ext #4256

ggerganov · 2023-11-29T15:19:33Z

Fused op:

soft_max_ext(x) = soft_max(x*scale + mask)

Should improve PP performance for large contexts ( > 2k)

Also:

~~multi-thread ggml_soft_max on CPU~~
CUDA and Metal soft max kernels use up to 1024 threads (based on dim 0) instead of always 32 and perform warp-based reduction
Metal rms_norm kernel changed to perform warp-based reduction as well

Perf CUDA V100

LLAMA_CUBLAS=1 make -j llama-bench && ./llama-bench -m ./models/openllama-7b-v2/ggml-model-q4_0.gguf -p 512,1024,4096,8192 -ngl 99

ggml_init_cublas: GGML_CUDA_FORCE_MMQ: no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
Device 0: Tesla V100-PCIE-16GB, compute capability 7.0

model	size	backend	ngl	test	master t/s	PR t/s	speedup
llama 7B Q4_0	3.56 GiB	CUDA	99	pp 512	2661.62 ± 71.30	2659.81 ± 65.20	0.999
llama 7B Q4_0	3.56 GiB	CUDA	99	pp 1024	2515.25 ± 39.75	2577.93 ± 46.42	1.025
llama 7B Q4_0	3.56 GiB	CUDA	99	pp 4096	1793.41 ± 24.16	2106.14 ± 29.02	1.174
llama 7B Q4_0	3.56 GiB	CUDA	99	pp 8192	1297.36 ± 12.95	1552.12 ± 18.71	1.196
llama 7B Q4_0	3.56 GiB	CUDA	99	tg 128	112.52 ± 0.64	114.98 ± 0.58	1.022

LLAMA_CUBLAS=1 make -j batched-bench && ./batched-bench ./models/openllama-7b-v2/ggml-model-q4_0.gguf 10880 0 99 0 2048 128 1,2,4,5

master

PP	TG	B	N_KV	T_PP s	S_PP t/s	T_TG s	S_TG t/s	T s	S t/s
2048	128	1	2176	0.982	2085.81	1.703	75.16	2.685	810.45
2048	128	2	4352	2.329	1758.55	2.587	98.94	4.917	885.17
2048	128	4	8704	6.415	1276.99	3.737	137.02	10.152	857.38
2048	128	5	10880	9.177	1115.80	5.052	126.67	14.230	764.60

PPL = 7.0773 +/- 0.04208

PR

PP	TG	B	N_KV	T_PP s	S_PP t/s	T_TG s	S_TG t/s	T s	S t/s
2048	128	1	2176	0.908	2255.56	1.645	77.80	2.553	852.25
2048	128	2	4352	1.994	2054.42	2.509	102.04	4.502	966.57
2048	128	4	8704	5.375	1524.08	3.553	144.12	8.928	974.95
2048	128	5	10880	7.669	1335.30	4.732	135.26	12.400	877.40

PPL = 7.0768 +/- 0.04207

Perf M2 Ultra

make -j llama-bench && ./llama-bench -m ./models/llama-7b-v2/ggml-model-q4_0.gguf -p 512,1024,4096,8192 -ngl 99

model	size	backend	ngl	test	master t/s	PR t/s	speedup
llama 7B Q4_0	3.56 GiB	Metal	99	pp 512	1238.73 ± 0.94	1275.95 ± 1.05	1.03
llama 7B Q4_0	3.56 GiB	Metal	99	pp 1024	1206.55 ± 0.34	1260.08 ± 0.96	1.04
llama 7B Q4_0	3.56 GiB	Metal	99	pp 4096	1080.81 ± 0.30	1154.05 ± 0.30	1.07
llama 7B Q4_0	3.56 GiB	Metal	99	pp 8192	933.17 ± 0.32	1023.16 ± 1.19	1.10
llama 7B Q4_0	3.56 GiB	Metal	99	tg 128	94.40 ± 0.04	96.23 ± 0.06	1.02

make -j batched-bench && ./batched-bench ./models/llama-7b-v2/ggml-model-q4_0.gguf 10880 0 99 0 2048 128 1,2,4,5

master

PP	TG	B	N_KV	T_PP s	S_PP t/s	T_TG s	S_TG t/s	T s	S t/s
2048	128	1	2176	1.770	1157.14	1.785	71.71	3.555	612.10
2048	128	2	4352	3.793	1079.79	7.648	33.47	11.442	380.36
2048	128	4	8704	8.774	933.72	9.399	54.48	18.172	478.98
2048	128	5	10880	11.817	866.53	10.285	62.22	22.103	492.25

PPL = 6.1212 +/- 0.03511

PR

PP	TG	B	N_KV	T_PP s	S_PP t/s	T_TG s	S_TG t/s	T s	S t/s
2048	128	1	2176	1.677	1221.20	1.741	73.54	3.418	636.70
2048	128	2	4352	3.512	1166.13	7.555	33.89	11.067	393.24
2048	128	4	8704	7.786	1052.16	9.099	56.27	16.885	515.49
2048	128	5	10880	10.243	999.71	10.004	63.98	20.247	537.38

PPL = 6.1212 +/- 0.03511

make -j llama-bench && ./llama-bench -m ./models/llama-7b-v2/ggml-model-q4_0.gguf -p 512 -n 128,256,512 -ngl 99

master

model	size	params	backend	ngl	test	t/s
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	pp 512	1237.37 ± 1.50
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 128	94.49 ± 0.03
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 256	93.79 ± 0.04
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 512	92.37 ± 0.18

build: 1f5cd83 (1577)

PR

model	size	params	backend	ngl	test	t/s
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	pp 512	1269.86 ± 1.40
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 128	98.14 ± 0.02
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 256	97.36 ± 0.06
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 512	95.78 ± 0.05

Perf M1 Pro

make -j llama-bench && ./llama-bench -m ./models/llama-7b-v2/ggml-model-q4_0.gguf -p 512 -n 128,256,512 -ngl 99

master

model	size	params	backend	ngl	test	t/s
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	pp 512	264.95 ± 0.92
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 128	36.52 ± 0.13
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 256	36.21 ± 0.01
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 512	35.54 ± 0.01

build: 1f5cd83 (1577)

PR

model	size	params	backend	ngl	test	t/s
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	pp 512	273.44 ± 0.90
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 128	37.18 ± 0.03
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 256	36.80 ± 0.03
llama 7B mostly Q4_0	3.56 GiB	6.74 B	Metal	99	tg 512	36.20 ± 0.01

build: c4db592 (1589)

ggml-ci

slaren · 2023-11-29T20:10:39Z

I see a ~15% speedup in large contexts with CUDA. We should probably check that the performance of the base softmax with CPU isn't affected too much by this, this adds conditionals in the inner loops.

ggml_init_cublas: GGML_CUDA_FORCE_MMQ: no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
Device 0: NVIDIA GeForce RTX 3090 Ti, compute capability 8.6

model	backend	ngl	test	master `1f5cd83`	PR `580fe20`	speedup
llama 7B mostly Q4_0	CUDA	99	pp 512	3761.61 ± 76.47	3900.34 ± 87.43	1.036
llama 7B mostly Q4_0	CUDA	99	pp 1024	3565.39 ± 37.27	3711.82 ± 27.71	1.041
llama 7B mostly Q4_0	CUDA	99	pp 4096	2523.32 ± 12.70	2711.08 ± 17.84	1.074
llama 7B mostly Q4_0	CUDA	99	pp 8192	1776.08 ± 8.35	1994.95 ± 10.97	1.123
llama 7B mostly Q4_0	CUDA	99	pp 16384	1124.91 ± 6.93	1306.34 ± 10.11	1.161
llama 7B mostly Q4_0	CUDA	99	tg 128	124.38 ± 0.81	127.00 ± 0.98	1.021

ggerganov · 2023-11-29T21:11:17Z

I think the CPU performance is degraded atm (did some quick tests earlier, haven't looked deeper yet). But I suspect it's something to do with the threading.

ggml-cuda.cu

ggerganov · 2023-11-30T16:23:14Z

ggml-cuda.cu

+
+    const int block_size = blockDim.x;
+
+    __shared__ float buf[CUDA_SOFT_MAX_BLOCK_SIZE];


Replaced the old WARP_SIZE kernel with a new that uses up to CUDA_SOFT_MAX_BLOCK_SIZE=512 threads per block (based on ne00). This seems to perform better on V100.

Any reason to prefer the old kernel?

Block wide reductions are a lot more expensive than warp reductions, so most of the kernels only distribute operations that need a reduction to a warp. A block reduction is usually done by doing a reduction in every warp and storing the result of each warp to shared memory, and then doing a reduction of the shared memory.

Updated to use warp-based block reduction with max of 32 warps. Results improved, more significantly for the largest contexts.

I'll try to do similar optimization for the Metal kernel

ggerganov

Should be ready to merge. Improved performance with Metal and CUDA. On CPU - not much difference

slaren · 2023-11-30T19:02:15Z

Much larger speedup now, but it is probably broken, perplexity returns only nans.

model	backend	ngl	test	master `1f5cd83`	PR `6b86bcf`	speedup
llama 7B mostly Q4_0	CUDA	99	pp 512	3761.61 ± 76.47	3848.99 ± 68.95	1.023
llama 7B mostly Q4_0	CUDA	99	pp 1024	3565.39 ± 37.27	3646.32 ± 39.29	1.022
llama 7B mostly Q4_0	CUDA	99	pp 4096	2523.32 ± 12.70	3039.41 ± 23.85	1.204
llama 7B mostly Q4_0	CUDA	99	pp 8192	1776.08 ± 8.35	2502.62 ± 15.85	1.409
llama 7B mostly Q4_0	CUDA	99	pp 16384	1124.91 ± 6.93	1745.35 ± 16.31	1.551
llama 7B mostly Q4_0	CUDA	99	tg 128	124.38 ± 0.81	127.22 ± 1.15	1.022

ggerganov · 2023-11-30T19:43:15Z

Thanks for checking - the shared arrays were not initialized so the warps were reducing random data. Should be fixed now

slaren · 2023-11-30T19:57:38Z

Tested again, perplexity looks good now with CUDA, and the speedup is still the same. Nice!

ggml-ci

LostRuins · 2023-12-02T16:41:07Z

This appears to break CLBlast prompt processing in some cases: #4296

This reverts commit ef47ec1.

This reverts commit a8e66ef.

commit 53b5ae02cb1b533b78302422951bcfdeca6e2738 Author: YellowRoseCx <[email protected]> Date: Tue Dec 12 12:08:29 2023 -0600 mixtral fan service commit 168b1d74e26d0321e2e89358303b6c33e8d7d33e Merge: f13295b de15d4a6 Author: YellowRoseCx <[email protected]> Date: Tue Dec 12 12:00:52 2023 -0600 Merge branch 'kcpp-rocm-mixtral2' into main2 commit de15d4a632939a685ec12fa17355298542facf15 Merge: 74acc54 ea4402b Author: YellowRoseCx <[email protected]> Date: Tue Dec 12 11:45:19 2023 -0600 Merge branch 'mixtral' into kcpp-rocm-mixtral commit ea4402b Author: Georgi Gerganov <[email protected]> Date: Tue Dec 12 17:03:38 2023 +0200 test-backend-ops : add one more sum_rows test commit a51bc0c Author: Georgi Gerganov <[email protected]> Date: Tue Dec 12 15:55:42 2023 +0200 metal : fix binary ops for ne10 % 4 != 0 commit 08eb991 Author: Georgi Gerganov <[email protected]> Date: Tue Dec 12 14:14:15 2023 +0200 metal : add cpy f16 -> f32 kernel commit a742d9f Author: slaren <[email protected]> Date: Tue Dec 12 12:46:33 2023 +0100 gguf-py : bump version commit 6a419f4 Author: Georgi Gerganov <[email protected]> Date: Tue Dec 12 13:04:33 2023 +0200 convert : support safetensors format commit 74acc54 Author: Concedo <[email protected]> Date: Tue Dec 12 10:53:34 2023 +0800 Revert "Hide hipBLAS (ROCm) if CuBLAS exists - vice versa" This reverts commit 4b854d4. commit f1cbfab Author: slaren <[email protected]> Date: Mon Dec 11 20:02:55 2023 +0100 convert : fix style commit 7dc75e3 Author: slaren <[email protected]> Date: Mon Dec 11 20:00:28 2023 +0100 convert : use 1e6 rope_freq_base for mixtral commit 296c945 Author: slaren <[email protected]> Date: Mon Dec 11 16:53:25 2023 +0100 cuda : fix mul_mat_id with multi gpu commit 33e50f1 Author: slaren <[email protected]> Date: Mon Dec 11 12:27:48 2023 +0100 test-backend-ops : disable MOE test with thread sanitizer commit ffda94c Author: slaren <[email protected]> Date: Mon Dec 11 12:15:31 2023 +0100 test-backend-ops : simplify and disable slow tests to avoid CI timeout commit 06581f2 Author: Concedo <[email protected]> Date: Mon Dec 11 16:54:42 2023 +0800 perf endpoint lets you monitor if the embedded horde worker has issues commit fce971d Author: Concedo <[email protected]> Date: Mon Dec 11 16:17:10 2023 +0800 do not build the clblast noavx2 binary if not on windows commit 8cbaed1 Author: Georgi Gerganov <[email protected]> Date: Mon Dec 11 08:55:16 2023 +0200 llama : fix hard-coded number of experts commit 4b854d4 Author: YellowRoseCx <[email protected]> Date: Sun Dec 10 22:49:35 2023 -0600 Hide hipBLAS (ROCm) if CuBLAS exists - vice versa commit b002981 Author: slaren <[email protected]> Date: Mon Dec 11 02:43:52 2023 +0100 test-backend-ops : fix dequantize block offset commit f1380d7 Author: slaren <[email protected]> Date: Sun Dec 10 22:58:31 2023 +0100 test-backend-ops : add cpy from f32 -> all types test commit 54d254b Author: slaren <[email protected]> Date: Sun Dec 10 21:52:11 2023 +0100 test-backend-ops : cleanup, add moe test for batches commit e2cf3b7 Author: henk717 <[email protected]> Date: Sun Dec 10 14:30:17 2023 +0100 koboldcpp.sh - The Mamba Multitool (LostRuins#554) * .sh script V1 * koboldcpp.sh polish * koboldcpp.sh dist generator * Include html's in dist * RWKV in Linux Dist * Lower dependency requirements * Eliminate wget dependency * More distinct binary name I know its technically amd64, but I don't want to cause confusion among nvidia users. * Use System OpenCL Unsure how this will behave in the pyinstaller build, but pocl ended up CPU only. With a bit of luck the pyinstaller uses the one from the actual system if compiled in a system without opencl, while conda now includes it for that specific system. * Add cblas dependency Missing this causes compile failures on some system's * ICD workaround Ideally we find a better solution, but conda forces ICD and needs this for the successful compile. However, pyinstaller then embeds the ICD causing it to be limited to the system it was compiled for. By temporarily removing the ICD pyinstaller can't find it and everything remains functional. Ideally we do this on a pyinstaller level, but I could not find any good options to do so yet. --------- Co-authored-by: root <root@DESKTOP-DQ1QRAG> commit 54ba263 Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 15:27:41 2023 +0200 test-backend-ops : make experts more evenly probable (test_moe) commit b0b83dd Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 14:30:38 2023 +0200 metal : fix ggml_mul_mat_id for F32 commit 65923a8 Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 14:17:46 2023 +0200 convert : determine n_ctx correctly commit 8614aa7 Author: slaren <[email protected]> Date: Sun Dec 10 13:12:11 2023 +0100 cuda : fix get_rows when ncols is odd commit cefebb3 Author: slaren <[email protected]> Date: Sun Dec 10 13:11:39 2023 +0100 test-backend-ops : add moe test commit e640cbe Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 13:57:54 2023 +0200 llama : add n_expert and n_expert_used to hparams + change quants commit d1259b7 Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 13:00:13 2023 +0200 llama : do not quantize expert gating tensors commit 6cfb31f Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 10:59:13 2023 +0200 metal : add indirect mat-vec kernels for all quantization types commit 016f9bb Author: Georgi Gerganov <[email protected]> Date: Sun Dec 10 09:38:21 2023 +0200 metal : fix ggml_get_rows to work with non-cont src1 commit 0710b0f Author: slaren <[email protected]> Date: Sat Dec 9 23:29:47 2023 +0100 llama : offload missing ffn_moe_silu commit 62b95f9 Author: slaren <[email protected]> Date: Sat Dec 9 22:39:34 2023 +0100 cuda : support non-contiguous src1 in get_rows commit 2e4db48 Author: slaren <[email protected]> Date: Sat Dec 9 22:38:22 2023 +0100 ggml : update get_rows f16 and q commit ac3f7d8 Author: slaren <[email protected]> Date: Sat Dec 9 19:19:03 2023 +0100 ggml : get_rows : support non-contiguos tensors with gaps, generalize up to 3D commit 8c5b66e Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 15:30:34 2023 +0200 metal : reduce the kernel launches for ggml_mul_mat_id commit 7e2006b Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 14:24:58 2023 +0200 metal : add/mul/div use general kernel when src1 not cont commit 06dfde3 Author: slaren <[email protected]> Date: Sat Dec 9 13:21:09 2023 +0100 llama : add basic support for offloading moe with CUDA commit 2cbcba8 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 14:18:42 2023 +0200 metal : add more general support for ggml_get_rows + tests commit 9064b1c Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 14:04:54 2023 +0200 ggml : fix ggml_get_rows to take into account ne02 / ne11 commit ee8fb39 Author: slaren <[email protected]> Date: Sat Dec 9 12:42:25 2023 +0100 ggml : add n_as argument to ggml_mul_mat_id commit 7372b62 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 13:18:58 2023 +0200 ggml : ggml_get_rows support 2D indexing [n_tokens, n_experts] (cpu only) commit 8b185b7 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 13:01:42 2023 +0200 llama : fix expert weighting in the FFN commit 7ea3695 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 12:45:15 2023 +0200 llama : first working version commit af1a096 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 12:07:39 2023 +0200 llama : fix cur -> cur_expert commit aedfad1 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 11:47:40 2023 +0200 llama : update graph to support MoE commit 861cd67 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 11:19:46 2023 +0200 ggml : sync latest ggml_mul_mat_id commit a3eefe9 Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 11:14:03 2023 +0200 llama : model loading commit d38e41e Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 10:59:37 2023 +0200 convert : fix n_ff typo commit dff8cbe Author: Georgi Gerganov <[email protected]> Date: Sat Dec 9 10:51:58 2023 +0200 convert : support Mixtral as LLAMA arch commit 7a69152 Author: Concedo <[email protected]> Date: Fri Dec 8 21:06:32 2023 +0800 lowvram var defaults commit 7418bca Author: Concedo <[email protected]> Date: Fri Dec 8 19:20:30 2023 +0800 up ver commit c47bc28 Author: Concedo <[email protected]> Date: Fri Dec 8 18:35:45 2023 +0800 slight refactor for noscript ui commit 7469f20 Author: Concedo <[email protected]> Date: Fri Dec 8 18:16:14 2023 +0800 use lowvram flag for offload qkv commit ec21fa7 Merge: 930cdfb fe680e3 Author: Concedo <[email protected]> Date: Fri Dec 8 17:42:26 2023 +0800 Merge branch 'master' into concedo_experimental # Conflicts: # .github/workflows/build.yml # .gitignore # CMakeLists.txt # Makefile # Package.swift # README.md # ggml-cuda.cu # llama.cpp # llama.h # scripts/sync-ggml.sh # tests/CMakeLists.txt commit 930cdfb Author: Concedo <[email protected]> Date: Fri Dec 8 16:53:30 2023 +0800 updated lite, added patch that links to noscript mode commit fe680e3 Author: Georgi Gerganov <[email protected]> Date: Thu Dec 7 22:26:54 2023 +0200 sync : ggml (new ops, tests, backend, etc.) (ggerganov#4359) * sync : ggml (part 1) * sync : ggml (part 2, CUDA) * sync : ggml (part 3, Metal) * ggml : build fixes ggml-ci * cuda : restore lost changes * cuda : restore lost changes (StableLM rope) * cmake : enable separable compilation for CUDA ggml-ci * ggml-cuda : remove device side dequantize * Revert "cmake : enable separable compilation for CUDA" This reverts commit 09e35d0. * cuda : remove assert for rope * tests : add test-backend-ops * ggml : fix bug in ggml_concat * ggml : restore `ggml_get_n_tasks()` logic in `ggml_graph_plan()` * ci : try to fix macOS * ggml-backend : remove backend self-registration * ci : disable Metal for macOS cmake build ggml-ci * metal : fix "supports family" call * metal : fix assert * metal : print resource path ggml-ci --------- Co-authored-by: slaren <[email protected]> commit bcc0eb4 Author: Georgi Gerganov <[email protected]> Date: Thu Dec 7 13:03:17 2023 +0200 llama : per-layer KV cache + quantum K cache (ggerganov#4309) * per-layer KV * remove unnecessary copies * less code duplication, offload k and v separately * llama : offload KV cache per-layer * llama : offload K shift tensors * llama : offload for rest of the model arches * llama : enable offload debug temporarily * llama : keep the KV related layers on the device * llama : remove mirrors, perform Device -> Host when partial offload * common : add command-line arg to disable KV cache offloading * llama : update session save/load * llama : support quantum K cache (ggerganov#4312) * llama : support quantum K cache (wip) * metal : add F32 -> Q8_0 copy kernel * cuda : add F32 -> Q8_0 copy kernel ggml-ci * cuda : use mmv kernel for quantum cache ops * llama : pass KV cache type through API * llama : fix build ggml-ci * metal : add F32 -> Q4_0 copy kernel * metal : add F32 -> Q4_1 copy kernel * cuda : wip * cuda : add F32 -> Q4_0 and F32 -> Q4_1 copy kernels * llama-bench : support type_k/type_v * metal : use mm kernel only for quantum KV cache * cuda : add comment * llama : remove memory_f16 and kv_f16 flags --------- Co-authored-by: slaren <[email protected]> * readme : add API change notice --------- Co-authored-by: slaren <[email protected]> commit 81bc921 Author: Hongyu Ouyang <[email protected]> Date: Thu Dec 7 02:25:22 2023 -0800 train : fix ggerganov#4227 (double free in examples/train-text-from-scratch/train-text-from-scratch.cpp) (ggerganov#4351) On commit b1108 (44c117f) xaedes added ggml_allocr * alloc = NULL; ... (many lines in between) if (alloc) { ggml_allocr_free(alloc); } Which is correct, but it's easy to lose context after many lines in between. On commit b1287 (0e76a899) xaedes made a big change. From here on, alloc is freed eagerly. alloc = ggml_allocr_new(...) ... (short lines of code) ggml_allocr_free(alloc) This happens a few times, but alloc is never set to NULL, and many lines below, we still have if (alloc) { ggml_allocr_free(alloc); } which causes a double-free. commit 05cd6e5 Author: Georgi Gerganov <[email protected]> Date: Wed Dec 6 20:21:59 2023 +0200 server : recognize cache_prompt parameter in OAI API (ggerganov#4347) commit c751152 Author: Concedo <[email protected]> Date: Thu Dec 7 00:52:25 2023 +0800 noscript mode is done commit 12002d8 Author: Concedo <[email protected]> Date: Wed Dec 6 17:51:08 2023 +0800 very basic noscript mode commit caa9249 Author: Georgi Gerganov <[email protected]> Date: Wed Dec 6 10:41:03 2023 +0200 common : fix compile warning commit da5eaef Author: stduhpf <[email protected]> Date: Wed Dec 6 09:08:17 2023 +0100 speculative : support `--color` (ggerganov#4343) * speculative: add some colors * minor : add braces --------- Co-authored-by: Georgi Gerganov <[email protected]> commit 5f6e0c0 Author: Marcus Dunn <[email protected]> Date: Tue Dec 5 10:55:12 2023 -1000 grammar : pre-computed pieces + reserve mem + less string copies (ggerganov#4330) * reserve space for codepoints * improvement for the appended 0 * used precomputed token text for grammar sample * reserve canidates_decoded * reserve canidates_grammar * remove candidates_decoded * Revert "remove candidates_decoded" This reverts commit 3773328. * changed decode_utf8 to take src by ref commit 5aa365d Author: Kerfuffle <[email protected]> Date: Tue Dec 5 10:19:18 2023 -0700 llama : allow overriding GGUF metadata when loading model (ggerganov#4092) * feat: Allow overriding GGUF metadata when loading model * Fix the one time GCC is stricter than clang about something * Step1 * Refactor... basically everything! * Nuke obsolete GetArrayLen struct * simplify std::string specialization * Various cleanups Add informational output when overrides are applied Warn user when an override with the wrong type is specified * Fix broken logic for parsing bool KV overrides Fix issue where overrides didn't apply when key missing in GGUF metadata Resolve merge changes * llama : rearrange model params * Update new GET_KEY call Add note that metadata KV overrides aren't reflected in initial metadata KV info dump --------- Co-authored-by: cebtenzzre <[email protected]> Co-authored-by: Georgi Gerganov <[email protected]> commit b6f952f Author: Concedo <[email protected]> Date: Tue Dec 5 21:08:10 2023 +0800 improved exit logic commit 52c8bc3 Author: MaggotHATE <[email protected]> Date: Tue Dec 5 15:05:51 2023 +0500 sampling : custom samplers order (ggerganov#4285) * Samplers sequence order w parameter * Cleaned commented code * Fixed formatting * Rewrote with unordered_map * Revert and rewrite, too many problems and safeguards would be needed * Fixed code style * Code style fixes according to review * More readable samplers input string, fixed help * Style fix in sampler_queue * Formatting fixes * Fixing whitespaces commit e4b76bb Author: kchro3 <[email protected]> Date: Mon Dec 4 23:29:46 2023 -0800 swift : revert compiler checks for swift package (ggerganov#4332) commit 23b5e12 Author: Daniel Bevenius <[email protected]> Date: Mon Dec 4 17:04:21 2023 +0100 simple : update error message for KV cache check (ggerganov#4324) This commit updates the error message that is printed when the KV cache is not big enough to hold all the prompt and generated tokens. Specifically it removes the reference to n_parallel and replaces it with n_len. Signed-off-by: Daniel Bevenius <[email protected]> commit d208995 Author: Miwa / Ensan <[email protected]> Date: Tue Dec 5 01:03:49 2023 +0900 swift : fix concatenation method to avoid invalid UTF8 stringfication (ggerganov#4325) commit 5c9f90c Author: Miwa / Ensan <[email protected]> Date: Mon Dec 4 22:43:45 2023 +0900 swift : fix prompt tokenization logic (ggerganov#4321) commit a5a5839 Author: Concedo <[email protected]> Date: Mon Dec 4 21:10:42 2023 +0800 handle accidentally selecting a kcpps file as model instead commit 4fa44e8 Author: Ikko Eltociear Ashimine <[email protected]> Date: Mon Dec 4 16:57:35 2023 +0900 grammar-parser : fix typo (ggerganov#4318) preceeding -> preceding commit 8602f5a Merge: ac36aee fbbc428 Author: Concedo <[email protected]> Date: Sun Dec 3 22:00:14 2023 +0800 Merge branch 'master' into concedo_experimental commit fbbc428 Author: Georgi Gerganov <[email protected]> Date: Sun Dec 3 15:56:35 2023 +0200 ggml : reuse ggml_get_n_tasks() in ggml_graph_plan() (ggerganov#4308) * ggml : fix soft max out-of-bounds access ggml-ci * ggml : reuse ggml_get_n_tasks() in ggml_graph_plan() ggml-ci commit ac36aee Merge: 48544cd 33e171d Author: Concedo <[email protected]> Date: Sun Dec 3 21:56:29 2023 +0800 Merge branch 'master' into concedo_experimental # Conflicts: # CMakeLists.txt # Makefile commit adf3de4 Author: Georgi Gerganov <[email protected]> Date: Sun Dec 3 15:56:22 2023 +0200 ggml : fix soft max out-of-bounds access (ggerganov#4307) ggml-ci commit 48544cd Author: Concedo <[email protected]> Date: Sun Dec 3 21:46:50 2023 +0800 Revert "Revert "ggml : add ggml_soft_max_ext (ggerganov#4256)"" This reverts commit a8e66ef. commit 33e171d Author: Ed Lee <[email protected]> Date: Sun Dec 3 01:10:43 2023 -0800 server : fix OpenAI API `stop` field to be optional (ggerganov#4299) (cherry picked from commit Mozilla-Ocho/llamafile@e8c92bc) commit 6949b50 Author: Rickard Edén <[email protected]> Date: Sun Dec 3 10:03:25 2023 +0100 py : add grammar to oai like api (ggerganov#4294) commit d7b800b Author: Georgi Gerganov <[email protected]> Date: Sun Dec 3 10:58:16 2023 +0200 llama : pad KV cache size (ggerganov#4280) * llama : pad KV cache size to 32 * metal : try to improve batched decoding commit 6570a20 Author: Concedo <[email protected]> Date: Sun Dec 3 15:44:53 2023 +0800 token count includes ids commit 5a7d312 Author: Georgi Gerganov <[email protected]> Date: Fri Dec 1 20:39:12 2023 +0200 llama : avoid using "optional" keyword (ggerganov#4283) commit d5a1cbd Author: Georgi Gerganov <[email protected]> Date: Fri Dec 1 20:35:03 2023 +0200 llama : support optional tensors (ggerganov#4283) commit b220222 Author: Miwa / Ensan <[email protected]> Date: Sat Dec 2 03:19:45 2023 +0900 swift : fix token_to_piece implementation (ggerganov#4278) * Fix token_to_piece implementation in Swift * Fix errors commit 511f52c Author: Jared Van Bortel <[email protected]> Date: Fri Dec 1 13:18:35 2023 -0500 build : enable libstdc++ assertions for debug builds (ggerganov#4275) commit 03562f3 Author: CausalLM <[email protected]> Date: Sat Dec 2 02:17:06 2023 +0800 llama : support attention bias on LLaMA architecture (ggerganov#4283) * Support attention_bias on LLaMA architecture QKVO bias, should fix InternLM (ggerganov#3133) and works for LLaMAfied Qwen models (ggerganov#3743 (comment)). * check existence of qkvo bias while loading llama models Tested on LLaMA2, CUDA and CPU. * Update llama.cpp commit 37c746d Author: Shijie <[email protected]> Date: Sat Dec 2 02:16:31 2023 +0800 llama : add Qwen support (ggerganov#4281) * enable qwen to llama.cpp * llama : do not GPU split bias tensors --------- Co-authored-by: Georgi Gerganov <[email protected]> commit 880f579 Author: Georgi Gerganov <[email protected]> Date: Fri Dec 1 18:42:11 2023 +0200 llama : fix integer overflow during quantization (ggerganov#4284) happens with multi-threaded quantization of Qwen-72B ggml-ci

* metal : implement soft_max_ext * cuda : implement soft_max_ext * ggml : implement soft_max_ext (CPU) * batched-bench : print threads ggml-ci * metal : simplify soft_max encoding ggml-ci * cuda : use 512 threads for soft_max instead of 32 * ggml : update soft max cpu * cuda : do warp-based block reduce * cuda : increase max block size to 1024 * cuda : fix warp reduction initialization of shared mem * metal : warp-based reduction for soft max kernel * metal : warp-based reduce for rms_norm * metal : simplify soft max kernel ggml-ci * alloc : fix build with debug

llama : restore prefix space in llama tokenizer (ggerganov#4081) gguf : fix potential infinite loops while parsing (ggerganov#4100) Co-authored-by: Bernhard Gstrein <[email protected]> Respect tokenizer.ggml.add_bos_token value when tokenizing (ggerganov#4040) * gguf-py: gguf-dump: Respect --no-tensor flag in JSON mode. * Respect add_bos_token GGUF metadata value * gguf-py: Try to fix SpecialVocab giving up too easily for the Nth time llama : fix data units (ggerganov#4101) * llama : fix data units ggml-ci * Revert "llama : fix data units" This reverts commit f5feac8. * llama : disambiguate data units ggml-ci cuda : get_row_rounding F32 (ggerganov#4095) * Fix ggerganov#4017 * Update ggml-cuda.cu Co-authored-by: Jared Van Bortel <[email protected]> * Update ggml-cuda.cu Co-authored-by: Jared Van Bortel <[email protected]> --------- Co-authored-by: Jared Van Bortel <[email protected]> finetune : zero the loraB initial vectors (ggerganov#4082) * finetune : zero the loraB initial vectors Without this, the first iteration is starting out far from the base model, instead of exactly on it. Zeroing loraB is what the paper recommends. loralib also zeroes at least one of the init vector pairs (though it departs from the paper in using a different distribution for the other vector, in some cases). * tabs to spaces * Use ggml_set_zero instead of adding a new function finetune : speed-up ggml_compute_forward_out_prod_f32 via BLAS (ggerganov#4079) * Remove logically superfluous assertions and order by dimension * Use cblas_sgemm() to implement ggml_compute_forward_out_prod() * Remove ggml_compute_forward_out_prod_use_blas(), fix compiling errors on cmake/zig, remove trailing whitespace * Add openBLAS support for sgemm() in compute_forward_out_prod() llama : add functions to get the model's metadata (ggerganov#4013) * llama : add functions to get the model's metadata * format -> std::to_string * better documentation train : move number of gpu layers argument parsing to common/train.cpp (ggerganov#4074) - introduces help entry for the argument - cuts '--gpu-layers' form in order to simplify usage and documentation. Signed-off-by: Jiri Podivin <[email protected]> Co-authored-by: Jiri Podivin <[email protected]> py : remove superfluous import statements (ggerganov#4076) Signed-off-by: Jiri Podivin <[email protected]> Co-authored-by: Jiri Podivin <[email protected]> llava : fix compilation warning that fread return value is not used (ggerganov#4069) common : improve yaml log escaping (ggerganov#4080) * logging: improve escaping in yaml output * logging: include review feedback py : Falcon HF compatibility (ggerganov#4104) Falcon HF compatibility convert : use 'model' value if it exists. This allows karpathy/tinyllamas to load (ggerganov#4089) Co-authored-by: Don Mahurin <@> examples : add tokenize (ggerganov#4039) tokenize : fix trailing whitespace build : support ppc64le build for make and CMake (ggerganov#3963) * build: support ppc64le build for make and CMake * build: keep __POWER9_VECTOR__ ifdef and extend with __powerpc64__ Co-authored-by: Georgi Gerganov <[email protected]> --------- Co-authored-by: Georgi Gerganov <[email protected]> llama : increase max nodes (ggerganov#4115) Clean up ggml-cuda.cu warnings when compiling with clang (for ROCM) (ggerganov#4124) * ggml-cuda.cu: Clean up warnings when compiling with clang * ggml-cuda.cu: Move static items into anonymous namespace * ggml-cuda.cu: Fix use of namespace start macro * Revert "ggml-cuda.cu: Fix use of namespace start macro" This reverts commit 26c1149. * Revert "ggml-cuda.cu: Move static items into anonymous namespace" This reverts commit e29757e. scripts : Remove missed baichuan convert script (ggerganov#4127) tokenize example: Respect normal add BOS token behavior (ggerganov#4126) Allow building with Makefile gguf-py : export chat templates (ggerganov#4125) * gguf-py : export chat templates * llama.cpp : escape new lines in gguf kv info prints * gguf-py : bump version * gguf-py : check chat_template type * gguf-py : initialize chat_template gitignore : tokenize common : comma should be semicolon (ggerganov#4137) server : relay error messages (ggerganov#4131) finetune : add --n-gpu-layers flag info to --help (ggerganov#4128) Revert "finetune : add --n-gpu-layers flag info to --help (ggerganov#4128)" This reverts commit 05e8301. speculative : fix prompt tokenization in speculative example (ggerganov#4025) * Support special tokens and not adding BOS to prompt in speculative * Adapt to new should_add_bos function * Ensure tgt and dft have same add_bos setting ci : add flake8 to github actions (python linting) (ggerganov#4129) Disabled rules: * E203 Whitespace before ':' - disabled because we often use 'C' Style where values are aligned * E211 Whitespace before '(' (E211) - disabled because we often use 'C' Style where values are aligned * E221 Multiple spaces before operator - disabled because we often use 'C' Style where values are aligned * E225 Missing whitespace around operator - disabled because it's broken so often it seems like a standard * E231 Missing whitespace after ',', ';', or ':' - disabled because we often use 'C' Style where values are aligned * E241 Multiple spaces after ',' - disabled because we often use 'C' Style where values are aligned * E251 Unexpected spaces around keyword / parameter equals - disabled because it's broken so often it seems like a standard * E261 At least two spaces before inline comment - disabled because it's broken so often it seems like a standard * E266 Too many leading '#' for block comment - sometimes used as "section" separator * E501 Line too long - disabled because it's broken so often it seems like a standard * E701 Multiple statements on one line (colon) - broken only in convert.py when defining abstract methods (we can use# noqa instead) * E704 Multiple statements on one line - broken only in convert.py when defining abstract methods (we can use# noqa instead) main : Add ChatML functionality to main example (ggerganov#4046) Co-authored-by: Sebastian Cramond <[email protected]> readme : update ROCm Windows instructions (ggerganov#4122) * Update README.md * Update README.md Co-authored-by: Jared Van Bortel <[email protected]> --------- Co-authored-by: Jared Van Bortel <[email protected]> finetune - update readme to mention llama support only (ggerganov#4148) stablelm : simplify + speedup generation (ggerganov#4153) docs : add llama-star arch idea examples : fix typo in parallel example doc comment (ggerganov#4181) Signed-off-by: Daniel Bevenius <[email protected]> readme : update hot topics llama : KV cache view API + better KV cache management (ggerganov#4170) * llama : keep track of used KV cells + better KV cache management * llama : zero KV cache used upon clear ggml-ci * llama : allow exporting a view of the KV cache (ggerganov#4180) * Allow exporting a view of the KV cache * Allow dumping the sequences per cell in common * Track max contiguous cells value and position as well * Fix max contiguous empty cells index calculation Make dump functions deal with lengths or sequences counts > 10 better * Fix off by one error in dump_kv_cache_view * Add doc comments for KV cache view functions Eliminate cell sequence struct; use llama_seq_id directly Minor cleanups * common : add -dkvc arg for enabling kv cache dumps --------- Co-authored-by: Kerfuffle <[email protected]> Fix incorrect format strings and uninitialized variables. (ggerganov#4133) * Fix incorrect format strings and uninitialized variables. * Address comments * Add the missing include statement readme : use PATH for Windows ROCm (ggerganov#4195) * Update README.md to use PATH for Windows ROCm * Update README.md * Update README.md main.swift : fix eos checking (ggerganov#4197) llama_token_eos(const struct llama_model *) is currently getting struct llama_context type variable context as a parameter. convert : fix tensors using grad in some models (ggerganov#4173) ggml-cuda : support stablelm rope (ggerganov#4156) * ggml-cuda : support stablelm rope * remove unused freq_base kernel parameter * add n_dims parameter to llm_build_k_shift, default to n_rot via overload * llama : fix llm_build_k_shift args --------- Co-authored-by: Georgi Gerganov <[email protected]> llama : set metal log callback correctly (ggerganov#4204) server : OAI API compatibility (ggerganov#4198) * Add openai-compatible POST /v1/chat/completions API endpoint to server example * fix code style * Update server README.md * Improve server README.md * Fix server.cpp code style according to review * server : some style changes * server : indentation * server : enable special tokens during tokenization by default * server : minor code style * server : change random string generator * straightforward /v1/models endpoint --------- Co-authored-by: kir-gadjello <[email protected]> Co-authored-by: Tobi Lütke <[email protected]> readme : update hot topics Update docs for yarn_ext_factor <0.0 as unspecified instead of NaN (ggerganov#4189) llama : grammar `reserve` space in `decode_utf8` (ggerganov#4210) * reserve space for codepoints * improvement for the appended 0 scripts : Use mmap in torch load (ggerganov#4202) * Use mmap in torch load, prefer .bin files when loading * Revert .bin > .safetensors preference metal : fix yarn (ggerganov#4220) get the correct n_orig_ctx in metal lookahead : add example for lookahead decoding (ggerganov#4207) * lookahead : init * lookahead : generate and store n-grams * lookahead : use loop instead recursion to generate n-grams * lookahead : initial working implementation * lookahead : filter repeating n-grams * lookahead : use deterministic init * lookahead : add to Makefile * lookahead : fix a bug in the seq_id of the lookahead tokens * lookahead : add comments --------- Co-authored-by: slaren <[email protected]> readme : update hot topics lookahead : support `-n -1` infinite generation ggml : fix -Warray-bounds warning with gcc (ggerganov#4231) examples : iOS example with swift ui (ggerganov#4159) * copy to llama.cpp as subdir * attempt enabling metal, fails * ggml metal compiles! * Update README.md * initial conversion to new format, utf8 errors? * bug fixes, but now has an invalid memory access :( * added O3, now has insufficient memory access * begin sync with master * update to match latest code, new errors * fixed it! * fix for loop conditionals, increase result size * fix current workflow errors * attempt a llama.swiftui workflow * Update .github/workflows/build.yml Co-authored-by: Georgi Gerganov <[email protected]> --------- Co-authored-by: Georgi Gerganov <[email protected]> readme : add Amica to UI list (ggerganov#4230) cmake : fix issue with version info not getting baked into LlamaConfig.cmake (ggerganov#3970) * Split CPP generation from build-info query * Remove blank lines * Add BUILD_SHARED_LIBS option ggml : re-enable BLAS for CPU when src0 != F32 + remove redundant full offload checks in llama.cpp (ggerganov#4240) * ggml : use blas even if src0 is not F32 * llama : use n_threads_batch only when n_tokens >= 32 ggml-ci * llama : revert n_threads_batch logic ggml-ci ggml : restore abort() in GGML_ASSERT (ggerganov#4242) readme : add FreeChat (ggerganov#4248) examples : add readme files py : fix oai proxy (ggerganov#3972) * fix oai proxy fix generation not stoped while bot stop talking in chat mode fix possible `slot_id` not exist response for cors (and pre flight) * oai proxy: workaround for some client (such as Chatbox) * use stop as separator to replace hardcoded `\n` llama : fix typical sampling (ggerganov#4261) Typical sampling was broken because after copying new_candidates into canditates, the "sorted" bool is left at "true", but the new data is no longer sorted according to probability. Patch to set "sorted" to false. Test: Generating with temp=0.0001 (approx. argmax) should generate the same sequence at typical>=1.0 and typical=0.9999 (approx. disabled, but enters the typical sampling codepath). convert.py : fix llama/llama2 conversion due to vocab_size=-1 (ggerganov#4258) llama : fix alignment of general.name in print meta (ggerganov#4254) * llama: fix alignment of general.name in print meta This commit fixes the alignment of the general.name field in the llm_load_print_meta function. Currently the output looks like this: ```console llm_load_print_meta: model ftype = mostly Q4_0 llm_load_print_meta: model params = 13.02 B llm_load_print_meta: model size = 6.86 GiB (4.53 BPW) llm_load_print_meta: general.name = LLaMA v2 ``` And with this commit it looks like this: ```console llm_load_print_meta: model ftype = mostly Q4_0 llm_load_print_meta: model params = 13.02 B llm_load_print_meta: model size = 6.86 GiB (4.53 BPW) llm_load_print_meta: general.name = LLaMA v2 ``` Signed-off-by: Daniel Bevenius <[email protected]> * llama: fix alignment of special tokens Signed-off-by: Daniel Bevenius <[email protected]> --------- Signed-off-by: Daniel Bevenius <[email protected]> readme : fix typo (ggerganov#4253) llama.cpp uses GitHub Actions, not Gitlab Actions. cmake : fix the metal file foder path (ggerganov#4217) batched.swift : update README.md (ggerganov#4214) docs: update how to run docker : add finetune option (ggerganov#4211) readme : fix (ggerganov#4135) * fix: readme * chore: resolve comments * chore: resolve comments main : pass LOG_TEE callback to llama.cpp log (ggerganov#4033) * main : Call llama_log_set to use LOG_TEE * tabs to spaces llava : ShareGPT4V compatibility (vision encoder only loading) (ggerganov#4172) * ShareGPT4 compatibility (vision encoder only loading) Load only a CLIP vision encoder (as supplied by ShareGPT finetunes) Corrects the argument parsing for --img_mean and --img_std (which were previously not parsed but attempted to access) Defines defaults for img_mean and img_std which are equal to the llava 1.5 CLIP encoder, so you do not have to provide them * Update convert-image-encoder-to-gguf.py build : fix build info generation and cleanup Makefile (ggerganov#3920) * cmake : fix joining of REAL_GIT_DIR * fix includes with help from include-what-you-use * make : remove unneeded deps and add test-rope target * fix C includes in C++ source files * Revert "fix includes with help from include-what-you-use" This reverts commit 635e9fa. make : fix Apple clang determination bug (ggerganov#4272) Co-authored-by: Will Findley <[email protected]> server : add single-client multi-prompt support (ggerganov#4232) * * add multiprompt support * * cleanup * * more cleanup * * remove atomicity of id_gen, and change lock_guard to unique_lock on completion requests * * remove all references to mutex_multitasks * Update examples/server/server.cpp Co-authored-by: Jared Van Bortel <[email protected]> * Update examples/server/server.cpp Co-authored-by: Jared Van Bortel <[email protected]> * Update examples/server/server.cpp Co-authored-by: Jared Van Bortel <[email protected]> * Update examples/server/server.cpp Co-authored-by: Jared Van Bortel <[email protected]> * * change to set --------- Co-authored-by: Jared Van Bortel <[email protected]> server : add --log-disable to disable logging to file (ggerganov#4260) * * add --log-disable to disable logging to file in the server example * * typo fix ggml : add ggml_soft_max_ext (ggerganov#4256) * metal : implement soft_max_ext * cuda : implement soft_max_ext * ggml : implement soft_max_ext (CPU) * batched-bench : print threads ggml-ci * metal : simplify soft_max encoding ggml-ci * cuda : use 512 threads for soft_max instead of 32 * ggml : update soft max cpu * cuda : do warp-based block reduce * cuda : increase max block size to 1024 * cuda : fix warp reduction initialization of shared mem * metal : warp-based reduction for soft max kernel * metal : warp-based reduce for rms_norm * metal : simplify soft max kernel ggml-ci * alloc : fix build with debug py : add requirements file for convert-hf-to-gguf.py (ggerganov#4277) This commit adds a requirements file for the convert-hf-to-gguf.py script, and also add the torch and transformers packages to it. The motivation for this is that currently running convert-hf-to-gguf.py will produce the following error: ```console $ python3 -m venv venv $ source venv/bin/activate (venv) $ pip install -r requirements.txt Collecting numpy==1.24.4 Collecting sentencepiece==0.1.98 Collecting gguf>=0.1.0 Installing collected packages: sentencepiece, numpy, gguf Successfully installed gguf-0.5.1 numpy-1.24.4 sentencepiece-0.1.98 (venv) $ python convert-hf-to-gguf.py --help Traceback (most recent call last): File "llama.cpp/convert-hf-to-gguf.py", line 16, in <module> import torch ModuleNotFoundError: No module named 'torch' ``` With this commit, and using requirements-hf-to-gguf.txt instead of requirements.txt, the script can be run and shows the help output. Signed-off-by: Daniel Bevenius <[email protected]> llama : fix integer overflow during quantization (ggerganov#4284) happens with multi-threaded quantization of Qwen-72B ggml-ci llama : add Qwen support (ggerganov#4281) * enable qwen to llama.cpp * llama : do not GPU split bias tensors --------- Co-authored-by: Georgi Gerganov <[email protected]> llama : support attention bias on LLaMA architecture (ggerganov#4283) * Support attention_bias on LLaMA architecture QKVO bias, should fix InternLM (ggerganov#3133) and works for LLaMAfied Qwen models (ggerganov#3743 (comment)). * check existence of qkvo bias while loading llama models Tested on LLaMA2, CUDA and CPU. * Update llama.cpp build : enable libstdc++ assertions for debug builds (ggerganov#4275) swift : fix token_to_piece implementation (ggerganov#4278) * Fix token_to_piece implementation in Swift * Fix errors llama : support optional tensors (ggerganov#4283) llama : avoid using "optional" keyword (ggerganov#4283) llama : pad KV cache size (ggerganov#4280) * llama : pad KV cache size to 32 * metal : try to improve batched decoding py : add grammar to oai like api (ggerganov#4294) server : fix OpenAI API `stop` field to be optional (ggerganov#4299) (cherry picked from commit Mozilla-Ocho/llamafile@e8c92bc) ggml : fix soft max out-of-bounds access (ggerganov#4307) ggml-ci ggml : reuse ggml_get_n_tasks() in ggml_graph_plan() (ggerganov#4308) * ggml : fix soft max out-of-bounds access ggml-ci * ggml : reuse ggml_get_n_tasks() in ggml_graph_plan() ggml-ci grammar-parser : fix typo (ggerganov#4318) preceeding -> preceding swift : fix prompt tokenization logic (ggerganov#4321) swift : fix concatenation method to avoid invalid UTF8 stringfication (ggerganov#4325) simple : update error message for KV cache check (ggerganov#4324) This commit updates the error message that is printed when the KV cache is not big enough to hold all the prompt and generated tokens. Specifically it removes the reference to n_parallel and replaces it with n_len. Signed-off-by: Daniel Bevenius <[email protected]> swift : revert compiler checks for swift package (ggerganov#4332) sampling : custom samplers order (ggerganov#4285) * Samplers sequence order w parameter * Cleaned commented code * Fixed formatting * Rewrote with unordered_map * Revert and rewrite, too many problems and safeguards would be needed * Fixed code style * Code style fixes according to review * More readable samplers input string, fixed help * Style fix in sampler_queue * Formatting fixes * Fixing whitespaces llama : allow overriding GGUF metadata when loading model (ggerganov#4092) * feat: Allow overriding GGUF metadata when loading model * Fix the one time GCC is stricter than clang about something * Step1 * Refactor... basically everything! * Nuke obsolete GetArrayLen struct * simplify std::string specialization * Various cleanups Add informational output when overrides are applied Warn user when an override with the wrong type is specified * Fix broken logic for parsing bool KV overrides Fix issue where overrides didn't apply when key missing in GGUF metadata Resolve merge changes * llama : rearrange model params * Update new GET_KEY call Add note that metadata KV overrides aren't reflected in initial metadata KV info dump --------- Co-authored-by: cebtenzzre <[email protected]> Co-authored-by: Georgi Gerganov <[email protected]> grammar : pre-computed pieces + reserve mem + less string copies (ggerganov#4330) * reserve space for codepoints * improvement for the appended 0 * used precomputed token text for grammar sample * reserve canidates_decoded * reserve canidates_grammar * remove candidates_decoded * Revert "remove candidates_decoded" This reverts commit 3773328. * changed decode_utf8 to take src by ref speculative : support `--color` (ggerganov#4343) * speculative: add some colors * minor : add braces --------- Co-authored-by: Georgi Gerganov <[email protected]> common : fix compile warning server : recognize cache_prompt parameter in OAI API (ggerganov#4347) train : fix ggerganov#4227 (double free in examples/train-text-from-scratch/train-text-from-scratch.cpp) (ggerganov#4351) On commit b1108 (44c117f) xaedes added ggml_allocr * alloc = NULL; ... (many lines in between) if (alloc) { ggml_allocr_free(alloc); } Which is correct, but it's easy to lose context after many lines in between. On commit b1287 (0e76a899) xaedes made a big change. From here on, alloc is freed eagerly. alloc = ggml_allocr_new(...) ... (short lines of code) ggml_allocr_free(alloc) This happens a few times, but alloc is never set to NULL, and many lines below, we still have if (alloc) { ggml_allocr_free(alloc); } which causes a double-free. llama : per-layer KV cache + quantum K cache (ggerganov#4309) * per-layer KV * remove unnecessary copies * less code duplication, offload k and v separately * llama : offload KV cache per-layer * llama : offload K shift tensors * llama : offload for rest of the model arches * llama : enable offload debug temporarily * llama : keep the KV related layers on the device * llama : remove mirrors, perform Device -> Host when partial offload * common : add command-line arg to disable KV cache offloading * llama : update session save/load * llama : support quantum K cache (ggerganov#4312) * llama : support quantum K cache (wip) * metal : add F32 -> Q8_0 copy kernel * cuda : add F32 -> Q8_0 copy kernel ggml-ci * cuda : use mmv kernel for quantum cache ops * llama : pass KV cache type through API * llama : fix build ggml-ci * metal : add F32 -> Q4_0 copy kernel * metal : add F32 -> Q4_1 copy kernel * cuda : wip * cuda : add F32 -> Q4_0 and F32 -> Q4_1 copy kernels * llama-bench : support type_k/type_v * metal : use mm kernel only for quantum KV cache * cuda : add comment * llama : remove memory_f16 and kv_f16 flags --------- Co-authored-by: slaren <[email protected]> * readme : add API change notice --------- Co-authored-by: slaren <[email protected]> sync : ggml (new ops, tests, backend, etc.) (ggerganov#4359) * sync : ggml (part 1) * sync : ggml (part 2, CUDA) * sync : ggml (part 3, Metal) * ggml : build fixes ggml-ci * cuda : restore lost changes * cuda : restore lost changes (StableLM rope) * cmake : enable separable compilation for CUDA ggml-ci * ggml-cuda : remove device side dequantize * Revert "cmake : enable separable compilation for CUDA" This reverts commit 09e35d0. * cuda : remove assert for rope * tests : add test-backend-ops * ggml : fix bug in ggml_concat * ggml : restore `ggml_get_n_tasks()` logic in `ggml_graph_plan()` * ci : try to fix macOS * ggml-backend : remove backend self-registration * ci : disable Metal for macOS cmake build ggml-ci * metal : fix "supports family" call * metal : fix assert * metal : print resource path ggml-ci --------- Co-authored-by: slaren <[email protected]> grammar : revert the replacement of llama_token_to_piece with id_to_token (ggerganov#4396) Update README.md (ggerganov#4388) Fix small typo. ggml : increased GGML_MAX_PARAMS to allow finetuning of 70b models (ggerganov#4424) server : fix local model name in server (ggerganov#4420) llama : document logits_all deprecation (ggerganov#4418) llama_context_params.logits_all is a parameter for controlling llama_eval. This documents that logits_all should not be used with llama_decode and llama_batch. build : target Windows 8 for standard mingw-w64 (ggerganov#4405) * build : target Windows 8 for standard mingw-w64 * make : fix missing console.o deps This was causing a link error with `make all` on Windows. english : use `typos` to fix comments and logs (ggerganov#4354) server : tweak default sampling parameters (ggerganov#4367) * Set a more typical Top P setting as the default * Update temp max llama : add Mixtral support (ggerganov#4406) * convert : support Mixtral as LLAMA arch * convert : fix n_ff typo * llama : model loading * ggml : sync latest ggml_mul_mat_id * llama : update graph to support MoE * llama : fix cur -> cur_expert * llama : first working version * llama : fix expert weighting in the FFN * ggml : ggml_get_rows support 2D indexing [n_tokens, n_experts] (cpu only) * ggml : add n_as argument to ggml_mul_mat_id * ggml : fix ggml_get_rows to take into account ne02 / ne11 * metal : add more general support for ggml_get_rows + tests * llama : add basic support for offloading moe with CUDA * metal : add/mul/div use general kernel when src1 not cont * metal : reduce the kernel launches for ggml_mul_mat_id * ggml : get_rows : support non-contiguos tensors with gaps, generalize up to 3D * ggml : update get_rows f16 and q * cuda : support non-contiguous src1 in get_rows * llama : offload missing ffn_moe_silu * metal : fix ggml_get_rows to work with non-cont src1 * metal : add indirect mat-vec kernels for all quantization types * llama : do not quantize expert gating tensors * llama : add n_expert and n_expert_used to hparams + change quants * test-backend-ops : add moe test * cuda : fix get_rows when ncols is odd * convert : determine n_ctx correctly * metal : fix ggml_mul_mat_id for F32 * test-backend-ops : make experts more evenly probable (test_moe) * test-backend-ops : cleanup, add moe test for batches * test-backend-ops : add cpy from f32 -> all types test * test-backend-ops : fix dequantize block offset * llama : fix hard-coded number of experts * test-backend-ops : simplify and disable slow tests to avoid CI timeout * test-backend-ops : disable MOE test with thread sanitizer * cuda : fix mul_mat_id with multi gpu * convert : use 1e6 rope_freq_base for mixtral * convert : fix style * convert : support safetensors format * gguf-py : bump version * metal : add cpy f16 -> f32 kernel * metal : fix binary ops for ne10 % 4 != 0 * test-backend-ops : add one more sum_rows test * ggml : do not use BLAS with ggml_mul_mat_id * convert-hf : support for mixtral-instruct (ggerganov#4428) * convert : typo fix, add additional hyperparameters, use LLaMA arch for Mixtral-instruct * convert : use sentencepiece tokenizer for Mixtral-instruct * convert : make flake8 happy * metal : fix soft_max kernels ref: ggerganov/ggml@1914017 * metal : limit kernels to not use more than the allowed threads --------- Co-authored-by: Georgi Gerganov <[email protected]> Co-authored-by: Radek Pilar <[email protected]>

ggerganov added 4 commits November 29, 2023 14:56

metal : implement soft_max_ext

e89597c

cuda : implement soft_max_ext

88519fb

ggml : implement soft_max_ext (CPU)

6a66f69

batched-bench : print threads

390a445

ggml-ci

ggerganov force-pushed the gg/soft-max-ext branch from b36b103 to 390a445 Compare November 29, 2023 15:26

metal : simplify soft_max encoding

580fe20

ggml-ci

ggerganov commented Nov 30, 2023

View reviewed changes

ggml-cuda.cu Outdated Show resolved Hide resolved

cuda : use 512 threads for soft_max instead of 32

ebd062b

ggerganov force-pushed the gg/soft-max-ext branch from fd70e7a to ebd062b Compare November 30, 2023 16:08

ggerganov commented Nov 30, 2023

View reviewed changes

ggml : update soft max cpu

c7c8dab

ggerganov commented Nov 30, 2023

View reviewed changes

ggerganov added 2 commits November 30, 2023 20:36

cuda : do warp-based block reduce

62532c0

cuda : increase max block size to 1024

6b86bcf

cuda : fix warp reduction initialization of shared mem

68e02c0

metal : warp-based reduction for soft max kernel

55717c9

ggerganov added 3 commits November 30, 2023 22:21

metal : warp-based reduce for rms_norm

c4db592

metal : simplify soft max kernel

d9c8fa3

ggml-ci

alloc : fix build with debug

eb594c0

ggerganov merged commit ef47ec1 into master Dec 1, 2023
36 checks passed

LostRuins mentioned this pull request Dec 2, 2023

CLBlast fails on context lengths above 2048 after merging #4256 #4296

Closed

LostRuins added a commit to LostRuins/koboldcpp that referenced this pull request Dec 2, 2023

Revert "ggml : add ggml_soft_max_ext (ggerganov#4256)"

a8e66ef

This reverts commit ef47ec1.

ggerganov mentioned this pull request Dec 3, 2023

ggml : fix soft max out-of-bounds access #4307

Merged

LostRuins added a commit to LostRuins/koboldcpp that referenced this pull request Dec 3, 2023

Revert "Revert "ggml : add ggml_soft_max_ext (ggerganov#4256)""

48544cd

This reverts commit a8e66ef.

mindkrypted mentioned this pull request Feb 14, 2024

Crippled Performance on Multi Gpu Due to Loading onto RAM #4498

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

ggml : add ggml_soft_max_ext #4256

ggml : add ggml_soft_max_ext #4256

ggerganov commented Nov 29, 2023 •

edited

Loading

slaren commented Nov 29, 2023 •

edited

Loading

ggerganov commented Nov 29, 2023

ggerganov Nov 30, 2023

slaren Nov 30, 2023

ggerganov Nov 30, 2023

ggerganov Nov 30, 2023

ggerganov left a comment

slaren commented Nov 30, 2023

ggerganov commented Nov 30, 2023

slaren commented Nov 30, 2023

LostRuins commented Dec 2, 2023


		const int block_size = blockDim.x;

		__shared__ float buf[CUDA_SOFT_MAX_BLOCK_SIZE];

ggml : add ggml_soft_max_ext #4256

ggml : add ggml_soft_max_ext #4256

Conversation

ggerganov commented Nov 29, 2023 • edited Loading

Perf CUDA V100

Perf M2 Ultra

Perf M1 Pro

slaren commented Nov 29, 2023 • edited Loading

ggerganov commented Nov 29, 2023

ggerganov Nov 30, 2023

Choose a reason for hiding this comment

slaren Nov 30, 2023

Choose a reason for hiding this comment

ggerganov Nov 30, 2023

Choose a reason for hiding this comment

ggerganov Nov 30, 2023

Choose a reason for hiding this comment

ggerganov left a comment

Choose a reason for hiding this comment

slaren commented Nov 30, 2023

ggerganov commented Nov 30, 2023

slaren commented Nov 30, 2023

LostRuins commented Dec 2, 2023

ggerganov commented Nov 29, 2023 •

edited

Loading

slaren commented Nov 29, 2023 •

edited

Loading