ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-03-03 10:30:27 +00:00
Code Issues Packages Projects Releases Wiki Activity

iq2_ks: WIP

Browse Source
This commit is contained in:
Iwan Kawrakow
2024-10-12 16:28:46 +03:00
parent aa36d90684
commit 70e7b758f5
4 changed files with 112 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
ggml/src/ggml-cuda/convert.cu
View File

@@ -749,8 +749,8 @@ static __global__ void dequantize_block_iq2_ks(const void * __restrict__ vx, dst
int64_t ii = blockIdx.x;
int64_t row = (QK_K * ii) / n_per_row;
const char * cx = (const char *)vx + row * row_size;
const float d = *(const float *)cx;
const block_iq2_ks * x = (const block_iq2_ks *)(cx + sizeof(float));
const float d = (float)*(const half *)cx;
const block_iq2_ks * x = (const block_iq2_ks *)(cx + sizeof(half));
const int64_t i = ii - (row*n_per_row)/QK_K;
const int tid = threadIdx.x;

4
ggml/src/ggml-cuda/iqk_mmvq.cu
View File

@@ -460,8 +460,8 @@ __device__ __forceinline__ float vec_dot_iq2_k_q8_1(
__device__ __forceinline__ float vec_dot_iq2_ks_q8_1(
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
float scale = *(const float *)vbq;
const block_iq2_ks * bq2 = (const block_iq2_ks *)((const char *)vbq + sizeof(float)) + kbx;
float scale = *(const half *)vbq;
const block_iq2_ks * bq2 = (const block_iq2_ks *)((const char *)vbq + sizeof(half)) + kbx;
int i4 = iqs/4; // 0...7. We will process q8 blocks 4*(i4/4), 4*(i4/4)+1, 4*(i4/4)+2, 4*(i4/4)+3
const int32_t * q8_1 = (const int *)bq8_1[4*(i4/4)+0].qs + 2*(i4%4);

2
ggml/src/ggml.c
View File

@@ -1204,7 +1204,7 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
.vec_dot = vec_dot_iq2_ks_q8_k,
.vec_dot_type = GGML_TYPE_Q8_K,
.nrows = 1,
.row_meta_size = 4,
.row_meta_size = 2,
},
[GGML_TYPE_IQ3_K] = {
.type_name = "iq3_k",

130
ggml/src/iqk/iqk_quantize.cpp
View File

@@ -759,9 +759,13 @@ namespace {
void quantize_row_iq2_ks_impl(const float * x, void * vy, int n_per_row, const float * quant_weights, float * all_scales, float * all_sw, int8_t * all_Ls) {
constexpr int kBlockSize = 32;
constexpr int kMax_i1 = 3*kBlockSize/4;
constexpr int kMin_i3 = kBlockSize/4;
//constexpr int kNtry = 5;
//constexpr float kStep = 1.f;
float * dptr = (float *)vy;
*dptr = 0;
ggml_half * dptr = (ggml_half *)vy;
*dptr = GGML_FP32_TO_FP16(0.f);
block_iq2_ks * y = (block_iq2_ks *)(dptr + 1);
@@ -770,6 +774,9 @@ void quantize_row_iq2_ks_impl(const float * x, void * vy, int n_per_row, const f
std::array<std::pair<float,int>, kBlockSize> pairs;
float val [4] = {float(iq2nl_values[0]), float(iq2nl_values[1]), float(iq2nl_values[2]), float(iq2nl_values[3])};
float sval[4] = {float(iq2nl_values[4]), float(iq2nl_values[5]), float(iq2nl_values[6]), float(iq2nl_values[7])};
const int8_t * shifted_values = iq2nl_values + 4;
const int nblock = n_per_row/QK_K;
@@ -801,6 +808,83 @@ void quantize_row_iq2_ks_impl(const float * x, void * vy, int n_per_row, const f
sw[ib] += weight[j];
pairs[j] = {xb[j], j};
}
//float amax = 0, max = 0;
//for (int j = 0; j < kBlockSize; ++j) {
// float ax = fabsf(xb[j]);
// if (ax > amax) {
// amax = ax; max = xb[j];
// }
//}
//if (!amax) {
// scales[ib] = 0;
// continue;
//}
//float d = kNtry > 0 ? -max/iq2nl_values[0] : max/iq2nl_values[0];
//float id = 1/d;
//float sumqx_p = 0, sumq2_p = 0;
//float sumqx_m = 0, sumq2_m = 0;
//for (int j = 0; j < kBlockSize; ++j) {
// float w = weight[j];
// float al = id*xb[j];
// int l = best_index_iq2nl(iq2nl_values, al);
// float q = iq2nl_values[l];
// sumqx_p += w*q*xb[j];
// sumq2_p += w*q*q;
// l = best_index_iq2nl(iq2nl_values, -al);
// q = iq2nl_values[l];
// sumqx_m += w*q*xb[j];
// sumq2_m += w*q*q;
//}
//d = sumqx_p/sumq2_p;
//float best = d*sumqx_p;
//if (sumq2_m > 0 && sumqx_m*sumqx_m > best*sumq2_m) {
// d = sumqx_m/sumq2_m; best = d*sumqx_m;
//}
//bool is_shifted = false;
//for (int itry = -kNtry; itry <= kNtry; ++itry) {
// id = (kStep*itry + iq2nl_values[0])/max;
// sumqx_p = sumq2_p = 0;
// sumqx_m = sumq2_m = 0;
// for (int j = 0; j < kBlockSize; ++j) {
// float w = weight[j];
// float al = id*xb[j];
// int l = best_index_iq2nl(iq2nl_values, al);
// float q = iq2nl_values[l];
// sumqx_p += w*q*xb[j];
// sumq2_p += w*q*q;
// l = best_index_iq2nl(iq2nl_values, -al);
// q = iq2nl_values[l];
// sumqx_m += w*q*xb[j];
// sumq2_m += w*q*q;
// }
// if (sumq2_p > 0 && sumqx_p*sumqx_p > best*sumq2_p) {
// d = sumqx_p/sumq2_p; best = d * sumqx_p; is_shifted = false;
// }
// if (sumq2_m > 0 && sumqx_m*sumqx_m > best*sumq2_m) {
// d = sumqx_m/sumq2_m; best = d * sumqx_m; is_shifted = false;
// }
// id = (kStep*itry + shifted_values[0])/max;
// sumqx_p = sumq2_p = 0;
// sumqx_m = sumq2_m = 0;
// for (int j = 0; j < kBlockSize; ++j) {
// float w = weight[j];
// float al = id*xb[j];
// int l = best_index_iq2nl(shifted_values, al);
// float q = shifted_values[l];
// sumqx_p += w*q*xb[j];
// sumq2_p += w*q*q;
// l = best_index_iq2nl(shifted_values, -al);
// q = shifted_values[l];
// sumqx_m += w*q*xb[j];
// sumq2_m += w*q*q;
// }
// if (sumq2_p > 0 && sumqx_p*sumqx_p > best*sumq2_p) {
// d = sumqx_p/sumq2_p; best = d * sumqx_p; is_shifted = true;
// }
// if (sumq2_m > 0 && sumqx_m*sumqx_m > best*sumq2_m) {
// d = sumqx_m/sumq2_m; best = d * sumqx_m; is_shifted = true;
// }
//}
std::sort(pairs.begin(), pairs.end());
sumx[0] = sumw[0] = 0;
for (int j = 0; j < kBlockSize; ++j) {
@@ -811,34 +895,34 @@ void quantize_row_iq2_ks_impl(const float * x, void * vy, int n_per_row, const f
float best = 0, d = 0;
bool is_shifted = false;
float sumqx, sumq2;
for (int i1 = 0; i1 < kBlockSize; ++i1) {
for (int i1 = 0; i1 < kMax_i1; ++i1) {
for (int i2 = i1; i2 < kBlockSize; ++i2) {
for (int i3 = i2; i3 < kBlockSize; ++i3) {
sumqx = (sumx[i1] - sumx[ 0])*iq2nl_values[0] + (sumx[i2] - sumx[i1])*iq2nl_values[1]
+ (sumx[i3] - sumx[i2])*iq2nl_values[2] + (sumx[kBlockSize] - sumx[i3])*iq2nl_values[3];
sumq2 = (sumw[i1] - sumw[ 0])*iq2nl_values[0]*iq2nl_values[0] + (sumw[i2] - sumw[i1])*iq2nl_values[1]*iq2nl_values[1]
+ (sumw[i3] - sumw[i2])*iq2nl_values[2]*iq2nl_values[2] + (sumw[kBlockSize] - sumw[i3])*iq2nl_values[3]*iq2nl_values[3];
for (int i3 = std::max(i2, kMin_i3); i3 < kBlockSize; ++i3) {
sumqx = (sumx[i1] - sumx[ 0])*val[0] + (sumx[i2] - sumx[i1])*val[1]
+ (sumx[i3] - sumx[i2])*val[2] + (sumx[kBlockSize] - sumx[i3])*val[3];
sumq2 = (sumw[i1] - sumw[ 0])*val[0]*val[0] + (sumw[i2] - sumw[i1])*val[1]*val[1]
+ (sumw[i3] - sumw[i2])*val[2]*val[2] + (sumw[kBlockSize] - sumw[i3])*val[3]*val[3];
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
d = sumqx/sumq2; best = d*sumqx; is_shifted = false;
}
sumqx = (sumx[i1] - sumx[ 0])*shifted_values[0] + (sumx[i2] - sumx[i1])*shifted_values[1]
+ (sumx[i3] - sumx[i2])*shifted_values[2] + (sumx[kBlockSize] - sumx[i3])*shifted_values[3];
sumq2 = (sumw[i1] - sumw[ 0])*shifted_values[0]*shifted_values[0] + (sumw[i2] - sumw[i1])*shifted_values[1]*shifted_values[1]
+ (sumw[i3] - sumw[i2])*shifted_values[2]*shifted_values[2] + (sumw[kBlockSize] - sumw[i3])*shifted_values[3]*shifted_values[3];
sumqx = (sumx[i1] - sumx[ 0])*sval[0] + (sumx[i2] - sumx[i1])*sval[1]
+ (sumx[i3] - sumx[i2])*sval[2] + (sumx[kBlockSize] - sumx[i3])*sval[3];
sumq2 = (sumw[i1] - sumw[ 0])*sval[0]*sval[0] + (sumw[i2] - sumw[i1])*sval[1]*sval[1]
+ (sumw[i3] - sumw[i2])*sval[2]*sval[2] + (sumw[kBlockSize] - sumw[i3])*sval[3]*sval[3];
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
d = sumqx/sumq2; best = d*sumqx; is_shifted = true;
}
sumqx = (sumx[i1] - sumx[ 0])*iq2nl_values[3] + (sumx[i2] - sumx[i1])*iq2nl_values[2]
+ (sumx[i3] - sumx[i2])*iq2nl_values[1] + (sumx[kBlockSize] - sumx[i3])*iq2nl_values[0];
sumq2 = (sumw[i1] - sumw[ 0])*iq2nl_values[3]*iq2nl_values[3] + (sumw[i2] - sumw[i1])*iq2nl_values[2]*iq2nl_values[2]
+ (sumw[i3] - sumw[i2])*iq2nl_values[1]*iq2nl_values[1] + (sumw[kBlockSize] - sumw[i3])*iq2nl_values[0]*iq2nl_values[0];
sumqx = (sumx[i1] - sumx[ 0])*val[3] + (sumx[i2 ] - sumx[i1])*val[2]
+ (sumx[i3] - sumx[i2])*val[1] + (sumx[kBlockSize] - sumx[i3])*val[0];
sumq2 = (sumw[i1] - sumw[ 0])*val[3]*val[3] + (sumw[i2 ] - sumw[i1])*val[2]*val[2]
+ (sumw[i3] - sumw[i2])*val[1]*val[1] + (sumw[kBlockSize] - sumw[i3])*val[0]*val[0];
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
d = sumqx/sumq2; best = d*sumqx; is_shifted = false;
}
sumqx = (sumx[i1] - sumx[ 0])*shifted_values[3] + (sumx[i2] - sumx[i1])*shifted_values[2]
+ (sumx[i3] - sumx[i2])*shifted_values[1] + (sumx[kBlockSize] - sumx[i3])*shifted_values[0];
sumq2 = (sumw[i1] - sumw[ 0])*shifted_values[3]*shifted_values[3] + (sumw[i2] - sumw[i1])*shifted_values[2]*shifted_values[2]
+ (sumw[i3] - sumw[i2])*shifted_values[1]*shifted_values[1] + (sumw[kBlockSize] - sumw[i3])*shifted_values[0]*shifted_values[0];
sumqx = (sumx[i1] - sumx[ 0])*sval[3] + (sumx[i2 ] - sumx[i1])*sval[2]
+ (sumx[i3] - sumx[i2])*sval[1] + (sumx[kBlockSize] - sumx[i3])*sval[0];
sumq2 = (sumw[i1] - sumw[ 0])*sval[3]*sval[3] + (sumw[i2 ] - sumw[i1])*sval[2]*sval[2]
+ (sumw[i3] - sumw[i2])*sval[1]*sval[1] + (sumw[kBlockSize] - sumw[i3])*sval[0]*sval[0];
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
d = sumqx/sumq2; best = d*sumqx; is_shifted = true;
}
@@ -893,7 +977,7 @@ void quantize_row_iq2_ks_impl(const float * x, void * vy, int n_per_row, const f
}
}
}
*dptr = 1.030f*(sumq2 > 0 ? sumqx/sumq2 : d);
*dptr = GGML_FP32_TO_FP16(1.030f*(sumq2 > 0 ? sumqx/sumq2 : d));
}
}
@@ -928,8 +1012,8 @@ void dequantize_row_iq2_ks(const block_iq2_ks * GGML_RESTRICT x, float * GGML_R
assert(k % QK_K == 0);
const int nb = k / QK_K;
const float * dptr = (const float *)x;
const float d = *dptr;
const ggml_half * dptr = (const ggml_half *)x;
const float d = GGML_FP16_TO_FP32(*dptr);
x = (const block_iq2_ks *)(dptr + 1);
for (int i = 0; i < nb; i++) {