Skip to content

Commit

Permalink
ggml : preserve old Q4 and Q5 formats
Browse files Browse the repository at this point in the history
  • Loading branch information
@ggerganov
ggerganov committed May 11, 2023
1 parent e116eb6 commit f9bbbe3
Show file tree
Hide file tree
Showing 4 changed files with 62 additions and 61 deletions.
27 changes: 13 additions & 14 deletions ggml-cuda.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -86,8 +86,8 @@ static __global__ void dequantize_block_q4_0(const void * vx, float * y) {
const int x0 = (x[i].qs[j] & 0xf) - 8;
const int x1 = (x[i].qs[j] >> 4) - 8;

y[i*qk + j + 0 ] = x0*d;
y[i*qk + j + qk/2] = x1*d;
y[i*qk + 2*j + 0] = x0*d;
y[i*qk + 2*j + 1] = x1*d;
}
}

Expand All @@ -105,8 +105,8 @@ static __global__ void dequantize_block_q4_1(const void * vx, float * y) {
const int x0 = (x[i].qs[j] & 0xf);
const int x1 = (x[i].qs[j] >> 4);

y[i*qk + j + 0 ] = x0*d + m;
y[i*qk + j + qk/2] = x1*d + m;
y[i*qk + 2*j + 0] = x0*d + m;
y[i*qk + 2*j + 1] = x1*d + m;
}
}

Expand All @@ -129,8 +129,8 @@ static __global__ void dequantize_block_q5_0(const void * vx, float * y) {
const int32_t x0 = ((x[i].qs[j] & 0xf) | xh_0) - 16;
const int32_t x1 = ((x[i].qs[j] >> 4) | xh_1) - 16;

y[i*qk + j + 0 ] = x0*d;
y[i*qk + j + qk/2] = x1*d;
y[i*qk + 2*j + 0] = x0*d;
y[i*qk + 2*j + 1] = x1*d;
}
}

Expand All @@ -154,24 +154,23 @@ static __global__ void dequantize_block_q5_1(const void * vx, float * y) {
const int x0 = (x[i].qs[j] & 0xf) | xh_0;
const int x1 = (x[i].qs[j] >> 4) | xh_1;

y[i*qk + j + 0 ] = x0*d + m;
y[i*qk + j + qk/2] = x1*d + m;
y[i*qk + 2*j + 0] = x0*d + m;
y[i*qk + 2*j + 1] = x1*d + m;
}
}

static __global__ void dequantize_block_q8_0(const void * vx, float * y) {
static const int qk = QK8_0;

const block_q8_0 * x = (const block_q8_0 *) vx;

const int i = blockIdx.x;

const float d = x[i].d;

const int8_t * pp = x[i].qs;

for (int l = 0; l < QK8_0; l++) {
const int8_t vi = pp[l];

y[i*QK8_0 + l] = vi*d;
for (int j = 0; j < qk/2; ++j) {
y[i*qk + 2*j + 0] = x[i].qs[j + 0 ]*d;
y[i*qk + 2*j + 1] = x[i].qs[j + qk/2]*d;
}
}

Expand Down
1 change: 1 addition & 0 deletions ggml-opencl.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -114,6 +114,7 @@ __kernel void dequantize_row_q8_0(__global struct block_q8_0* blocks, __global f
const uint i = get_global_id(0) / 32;
const uint l = get_local_id(0);

// TODO: this is broken
result[i*32 + l] = blocks[i].qs[l] * blocks[i].d;
}

Expand Down
91 changes: 47 additions & 44 deletions ggml.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -751,8 +751,8 @@ static void quantize_row_q4_0_reference(const float * restrict x, block_q4_0 * r
y[i].d = d;

for (int j = 0; j < qk/2; ++j) {
const float x0 = x[i*qk + 0 + j]*id;
const float x1 = x[i*qk + qk/2 + j]*id;
const float x0 = x[i*qk + 2*j + 0]*id;
const float x1 = x[i*qk + 2*j + 1]*id;

const uint8_t xi0 = MIN(15, (int8_t)(x0 + 8.5f));
const uint8_t xi1 = MIN(15, (int8_t)(x1 + 8.5f));
Expand Down Expand Up @@ -792,8 +792,8 @@ static void quantize_row_q4_1_reference(const float * restrict x, block_q4_1 * r
y[i].m = min;

for (int j = 0; j < qk/2; ++j) {
const float x0 = (x[i*qk + 0 + j] - min)*id;
const float x1 = (x[i*qk + qk/2 + j] - min)*id;
const float x0 = (x[i*qk + 2*j + 0] - min)*id;
const float x1 = (x[i*qk + 2*j + 1] - min)*id;

const uint8_t xi0 = MIN(15, (int8_t)(x0 + 0.5f));
const uint8_t xi1 = MIN(15, (int8_t)(x1 + 0.5f));
Expand Down Expand Up @@ -835,8 +835,8 @@ static void quantize_row_q5_0_reference(const float * restrict x, block_q5_0 * r
uint32_t qh = 0;

for (int j = 0; j < qk/2; ++j) {
const float x0 = x[i*qk + 0 + j]*id;
const float x1 = x[i*qk + qk/2 + j]*id;
const float x0 = x[i*qk + 2*j + 0]*id;
const float x1 = x[i*qk + 2*j + 1]*id;

const uint8_t xi0 = MIN(31, (int8_t)(x0 + 16.5f));
const uint8_t xi1 = MIN(31, (int8_t)(x1 + 16.5f));
Expand Down Expand Up @@ -883,8 +883,8 @@ static void quantize_row_q5_1_reference(const float * restrict x, block_q5_1 * r
uint32_t qh = 0;

for (int j = 0; j < qk/2; ++j) {
const float x0 = (x[i*qk + 0 + j] - min)*id;
const float x1 = (x[i*qk + qk/2 + j] - min)*id;
const float x0 = (x[i*qk + 2*j + 0] - min)*id;
const float x1 = (x[i*qk + 2*j + 1] - min)*id;

const uint8_t xi0 = (uint8_t)(x0 + 0.5f);
const uint8_t xi1 = (uint8_t)(x1 + 0.5f);
Expand Down Expand Up @@ -922,10 +922,12 @@ static void quantize_row_q8_0_reference(const float * restrict x, block_q8_0 * r

y[i].d = d;

for (int j = 0; j < QK8_0; ++j) {
const float v0 = x[i*QK8_0 + j]*id;
for (int j = 0; j < QK8_0/2; ++j) {
const float v0 = x[i*QK8_0 + 2*j + 0]*id;
const float v1 = x[i*QK8_0 + 2*j + 1]*id;

y[i].qs[j] = roundf(v0);
y[i].qs[ j] = v0 + 0.5f;
y[i].qs[QK8_0/2 + j] = v1 + 0.5f;
}
}
}
Expand All @@ -943,12 +945,12 @@ static void quantize_row_q8_0(const float * restrict x, void * restrict vy, int
float32x4_t asrcv[8];
float32x4_t amaxv[8];

for (int l = 0; l < 8; l++) srcv[l] = vld1q_f32(x + i*32 + 4*l);
for (int l = 0; l < 8; l++) asrcv[l] = vabsq_f32(srcv[l]);
for (int j = 0; j < 8; j++) srcv[j] = vld1q_f32(x + i*32 + 4*j);
for (int j = 0; j < 8; j++) asrcv[j] = vabsq_f32(srcv[j]);

for (int l = 0; l < 4; l++) amaxv[2*l] = vmaxq_f32(asrcv[2*l], asrcv[2*l+1]);
for (int l = 0; l < 2; l++) amaxv[4*l] = vmaxq_f32(amaxv[4*l], amaxv[4*l+2]);
for (int l = 0; l < 1; l++) amaxv[8*l] = vmaxq_f32(amaxv[8*l], amaxv[8*l+4]);
for (int j = 0; j < 4; j++) amaxv[2*j] = vmaxq_f32(asrcv[2*j], asrcv[2*j+1]);
for (int j = 0; j < 2; j++) amaxv[4*j] = vmaxq_f32(amaxv[4*j], amaxv[4*j+2]);
for (int j = 0; j < 1; j++) amaxv[8*j] = vmaxq_f32(amaxv[8*j], amaxv[8*j+4]);

const float amax = vmaxvq_f32(amaxv[0]);

Expand All @@ -957,14 +959,14 @@ static void quantize_row_q8_0(const float * restrict x, void * restrict vy, int

y[i].d = d;

for (int l = 0; l < 8; l++) {
const float32x4_t v = vmulq_n_f32(srcv[l], id);
for (int j = 0; j < 8; j++) {
const float32x4_t v = vmulq_n_f32(srcv[j], id);
const int32x4_t vi = vcvtnq_s32_f32(v);

y[i].qs[4*l + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[4*l + 1] = vgetq_lane_s32(vi, 1);
y[i].qs[4*l + 2] = vgetq_lane_s32(vi, 2);
y[i].qs[4*l + 3] = vgetq_lane_s32(vi, 3);
y[i].qs[ 2*j + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[16 + 2*j + 0] = vgetq_lane_s32(vi, 1);
y[i].qs[ 2*j + 1] = vgetq_lane_s32(vi, 2);
y[i].qs[16 + 2*j + 1] = vgetq_lane_s32(vi, 3);
}
}
#elif defined(__AVX2__) || defined(__AVX__)
Expand Down Expand Up @@ -1080,11 +1082,11 @@ static void quantize_row_q8_1_reference(const float * restrict x, block_q8_1 * r
int sum1 = 0;

for (int j = 0; j < QK8_1/2; ++j) {
const float v0 = x[i*QK8_1 + j]*id;
const float v1 = x[i*QK8_1 + QK8_1/2 + j]*id;
const float v0 = x[i*QK8_1 + 2*j + 0]*id;
const float v1 = x[i*QK8_1 + 2*j + 1]*id;

y[i].qs[ j] = roundf(v0);
y[i].qs[QK8_1/2 + j] = roundf(v1);
y[i].qs[ j] = v0 + 0.5f;
y[i].qs[QK8_1/2 + j] = v1 + 0.5f;

sum0 += y[i].qs[ j];
sum1 += y[i].qs[QK8_1/2 + j];
Expand Down Expand Up @@ -1129,10 +1131,10 @@ static void quantize_row_q8_1(const float * restrict x, void * restrict vy, int
const float32x4_t v = vmulq_n_f32(srcv[j], id);
const int32x4_t vi = vcvtnq_s32_f32(v);

y[i].qs[4*j + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[4*j + 1] = vgetq_lane_s32(vi, 1);
y[i].qs[4*j + 2] = vgetq_lane_s32(vi, 2);
y[i].qs[4*j + 3] = vgetq_lane_s32(vi, 3);
y[i].qs[ 2*j + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[16 + 2*j + 0] = vgetq_lane_s32(vi, 1);
y[i].qs[ 2*j + 1] = vgetq_lane_s32(vi, 2);
y[i].qs[16 + 2*j + 1] = vgetq_lane_s32(vi, 3);

accv0 = vaddq_s32(accv0, vi);
}
Expand All @@ -1142,10 +1144,10 @@ static void quantize_row_q8_1(const float * restrict x, void * restrict vy, int
const float32x4_t v = vmulq_n_f32(srcv[j], id);
const int32x4_t vi = vcvtnq_s32_f32(v);

y[i].qs[4*j + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[4*j + 1] = vgetq_lane_s32(vi, 1);
y[i].qs[4*j + 2] = vgetq_lane_s32(vi, 2);
y[i].qs[4*j + 3] = vgetq_lane_s32(vi, 3);
y[i].qs[ 2*j + 0] = vgetq_lane_s32(vi, 0);
y[i].qs[16 + 2*j + 0] = vgetq_lane_s32(vi, 1);
y[i].qs[ 2*j + 1] = vgetq_lane_s32(vi, 2);
y[i].qs[16 + 2*j + 1] = vgetq_lane_s32(vi, 3);

accv1 = vaddq_s32(accv1, vi);
}
Expand Down Expand Up @@ -1271,8 +1273,8 @@ static void dequantize_row_q4_0(const block_q4_0 * restrict x, float * restrict
const int x0 = (x[i].qs[j] & 0x0F) - 8;
const int x1 = (x[i].qs[j] >> 4) - 8;

y[i*qk + j + 0 ] = x0*d;
y[i*qk + j + qk/2] = x1*d;
y[i*qk + 2*j + 0] = x0*d;
y[i*qk + 2*j + 1] = x1*d;
}
}
}
Expand All @@ -1292,8 +1294,8 @@ static void dequantize_row_q4_1(const block_q4_1 * restrict x, float * restrict
const int x0 = (x[i].qs[j] & 0x0F);
const int x1 = (x[i].qs[j] >> 4);

y[i*qk + j + 0 ] = x0*d + m;
y[i*qk + j + qk/2] = x1*d + m;
y[i*qk + 2*j + 0] = x0*d + m;
y[i*qk + 2*j + 1] = x1*d + m;
}
}
}
Expand All @@ -1318,8 +1320,8 @@ static void dequantize_row_q5_0(const block_q5_0 * restrict x, float * restrict
const int32_t x0 = ((x[i].qs[j] & 0x0F) | xh_0) - 16;
const int32_t x1 = ((x[i].qs[j] >> 4) | xh_1) - 16;

y[i*qk + j + 0 ] = x0*d;
y[i*qk + j + qk/2] = x1*d;
y[i*qk + 2*j + 0] = x0*d;
y[i*qk + 2*j + 1] = x1*d;
}
}
}
Expand All @@ -1345,8 +1347,8 @@ static void dequantize_row_q5_1(const block_q5_1 * restrict x, float * restrict
const int x0 = (x[i].qs[j] & 0x0F) | xh_0;
const int x1 = (x[i].qs[j] >> 4) | xh_1;

y[i*qk + j + 0 ] = x0*d + m;
y[i*qk + j + qk/2] = x1*d + m;
y[i*qk + 2*j + 0] = x0*d + m;
y[i*qk + 2*j + 1] = x1*d + m;
}
}
}
Expand All @@ -1363,8 +1365,9 @@ static void dequantize_row_q8_0(const void * restrict vx, float * restrict y, in
for (int i = 0; i < nb; i++) {
const float d = x[i].d;

for (int j = 0; j < qk; ++j) {
y[i*qk + j] = x[i].qs[j]*d;
for (int j = 0; j < qk/2; ++j) {
y[i*qk + 2*j + 0] = x[i].qs[j + 0 ]*d;
y[i*qk + 2*j + 1] = x[i].qs[j + qk/2]*d;
}
}
}
Expand Down
4 changes: 1 addition & 3 deletions llama.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -919,9 +919,7 @@ static void llama_model_load_internal(
}

if (file_version != LLAMA_FILE_VERSION_GGJT_V2) {
if (hparams.ftype != LLAMA_FTYPE_ALL_F32 &&
hparams.ftype != LLAMA_FTYPE_MOSTLY_F16 &&
hparams.ftype != LLAMA_FTYPE_MOSTLY_Q8_0) {
if (hparams.ftype == LLAMA_FTYPE_MOSTLY_Q8_0) {
throw format("this format is no longer supported (see https://github.com/ggerganov/llama.cpp/pull/1305)");
}
}
Expand Down

0 comments on commit f9bbbe3

Please sign in to comment.