nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Even better Q4_0 KV cache (#1547)

This commit is contained in:
Kawrakow 2026-03-30 10:06:54 +02:00 committed by GitHub
parent 028fc79710
commit b9a2ce4677
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 14 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
ggml/src/ggml-cuda/cpy-utils.cuh
View File

@ -39,16 +39,28 @@ static __device__ void quantize_f32_q4_0_block(const float * __restrict__ x, blo
y->d = d;
float sumqx = 0, sumq2 = 0;
for (int j = 0; j < QK4_0/2; ++j) {
const float x0 = x[0 + j]*id;
const float x1 = x[QK4_0/2 + j]*id;
const float v0 = x[0 + j];
const float v1 = x[QK4_0/2 + j];
const float x0 = v0*id;
const float x1 = v1*id;
const uint8_t xi0 = min(15, (int8_t)(x0 + 8.5f));
const uint8_t xi1 = min(15, (int8_t)(x1 + 8.5f));
float q0 = xi0 - 8;
float q1 = xi1 - 8;
float w0 = v0*v0;
float w1 = v1*v1;
sumqx += w0*q0*v0 + w1*q1*v1;
sumq2 += w0*q0*q0 + w1*q1*q1;
y->qs[j] = xi0;
y->qs[j] |= xi1 << 4;
}
if (sumq2 > 0) {
y->d = sumqx/sumq2;
}
}
static __device__ void quantize_f32_q4_1_block(const float * __restrict__ x, block_q4_1 * __restrict__ y) {