From 70e7b758f5c964a3de7432c11d0edab059dbb355 Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Sat, 12 Oct 2024 16:28:46 +0300 Subject: [PATCH] iq2_ks: WIP --- ggml/src/ggml-cuda/convert.cu | 4 +- ggml/src/ggml-cuda/iqk_mmvq.cu | 4 +- ggml/src/ggml.c | 2 +- ggml/src/iqk/iqk_quantize.cpp | 130 +++++++++++++++++++++++++++------ 4 files changed, 112 insertions(+), 28 deletions(-) diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 4c6aa16e..1e4421b1 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -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; diff --git a/ggml/src/ggml-cuda/iqk_mmvq.cu b/ggml/src/ggml-cuda/iqk_mmvq.cu index 79128484..9ca219e4 100644 --- a/ggml/src/ggml-cuda/iqk_mmvq.cu +++ b/ggml/src/ggml-cuda/iqk_mmvq.cu @@ -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); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index d945a34e..a9f795ae 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -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", diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp index c377f6ce..502f4525 100644 --- a/ggml/src/iqk/iqk_quantize.cpp +++ b/ggml/src/iqk/iqk_quantize.cpp @@ -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, 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++) {