From 82380acf103c825b9e266db0e1a650489e79582d Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Mon, 11 Mar 2024 13:12:33 +0200 Subject: [PATCH] iq1_s: we can do even better Spent one of the 4 scale bits on a signs of a 0.125 shift. I.e., quants are now -1 + delta, delta, 1 + delta, where delta is +/- 0.125. CUDA works, same performance as before. PPL(LLaMA-v2-7B) is now 11.85! --- ggml-common.h | 1 + ggml-cuda.cu | 36 ++++++++++++++++-------------------- ggml-quants.c | 40 +++++++++++++++++++++++++++------------- 3 files changed, 44 insertions(+), 33 deletions(-) diff --git a/ggml-common.h b/ggml-common.h index 5dd918081..2402b773c 100644 --- a/ggml-common.h +++ b/ggml-common.h @@ -645,6 +645,7 @@ GGML_TABLE_BEGIN(uint32_t, iq3s_grid, 512) GGML_TABLE_END() #define NGRID_IQ1S 2048 +#define IQ1S_DELTA 0.125f #if defined(GGML_COMMON_IMPL_C) GGML_TABLE_BEGIN(uint64_t, iq1s_grid, NGRID_IQ1S) 0xffffffffffffffff, 0xffffffffffffff01, 0xffffffffffff0000, 0xffffffffffff01ff, diff --git a/ggml-cuda.cu b/ggml-cuda.cu index d2945d3c2..01b1f15ee 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -1722,22 +1722,15 @@ static __global__ void dequantize_block_iq1_s(const void * __restrict__ vx, dst_ const int il = tid/8; // 0...3 const int ib = tid%8; // 0...7 dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const float d = (float)x[i].d * (2*((x[i].qh[ib] >> 12) & 0xf) + 1); -#if __CUDA_ARCH__ >= MIN_CC_DP4A // lowest compute capability for integer intrinsics - int grid32[2]; const int8_t * q = (const int8_t *)grid32; - grid32[0] = *((const int *)(iq1s_grid_gpu + (x[i].qs[4*ib+il] | (((x[i].qh[ib] >> 3*il) & 7) << 8)))); - grid32[1] = __vsub4((grid32[0] >> 4) & 0x0f0f0f0f, 0x01010101); - grid32[0] = __vsub4(grid32[0] & 0x0f0f0f0f, 0x01010101); + const float delta = x[i].qh[ib] & 0x8000 ? -1 - IQ1S_DELTA : -1 + IQ1S_DELTA; + const float d = (float)x[i].d * (2*((x[i].qh[ib] >> 12) & 7) + 1); + uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32; + grid32[0] = iq1s_grid_gpu[x[i].qs[4*ib+il] | (((x[i].qh[ib] >> 3*il) & 7) << 8)]; + grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f; + grid32[0] &= 0x0f0f0f0f; for (int j = 0; j < 8; ++j) { - y[j] = d * q[j]; + y[j] = d * (q[j] + delta); } -#else - const uint8_t * grid = (const uint8_t *)(iq1s_grid_gpu + (x[i].qs[4*ib+il] | (((x[i].qh[ib] >> 3*il) & 7) << 8))); - for (int j = 0; j < 4; ++j) { - y[j+0] = d * ((grid[j] & 0xf) - 1); - y[j+4] = d * ((grid[j] >> 4) - 1); - } -#endif #else assert(false); #endif @@ -4560,22 +4553,25 @@ static __device__ __forceinline__ float vec_dot_iq1_s_q8_1( const int * q8 = (const int *)bq8_1[ib32].qs; for (int l = 0; l < 4; ++l) { const int * grid = (const int *)(iq1s_grid_gpu + (bq1->qs[4*ib32+l] | (((bq1->qh[ib32] >> 3*l) & 7) << 8))); - int grid0 = __vsub4(grid[0] & 0x0f0f0f0f, 0x01010101); - int grid1 = __vsub4((grid[0] >> 4) & 0x0f0f0f0f, 0x01010101); + int grid0 = grid[0] & 0x0f0f0f0f; + int grid1 = (grid[0] >> 4) & 0x0f0f0f0f; sumi = __dp4a(q8[2*l+1], grid1, __dp4a(q8[2*l+0], grid0, sumi)); } #else - const int8_t * q8 = bq8_1[ib32].qs; + const int8_t * q8 = bq8_1[ib32].qs; for (int l = 0; l < 4; ++l) { const uint8_t * grid = (const uint8_t *)(iq1s_grid_gpu + (bq1->qs[4*ib32+l] | (((bq1->qh[ib32] >> 3*l) & 7) << 8))); for (int j = 0; j < 4; ++j) { - sumi += q8[j] * ((grid[j] & 0xf) - 1) + q8[j+4] * ((grid[j] >> 4) - 1); + sumi += q8[j] * (grid[j] & 0xf) + q8[j+4] * (grid[j] >> 4); } q8 += 8; } #endif - const float d = (float)bq1->d * __low2float(bq8_1[ib32].ds); - return d * sumi * (2*(bq1->qh[ib32] >> 12) + 1); + const float delta = bq1->qh[ib32] & 0x8000 ? -1-IQ1S_DELTA : -1+IQ1S_DELTA; + const float d1q = (float)bq1->d * (2*((bq1->qh[ib32] >> 12) & 7) + 1); + const float d = d1q * __low2float (bq8_1[ib32].ds); + const float m = d1q * __high2float(bq8_1[ib32].ds); + return d * sumi + m * delta; #else assert(false); return 0.f; diff --git a/ggml-quants.c b/ggml-quants.c index f9a3d9fd2..a269003f6 100644 --- a/ggml-quants.c +++ b/ggml-quants.c @@ -11438,7 +11438,7 @@ static int iq1_find_best_neighbour(const uint16_t * restrict neighbours, const u } static int iq1_find_best_neighbour2(const uint16_t * restrict neighbours, const uint64_t * restrict grid, - const float * restrict xval, const float * restrict weight, float scale, int8_t * restrict L, int ngrid) { + const float * restrict xval, const float * restrict weight, float scale, const float * restrict xg, int8_t * restrict L, int ngrid) { int num_neighbors = neighbours[0]; GGML_ASSERT(num_neighbors > 0); float best_score = FLT_MAX; @@ -11447,7 +11447,7 @@ static int iq1_find_best_neighbour2(const uint16_t * restrict neighbours, const const int8_t * pg = (const int8_t *)(grid + neighbours[j]); float d2 = 0; for (int i = 0; i < 8; ++i) { - float q = (pg[i] - 3)/2; + float q = xg[(pg[i] - 1)/2]; float w = weight[i]; float diff = scale*q - xval[i]; d2 += w*diff*diff; @@ -11463,7 +11463,7 @@ static int iq1_find_best_neighbour2(const uint16_t * restrict neighbours, const float d2 = 0; for (int j = 0; j < 8; ++j) { float w = weight[j]; - float q = (grid_i[j] - 3)/2; + float q = xg[(grid_i[j] - 1)/2]; float diff = scale*q - xval[i]; d2 += w*diff*diff; } @@ -11480,7 +11480,7 @@ static int iq1_find_best_neighbour2(const uint16_t * restrict neighbours, const const int8_t * pg = (const int8_t *)(grid + neighbours[j]); float sumqx = 0, sumq2 = 0; for (int i = 0; i < 8; ++i) { - float q = (pg[i] - 3)/2; + float q = xg[(pg[i] - 1)/2]; float w = weight[i]; sumqx += w*q*xval[i]; sumq2 += w*q*q; @@ -11519,6 +11519,9 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy block_iq1_s * y = vy; + const float x_p[3] = {-1 + IQ1S_DELTA, IQ1S_DELTA, 1 + IQ1S_DELTA}; + const float x_m[3] = {-1 - IQ1S_DELTA, -IQ1S_DELTA, 1 - IQ1S_DELTA}; + float scales[QK_K/IQ1S_BLOCK_SIZE]; float weight[IQ1S_BLOCK_SIZE]; int8_t L[IQ1S_BLOCK_SIZE]; @@ -11527,6 +11530,7 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy float pairs[2*IQ1S_BLOCK_SIZE]; int * idx = (int *)(pairs + 1); uint16_t index[IQ1S_BLOCK_SIZE/8]; + int8_t shifts[QK_K/IQ1S_BLOCK_SIZE]; for (int ibl = 0; ibl < nbl; ++ibl) { @@ -11572,25 +11576,33 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy } } float best_score = 0, scale = max; - int besti1 = 0, besti2 = 0; + int besti1 = -1, besti2 = -1, best_shift = 0; for (int i1 = 0; i1 <= IQ1S_BLOCK_SIZE; ++i1) { for (int i2 = i1; i2 <= IQ1S_BLOCK_SIZE; ++i2) { - float sumqx = -(sumx[i1] - sumx[0]) + (sumx[IQ1S_BLOCK_SIZE] - sumx[i2]); - float sumq2 = (sumw[i1] - sumw[0]) + (sumw[IQ1S_BLOCK_SIZE] - sumw[i2]); + float sumqx = (sumx[i1] - sumx[0])*x_p[0] + (sumx[i2] - sumx[i1])*x_p[1] + (sumx[IQ1S_BLOCK_SIZE] - sumx[i2])*x_p[2]; + float sumq2 = (sumw[i1] - sumw[0])*x_p[0]*x_p[0] + (sumw[i2] - sumw[i1])*x_p[1]*x_p[1] + (sumw[IQ1S_BLOCK_SIZE] - sumw[i2])*x_p[2]*x_p[2]; if (sumq2 > 0 && sumqx*sumqx > best_score*sumq2) { scale = sumqx/sumq2; best_score = scale*sumqx; - besti1 = i1; besti2 = i2; + besti1 = i1; besti2 = i2; best_shift = 1; + } + sumqx = (sumx[i1] - sumx[0])*x_m[0] + (sumx[i2] - sumx[i1])*x_m[1] + (sumx[IQ1S_BLOCK_SIZE] - sumx[i2])*x_m[2]; + sumq2 = (sumw[i1] - sumw[0])*x_m[0]*x_m[0] + (sumw[i2] - sumw[i1])*x_m[1]*x_m[1] + (sumw[IQ1S_BLOCK_SIZE] - sumw[i2])*x_m[2]*x_m[2]; + if (sumq2 > 0 && sumqx*sumqx > best_score*sumq2) { + scale = sumqx/sumq2; best_score = scale*sumqx; + besti1 = i1; besti2 = i2; best_shift = -1; } } } + GGML_ASSERT(besti1 >= 0 && besti2 >= 0 && best_shift != 0); for (int j = 0; j < besti1; ++j) L[idx[2*j]] = 0; for (int j = besti1; j < besti2; ++j) L[idx[2*j]] = 1; for (int j = besti2; j < IQ1S_BLOCK_SIZE; ++j) L[idx[2*j]] = 2; if (scale < 0) { for (int j = 0; j < IQ1S_BLOCK_SIZE; ++j) L[j] = 2 - L[j]; - scale = -scale; + scale = -scale; best_shift = -best_shift; } bool all_on_grid = true; + const float * xx = best_shift == 1 ? x_p : x_m; for (int k = 0; k < IQ1S_BLOCK_SIZE/8; ++k) { uint16_t u = 0; for (int j = 0; j < 8; ++j) u |= (L[8*k+j] << 2*j); @@ -11598,7 +11610,7 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy if (grid_index < 0) { all_on_grid = false; const uint16_t * neighbours = kneighbors_q2xs - kmap_q2xs[u] - 1; - grid_index = iq1_find_best_neighbour2(neighbours, kgrid_q2xs, xb + 8*k, weight + 8*k, scale, L + 8*k, NGRID_IQ1S); + grid_index = iq1_find_best_neighbour2(neighbours, kgrid_q2xs, xb + 8*k, weight + 8*k, scale, xx, L + 8*k, NGRID_IQ1S); GGML_ASSERT(grid_index >= 0); } index[k] = grid_index; @@ -11609,7 +11621,7 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy const int8_t * pg = (const int8_t *)(kgrid_q2xs + index[k]); for (int j = 0; j < 8; ++j) { float w = weight[8*k + j]; - float q = (pg[j] - 3)/2; + float q = xx[(pg[j] - 1)/2]; sumqx += w*q*xb[8*k+j]; sumq2 += w*q*q; } @@ -11624,6 +11636,7 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy y[ibl].qh[ib] = h; GGML_ASSERT(scale >= 0); scales[ib] = scale; + shifts[ib] = best_shift; max_scale = MAX(max_scale, scale); } @@ -11632,12 +11645,13 @@ static void quantize_row_iq1_s_impl(const float * restrict x, void * restrict vy continue; } - float d = max_scale/31; + float d = max_scale/15; y[ibl].d = GGML_FP32_TO_FP16(d*1.125f); // 1.085f is another fudge factor. Don't ask me why it is needed. float id = 1/d; for (int ib = 0; ib < QK_K/IQ1S_BLOCK_SIZE; ++ib) { int l = nearest_int(0.5f*(id*scales[ib]-1)); - l = MAX(0, MIN(15, l)); + l = MAX(0, MIN(7, l)); + if (shifts[ib] == -1) l |= 8; y[ibl].qh[ib] |= (l << 12); } }