From cce49832c1b81b4e535e78ff308417ef3a386b18 Mon Sep 17 00:00:00 2001 From: Kawrakow Date: Wed, 2 Oct 2024 15:22:13 +0300 Subject: Adding Q6_0 (#77) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Adding q6_0 - basics + AVX2/Zen4 working * Adding q6_0: CUDA dequantize works, but not mmvq * Adding q6_0: CUDA mmvq works * Adding q6_0: CUDA cpy, so Q6_0 can be used for KV-cache * Add q6_0 to CPU flash attention Disappointing result: for LlaMA-3.2-1B, q6_0 K- and V-cache gives about the same PPL as q8_0 K-cache and q4_0 V-cache, while needing the exact same RAM. I.e., what was the point? * q6_0: slightly better kv-cache result Better than q8_0+q4_0, but not as good as q8_0+iq4_nl * q6_0: works on ARM_NEON * q6_0: dequantize works on Metal, but not vector dot product * q6_0: it now works on Metal Outperforms q5_0 by a significant margin. E.g. | model | size | params | backend | ngl | threads | test | t/s | | ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | ------------: | ---------------: | | llama 8B Q6_0 | 6.08 GiB | 8.03 B | Metal | 100 | 4 | tg128 | 44.02 ± 0.08 | | llama 8B Q5_0 | 5.21 GiB | 8.03 B | Metal | 100 | 4 | tg128 | 40.13 ± 0.12 | | llama 8B Q6_0 | 6.08 GiB | 8.03 B | Metal | 100 | 4 | pp512 | 500.55 ± 0.32 | | llama 8B Q5_0 | 5.21 GiB | 8.03 B | Metal | 100 | 4 | pp512 | 448.02 ± 0.27 | * q6_0: can now be used for kv-cache on Metal --------- Co-authored-by: Iwan Kawrakow --- ggml/src/ggml-common.h | 11 +++++++++++ 1 file changed, 11 insertions(+) (limited to 'ggml/src/ggml-common.h') diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index bb0c4864..02ecf071 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -95,6 +95,9 @@ typedef sycl::half2 ggml_half2; #define QI5_1 (QK5_1 / (4 * QR5_1)) #define QR5_1 2 +#define QI6_0 (QK6_0 / (4 * QR6_0)) +#define QR6_0 2 + #define QI8_0 (QK8_0 / (4 * QR8_0)) #define QR8_0 1 @@ -187,6 +190,14 @@ typedef struct { } block_q5_1; static_assert(sizeof(block_q5_1) == 2 * sizeof(ggml_half) + sizeof(uint32_t) + QK5_1 / 2, "wrong q5_1 block size/padding"); +#define QK6_0 32 +typedef struct { + ggml_half d; // delta + uint8_t qh[QK6_0/4]; // 5+6-th bit of quants + uint8_t qs[QK6_0/2]; // nibbles / quants +} block_q6_0; +static_assert(sizeof(block_q6_0) == sizeof(ggml_half) + QK6_0/2 + QK6_0/4, "wrong q6_0 block size/padding"); + #define QK8_0 32 typedef struct { ggml_half d; // delta -- cgit v1.2.3