ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-02-25 23:54:10 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
2c9aaae809c4b6219152e57db5659ef9558246d5
ik_llama.cpp/ggml/src/iqk/iqk_quantize.cpp
Iwan Kawrakow 2c9aaae809 Fix Makefile 
I always use cmake, so had forgotten to pay attention to the
Makefile.
2024-08-09 17:29:32 +03:00

1985 lines
78 KiB
C++
Raw Blame History

//
// Copyright (C) 2024 Iwan Kawrakow
// MIT license
// SPDX-License-Identifier: MIT
//
#if GGML_USE_IQK_MULMAT
#include "iqk_mul_mat.h"
#endif
#include "ggml-quants.h"
#include "ggml-impl.h"
#define GGML_COMMON_IMPL_C
#include "ggml-common.h"
#include "iqk_quantize.h"
#include <vector>
#include <utility>
#include <cstdint>
#include <cmath>
#include <array>
#include <algorithm>
#include <cstring>
namespace {
inline int nearest_int(float fval) {
assert(fval <= 4194303.f);
float val = fval + 12582912.f;
int i; memcpy(&i, &val, sizeof(int));
return (i & 0x007fffff) - 0x00400000;
}
struct IQ1BNQuantizer {
int8_t L[QK_IQ1BN];
void quantize_one_row_1bn(const float * src, block_iq1_bn * y, int n_per_row, const float * imatrix);
void quantize_one_row_2bn(const float * src, block_iq2_bn * y, int n_per_row, const float * imatrix);
static inline float row_max(int n_per_row, const float * src) {
float max_in_row = 0;
for (int j = 0; j < n_per_row; ++j) {
float ax = fabsf(src[j]);
max_in_row = std::max(max_in_row, ax);
}
return max_in_row;
}
// The Makefile has issues dwaling with this?
//static constexpr uint8_t k_mult[5] = {81, 27, 9, 3, 1};
static const uint8_t k_mult[5];
};
const uint8_t IQ1BNQuantizer::k_mult[5] = {81, 27, 9, 3, 1};
void IQ1BNQuantizer::quantize_one_row_1bn(const float * src, block_iq1_bn * y, int n_per_row, const float * imatrix) {
static const int k_nb[6] = {1, 3, 9, 27, 81, 243};
(void)imatrix;
const int nblock = n_per_row/QK_IQ1BN;
for (int ib = 0; ib < nblock; ++ib) {
std::memset(&y[ib], 0, sizeof(block_iq1_bn));
auto xb = src + ib*QK_IQ1BN;
int v13 = 0;
for (int i16 = 0; i16 < QK_IQ1BN/16; ++i16) {
for (int k = 0; k < 3; ++k) {
int idx = 0;
for (int j = 0; j < 5; ++j) {
float v = xb[16*i16 + 5*k + j];
int q = fabsf(v) < 1e-6f ? 1 : v < 0 ? 0 : 2;
idx += k_nb[j]*q;
}
idx = (256*idx + k_nb[5] - 1)/k_nb[5];
y[ib].ql[3*i16 + k] = idx;
}
float v = xb[16*i16 + 15];
int q = fabsf(v) < 1e-6f ? 1 : v < 0 ? 0 : 2;
v13 += k_nb[i16]*q;
}
y[ib].extra = (256*v13 + k_nb[5] - 1)/k_nb[5];
}
}
void IQ1BNQuantizer::quantize_one_row_2bn(const float * src, block_iq2_bn * y, int n_per_row, const float * imatrix) {
(void)imatrix;
const int nblock = n_per_row/QK_IQ1BN;
constexpr int Nj = QK_IQ1BN/4;
for (int ib = 0; ib < nblock; ++ib) {
auto xb = src + QK_IQ1BN*ib;
for (int j = 0; j < QK_IQ1BN; ++j) {
L[j] = fabsf(xb[j]) < 1e-6f ? 1 : xb[j] < 0 ? 0 : 2;
}
for (int j = 0; j < Nj; ++j) {
y[ib].qs[j] = L[j] | (L[j + Nj] << 2) | (L[j + 2*Nj] << 4) | (L[j + 3*Nj] << 6);
}
}
}
}
size_t quantize_iq1_bn(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
IQ1BNQuantizer iq1bn;
int nblock = n_per_row/QK_IQ1BN;
block_iq1_bn * y = (block_iq1_bn *)dst;
for (int row = 0; row < nrows; ++row) {
iq1bn.quantize_one_row_1bn(src + row*n_per_row, y, n_per_row, imatrix);
y += nblock;
}
return sizeof(block_iq1_bn)*nblock*nrows;
}
void quantize_row_iq1_bn_ref(const float * x, block_iq1_bn * y, int64_t k) {
quantize_iq1_bn(x, y, 1, k, nullptr);
}
void quantize_row_iq1_bn(const float * x, void * y, int64_t k) {
quantize_iq1_bn(x, y, 1, k, nullptr);
}
void dequantize_row_iq1_bn(const block_iq1_bn * x, float * y, int64_t k) {
assert(k%QK_IQ1BN == 0);
int nblock = k / QK_IQ1BN;
for (int i = 0; i < nblock; ++i) {
uint8_t extra = x[i].extra;
auto ql = x[i].ql;
for (int i16 = 0; i16 < QK_IQ1BN/16; ++i16) {
for (int k = 0; k < 3; ++k) {
for (int j = 0; j < 5; ++j) {
uint8_t v = ql[k]*IQ1BNQuantizer::k_mult[j];
int8_t vs = ((v + (v >> 1)) >> 7);
*y++ = vs - 1;
}
}
ql += 3;
uint8_t v = extra*IQ1BNQuantizer::k_mult[i16];
int8_t vs = ((v + (v >> 1)) >> 7);
*y++ = vs - 1;
}
}
}
size_t quantize_iq2_bn(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
IQ1BNQuantizer iq1bn;
int nblock = n_per_row/QK_IQ1BN;
block_iq2_bn * y = (block_iq2_bn *)dst;
for (int row = 0; row < nrows; ++row) {
iq1bn.quantize_one_row_2bn(src + row*n_per_row, y, n_per_row, imatrix);
y += nblock;
}
return sizeof(block_iq2_bn)*nblock*nrows;
}
void quantize_row_iq2_bn_ref(const float * x, block_iq2_bn * y, int64_t k) {
quantize_iq2_bn(x, y, 1, k, nullptr);
}
void quantize_row_iq2_bn(const float * x, void * y, int64_t k) {
quantize_iq2_bn(x, y, 1, k, nullptr);
}
void dequantize_row_iq2_bn(const block_iq2_bn * x, float * y, int64_t k) {
assert(k%QK_IQ1BN == 0);
int nblock = k / QK_IQ1BN;
auto d1 = 1.f, d2 = 0.25f, d3 = d2*0.25f, d4 = d3*0.25f;
auto m = -1.f;
constexpr int Nj = QK_IQ1BN/4;
for (int i = 0; i < nblock; ++i) {
for (int j = 0; j < Nj; ++j) {
y[j+ 0] = d1*(x[i].qs[j] & 0x03) + m;
y[j+1*Nj] = d2*(x[i].qs[j] & 0x0c) + m;
y[j+2*Nj] = d3*(x[i].qs[j] & 0x30) + m;
y[j+3*Nj] = d4*(x[i].qs[j] & 0xc0) + m;
}
y += QK_IQ1BN;
}
}
namespace {
inline int8_t iq1bn_dequant(uint8_t q, int i) {
uint8_t v = IQ1BNQuantizer::k_mult[i]*q;
//int8_t vs = (v + (v << 1)) >> 8;
int8_t vs = 3*v >> 8;
return vs - 1;
}
}
static const int8_t iq1bn_values[1280] = {
-1, -1, -1, -1, -1, 0, 0, 0, 0, 0, 0, -1, -1, -1, -1, 1, -1, -1, -1, -1, -1, 0, -1, -1, -1, 0, 0, -1, -1, -1, 1, 0,
-1, -1, -1, -1, 1, -1, -1, -1, 0, 1, -1, -1, -1, 1, 1, -1, -1, -1, -1, -1, 0, -1, -1, 0, -1, 0, -1, -1, 1, -1, 0, -1,
-1, -1, 0, 0, -1, -1, 0, 0, 0, -1, -1, 1, 0, 0, -1, -1, -1, 1, 0, -1, -1, 0, 1, 0, -1, -1, 1, 1, 0, -1, -1, -1,
-1, 1, -1, -1, 0, 0, 0, 0, 0, 0, -1, 1, -1, -1, 1, -1, 1, -1, -1, -1, 0, 1, -1, -1, 0, 0, 1, -1, -1, 1, 0, 1,
-1, -1, -1, 1, 1, -1, -1, 0, 1, 1, -1, -1, 1, 1, 1, -1, -1, -1, -1, -1, 0, -1, 0, -1, -1, 0, -1, 1, -1, -1, 0, -1,
-1, 0, -1, 0, -1, 0, 0, -1, 0, -1, 1, 0, -1, 0, -1, -1, 1, -1, 0, -1, 0, 1, -1, 0, -1, 1, 1, -1, 0, -1, -1, -1,
0, 0, -1, 0, -1, 0, 0, -1, 0, 0, 0, 0, 0, 1, -1, 0, 0, -1, -1, 0, 0, 0, -1, 0, 0, 0, 0, -1, 1, 0, 0, 0,
-1, -1, 1, 0, 0, -1, 0, 1, 0, 0, -1, 1, 1, 0, 0, -1, -1, -1, 1, 0, -1, 0, -1, 1, 0, -1, 1, -1, 1, 0, -1, -1,
0, 1, 0, -1, 0, 0, 1, 0, -1, 1, 0, 1, 0, -1, -1, 1, 1, 0, -1, 0, 1, 1, 0, -1, 1, 1, 1, 0, -1, -1, -1, -1,
1, -1, 0, -1, -1, 1, -1, 1, -1, -1, 1, -1, 0, 0, 0, 0, 0, -1, 0, -1, 1, -1, 0, 0, -1, 1, -1, 1, 0, -1, 1, -1,
-1, 1, -1, 1, -1, 0, 1, -1, 1, -1, 1, 1, -1, 1, -1, -1, -1, 0, 1, -1, 0, -1, 0, 1, -1, 1, -1, 0, 1, -1, -1, 0,
0, 1, -1, 0, 0, 0, 1, -1, 1, 0, 0, 1, -1, -1, 1, 0, 1, -1, 0, 1, 0, 1, -1, 1, 1, 0, 1, -1, -1, -1, 1, 1,
-1, 0, -1, 1, 1, -1, 1, -1, 1, 1, -1, 0, 0, 0, 0, 0, -1, 0, 1, 1, -1, 0, 0, 1, 1, -1, 1, 0, 1, 1, -1, -1,
1, 1, 1, -1, 0, 1, 1, 1, -1, 1, 1, 1, 1, -1, -1, -1, -1, -1, 0, 0, -1, -1, -1, 0, 1, -1, -1, -1, 0, -1, 0, -1,
-1, 0, 0, 0, -1, -1, 0, 1, 0, -1, -1, 0, -1, 1, -1, -1, 0, 0, 1, -1, -1, 0, 1, 1, -1, -1, 0, -1, -1, 0, -1, 0,
0, -1, 0, -1, 0, 1, -1, 0, -1, 0, -1, 0, 0, -1, 0, 0, 0, 0, 0, 0, 0, 0, 0, -1, 0, 1, 0, 0, -1, 0, -1, 1,
0, -1, 0, 0, 1, 0, -1, 0, 1, 1, 0, -1, 0, -1, -1, 1, -1, 0, 0, -1, 1, -1, 0, 1, -1, 1, -1, 0, -1, 0, 1, -1,
0, 0, 0, 1, -1, 0, 1, 0, 1, -1, 0, -1, 1, 1, -1, 0, 0, 1, 1, -1, 0, 1, 1, 1, -1, 0, -1, -1, -1, 0, 0, 0,
-1, -1, 0, 0, 1, -1, -1, 0, 0, -1, 0, -1, 0, 0, 0, 0, -1, 0, 0, 0, 0, 0, 0, 0, 1, 0, -1, 0, 0, -1, 1, -1,
0, 0, 0, 1, -1, 0, 0, 1, 1, -1, 0, 0, -1, -1, 0, 0, 0, 0, -1, 0, 0, 0, 1, -1, 0, 0, 0, -1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 1, 0, 0, 0, 0, -1, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, -1, -1, 1, 0, 0, 0, -1,
1, 0, 0, 1, -1, 1, 0, 0, -1, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, -1, 1, 1, 0,
0, 0, 1, 1, 0, 0, 1, 1, 1, 0, 0, -1, -1, -1, 1, 0, 0, -1, -1, 1, 0, 1, -1, -1, 1, 0, -1, 0, -1, 1, 0, 0,
0, -1, 1, 0, 1, 0, -1, 1, 0, -1, 1, -1, 1, 0, 0, 1, -1, 1, 0, 1, 1, -1, 1, 0, -1, -1, 0, 1, 0, 0, -1, 0,
1, 0, 1, -1, 0, 1, 0, -1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, -1, 1, 0, 1, 0,
0, 1, 0, 1, 0, 1, 1, 0, 1, 0, -1, -1, 1, 1, 0, 0, -1, 1, 1, 0, 1, -1, 1, 1, 0, -1, 0, 1, 1, 0, 0, 0,
1, 1, 0, 1, 0, 1, 1, 0, -1, 1, 1, 1, 0, 0, 1, 1, 1, 0, 1, 1, 1, 1, 0, -1, -1, -1, -1, 1, 0, -1, -1, -1,
1, 1, -1, -1, -1, 1, -1, 0, -1, -1, 1, 0, 0, -1, -1, 1, 1, 0, -1, -1, 1, -1, 1, -1, -1, 1, 0, 0, 0, 0, 0, 0,
1, -1, -1, 1, 1, 1, -1, -1, 1, -1, -1, 0, -1, 1, 0, -1, 0, -1, 1, 1, -1, 0, -1, 1, -1, 0, 0, -1, 1, 0, 0, 0,
-1, 1, 1, 0, 0, -1, 1, -1, 1, 0, -1, 1, 0, 1, 0, -1, 1, 1, 1, 0, -1, 1, -1, -1, 1, -1, 1, 0, -1, 1, -1, 1,
1, -1, 1, -1, 1, -1, 0, 1, -1, 1, 0, 0, 1, -1, 1, 1, 0, 1, -1, 1, -1, 1, 1, -1, 1, 0, 0, 0, 0, 0, 0, 1,
1, -1, 1, 1, 1, 1, -1, 1, -1, -1, -1, 0, 1, 0, -1, -1, 0, 1, 1, -1, -1, 0, 1, -1, 0, -1, 0, 1, 0, 0, -1, 0,
1, 1, 0, -1, 0, 1, -1, 1, -1, 0, 1, 0, 1, -1, 0, 1, 1, 1, -1, 0, 1, -1, -1, 0, 0, 1, 0, -1, 0, 0, 1, 1,
-1, 0, 0, 1, -1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 1, 0, 0, 0, 1, -1, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0,
0, 0, 1, 1, 0, 0, 1, -1, -1, 1, 0, 1, 0, -1, 1, 0, 1, 1, -1, 1, 0, 1, -1, 0, 1, 0, 1, 0, 0, 1, 0, 1,
1, 0, 1, 0, 1, -1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 1, 1, 0, 1, -1, -1, -1, 1, 1, 0, -1, -1, 1, 1, 1, -1,
-1, 1, 1, -1, 0, -1, 1, 1, 0, 0, -1, 1, 1, 1, 0, -1, 1, 1, -1, 1, -1, 1, 1, 0, 1, -1, 1, 1, 1, 1, -1, 1,
1, 0, 0, 0, 0, 0, -1, -1, 0, 1, 1, 0, -1, 0, 1, 1, 1, -1, 0, 1, 1, -1, 0, 0, 1, 1, 0, 0, 0, 1, 1, 1,
0, 0, 1, 1, -1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 1, 0, 1, 1, -1, -1, 1, 1, 1, 0, -1, 1, 1, 1, 1, -1, 1,
1, 1, -1, 0, 1, 1, 1, 0, 0, 1, 1, 1, 1, 0, 1, 1, 1, -1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1,
};
void ggml_vec_dot_iq1_bn_q8_K64(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
GGML_UNUSED(bs);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(nrc);
static_assert(QK_IQ1BN == 64, "This dot product implementation for iq1_bn requires a block size of 64");
#if GGML_USE_IQK_MULMAT
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ1_BN, vx, 0, GGML_TYPE_Q8_K64, vy, 0, s, 0, 0, 1)) {
return;
}
#endif
const block_iq1_bn * x = (const block_iq1_bn *)vx;
const float * d8 = (const float *)vy;
const int8_t * q8 = (const int8_t *)(d8 + 4);
int nblock = n / QK_IQ1BN;
int sumi[8] = {};
int8_t q1[16];
for (int ii = 0; ii < nblock; ii += 32) {
int16_t sum16[8] = {};
int nb = std::min(ii + 32, nblock);
for (int i = ii; i < nb; ++i) {
auto ql = x[i].ql;
const int8_t * extra = iq1bn_values + 5*x[i].extra;
for (int i16 = 0; i16 < QK_IQ1BN/16; ++i16) {
for (int k = 0; k < 3; ++k) {
uint8_t q = *ql++;
const int8_t * vs = iq1bn_values + 5*q;
for (int j = 0; j < 5; ++j) q1[5*k+j] = vs[j];
}
q1[15] = extra[i16];
// We collect 8 q8 values per block into each element of sum16
// => 32 x 8 = 256 values in each loop over i, so this cannot overflow the int16_t range
// (q8 is in -127...127, and hence the sum is in -32512...32512
for (int j = 0; j < 8; ++j) sum16[j] += q8[2*j+0]*q1[2*j+0] + q8[2*j+1]*q1[2*j+1];
q8 += 16;
}
}
for (int j = 0; j < 8; ++j) sumi[j] += sum16[j];
}
*s = d8[0] * (sumi[0] + sumi[1]) + d8[1] * (sumi[2] + sumi[3]) + d8[2] * (sumi[4] + sumi[5]) + d8[3] * (sumi[6] + sumi[7]);
}
void ggml_vec_dot_iq2_bn_q8_K64(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
GGML_ASSERT(nrc == 1);
GGML_UNUSED(bs);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(nrc);
static_assert(QK_IQ1BN == 64, "This dot product implementation for iq2_bn requires a block size of 64");
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ2_BN, vx, 0, GGML_TYPE_Q8_K64, vy, 0, s, 0, 0, 1)) {
return;
}
constexpr int Nj = QK_IQ1BN/4;
const block_iq2_bn * x = (const block_iq2_bn *)vx;
int nblock = n / QK_IQ1BN;
const float * d = (const float *)vy;
const int8_t * q8 = (const int8_t *)(d + 4);
int sum[16] = { };
int sum0[4] = { };
for (int i = 0; i < nblock; ++i) {
for (int j = 0; j < Nj/4; ++j) {
for (int l = 0; l < 4; ++l) {
sum[4*j + 0] += q8[4*j + l + 0] * (x[i].qs[4*j+l] & 0x03);
sum[4*j + 1] += q8[4*j + l + 1*Nj] * (x[i].qs[4*j+l] & 0x0c);
sum[4*j + 2] += q8[4*j + l + 2*Nj] * (x[i].qs[4*j+l] & 0x30);
sum[4*j + 3] += q8[4*j + l + 3*Nj] * (x[i].qs[4*j+l] & 0xc0);
sum0[j] += q8[4*j + l] + q8[4*j + l + 1*Nj] + q8[4*j + l + 2*Nj] + q8[4*j + l + 3*Nj];
}
}
q8 += QK_IQ1BN;
}
float sumf = 0;
for (int j = 0; j < 4; ++j) {
sumf += d[j] * (sum[4*j + 0] + 0.25f*sum[4*j + 1] + 0.0625*sum[4*j + 2] + 0.015625*sum[4*j + 3] - sum0[j]);
}
*s = sumf;
}
void quantize_row_q8_K64_ref(const float * x, block_q8_K64 * y, int64_t k) {
float * dptr = (float *)y;
auto qs = (int8_t *)(dptr + 4);
#ifdef __ARM_NEON
static const uint8_t k_shuffle[16] = {0, 4, 8, 12, 16, 20, 24, 28, 32, 36, 40, 44, 48, 52, 56, 60};
auto shuffle = vld1q_u8(k_shuffle);
float32x4_t max[4] = { };
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
auto val = vld1q_f32(x + j + 4*i);
val = vabsq_f32(val);
max[i] = vmaxq_f32(max[i], val);
}
}
float32x4_t vid[4];
for (int i = 0; i < 4; ++i) {
dptr[i] = vmaxvq_f32(max[i])/127;
float id = dptr[i] > 0 ? 1/dptr[i] : 0.f;
vid[i] = vdupq_n_f32(id);
}
int8x16x4_t q;
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
auto val = vld1q_f32(x + j + 4*i);
val = vmulq_f32(vid[i], val);
q.val[i] = vreinterpretq_s8_s32(vcvtnq_s32_f32(val));
}
auto qi = vqtbl4q_s8(q, shuffle);
vst1q_s8(qs, qi);
qs += 16;
}
#elif defined __AVX__
__m128 max[4] = {};
__m128 sign_bit = _mm_set1_ps(-0.f);
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
auto val = _mm_loadu_ps(x + j + 4*i);
val = _mm_andnot_ps(sign_bit, val);
max[i] = _mm_max_ps(max[i], val);
}
}
__m128 vid[4];
for (int i = 0; i < 4; ++i) {
max[i] = _mm_max_ps(max[i], _mm_movehl_ps(max[i], max[i]));
max[i] = _mm_max_ss(max[i], _mm_movehdup_ps(max[i]));
float maxi = _mm_cvtss_f32(max[i]);
dptr[i] = maxi/127;
float id = dptr[i] > 0 ? 1/dptr[i] : 0.f;
vid[i] = _mm_set1_ps(id);
}
__m128i q[4];
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
auto val = _mm_loadu_ps(x + j + 4*i);
val = _mm_round_ps(_mm_mul_ps(vid[i], val), _MM_ROUND_NEAREST);
q[i] = _mm_cvtps_epi32(val);
}
auto q1 = _mm_packs_epi32(q[0], q[1]);
auto q2 = _mm_packs_epi32(q[2], q[3]);
auto qi = _mm_packs_epi16(q1, q2);
_mm_storeu_si128((__m128i *)qs, qi);
qs += 16;
}
#else
float aux[4] = {0.f, 0.f, 0.f, 0.f};
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
for (int l = 0; l < 4; ++l) {
float ax = fabsf(x[j+4*i+l]);
aux[i] = std::max(aux[i], ax);
}
}
}
for (int i = 0; i < 4; ++i) {
dptr[i] = aux[i]/127;
aux[i] = dptr[i] > 0 ? 1/dptr[i] : 0.f;
}
for (int j = 0; j < k; j += 16) {
for (int i = 0; i < 4; ++i) {
for (int l = 0; l < 4; ++l) qs[j+4*i+l] = nearest_int(aux[i]*x[j+4*i+l]);
}
}
#endif
}
void quantize_row_q8_K64(const float * x, void * y, int64_t k) {
quantize_row_q8_K64_ref(x, (block_q8_K64 *)y, k);
}
//
// ============================================== iq2_K
//
namespace {
inline int best_index_iq2nl(const int8_t * values, float x) {
int idx = x < values[1] ? 0 : x > values[2] ? 2 : 1;
return x - values[idx] < values[idx+1] - x ? idx : idx + 1;
}
void quantize_row_iq2_k_impl(const float * x, void * vy, int n_per_row, const float * quant_weights) {
constexpr int kBlockSize = 16;
block_iq2_k * y = (block_iq2_k *)vy;
float scales[QK_K/kBlockSize];
float weight[kBlockSize];
float sumx[kBlockSize+1], sumw[kBlockSize+1];
std::array<std::pair<float,int>, kBlockSize> pairs;
const int8_t * shifted_values = iq2nl_values + 4;
for (int ibl = 0; ibl < n_per_row/QK_K; ++ibl) {
memset(&y[ibl], 0, sizeof(block_iq2_k));
y[ibl].d = GGML_FP32_TO_FP16(0.f);
const float * xbl = x + ibl*QK_K;
float sumx2 = 0;
for (int j = 0; j < QK_K; ++j) sumx2 += xbl[j]*xbl[j];
const float sigma2 = 1.5f*sumx2/QK_K;
uint16_t extra = 0;
float max_abs_scale = 0;
for (int ib = 0; ib < QK_K/kBlockSize; ++ib) {
const float * xb = xbl + kBlockSize*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*kBlockSize;
for (int j = 0; j < kBlockSize; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < kBlockSize; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
for (int j = 0; j < kBlockSize; ++j) pairs[j] = {xb[j], j};
std::sort(pairs.begin(), pairs.end());
sumx[0] = sumw[0] = 0;
for (int j = 0; j < kBlockSize; ++j) {
int jj = pairs[j].second;
sumw[j+1] = sumw[j] + weight[jj];
sumx[j+1] = sumx[j] + weight[jj]*xb[jj];
}
float best = 0, d = 0;
bool is_shifted = false;
float sumqx, sumq2;
for (int i1 = 0; i1 < kBlockSize; ++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];
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];
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];
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];
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
d = sumqx/sumq2; best = d*sumqx; is_shifted = true;
}
}
}
}
scales[ib] = d;
if (is_shifted) extra |= (1 << ib);
float abs_scale = fabsf(scales[ib]);
max_abs_scale = MAX(max_abs_scale, abs_scale);
}
if (!max_abs_scale) continue;
float d = max_abs_scale/15;
y[ibl].d = GGML_FP32_TO_FP16(d);
y[ibl].extra = extra;
float id = 1/d;
float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < QK_K/kBlockSize; ++ib) {
int ls = nearest_int(0.5f*(id*scales[ib]+15));
ls = MAX(0, MIN(15, ls));
y[ibl].scales[ib/2] |= (ls << 4*(ib%2));
ls = 2*ls - 15;
float dl = d * ls;
if (dl) {
const int8_t * block_values = y[ibl].extra & (1 << ib) ? shifted_values : iq2nl_values;
const float * xb = xbl + kBlockSize*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*kBlockSize;
for (int j = 0; j < kBlockSize; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < kBlockSize; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float idl = 1/dl;
int ib32 = ib/2;
int offset = 16*(ib%2);
uint8_t * qs = y[ibl].qs + 32*(ib32/4) + offset;
for (int j = 0; j < 16; ++j) {
const float al = idl*xb[j];
int ibest = best_index_iq2nl(block_values, al);
qs[j] |= (ibest << 2*(ib32%4));
float w = weight[j];
float q = block_values[ibest]*ls;
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
}
}
if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(sumqx/sumq2);
}
}
}
void quantize_row_iq2_k_ref(const float * GGML_RESTRICT x, block_iq2_k * GGML_RESTRICT y, int64_t k) {
assert(k % QK_K == 0);
quantize_iq2_k(x, (void *)y, 1, k, nullptr);
}
void quantize_row_iq2_k(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) {
assert(k % QK_K == 0);
block_iq2_k * y = (block_iq2_k *)vy;
quantize_row_iq2_k_ref(x, y, k);
}
size_t quantize_iq2_k(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
GGML_ASSERT(n_per_row%QK_K == 0);
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
for (int64_t row = 0; row < nrows; ++row) {
quantize_row_iq2_k_impl(src, (void *)qrow, n_per_row, imatrix);
src += n_per_row;
qrow += nblock*sizeof(block_iq2_k);
}
return nrows * nblock * sizeof(block_iq2_k);
}
void dequantize_row_iq2_k(const block_iq2_k * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
assert(k % QK_K == 0);
const int nb = k / QK_K;
for (int i = 0; i < nb; i++) {
const float d = GGML_FP16_TO_FP32(x[i].d);
const uint8_t * qs = x[i].qs;
uint16_t extra = x[i].extra;
int shift = 0;
for (int ib32 = 0; ib32 < QK_K/32; ++ib32) {
float dl1 = d * (2*(x[i].scales[ib32] & 0xf) - 15);
float dl2 = d * (2*(x[i].scales[ib32] >> 4) - 15);
const int8_t * values1 = extra & 1 ? iq2nl_values + 4 : iq2nl_values;
const int8_t * values2 = extra & 2 ? iq2nl_values + 4 : iq2nl_values;
extra >>= 2;
for (int j = 0; j < 16; ++j) {
y[j+ 0] = dl1 * values1[(qs[j+ 0] >> shift) & 3];
y[j+16] = dl2 * values2[(qs[j+16] >> shift) & 3];
}
y += 32;
shift += 2;
if (shift == 8) { qs += 32; shift = 0; }
}
}
}
void vec_dot_iq2_k_q8_k(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
assert(n % QK_K == 0);
assert(nrc == 1);
GGML_UNUSED(nrc);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(bs);
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ2_K, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
const int nb = n / QK_K;
const block_iq2_k * x = (const block_iq2_k *)vx;
const block_q8_K * y = (const block_q8_K *)vy;
}
//
// ============================================== iq3_k
//
namespace {
const int8_t iq3nl_index[111] = {
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 8, 8, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 9,
9, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 10, 10, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 11, 11, 4, 4, 4, 4,
4, 4, 4, 4, 4, 4, 12, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 13, 13, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6,
6, 6, 6, 6, 14, 14, 7, 7, 7, 7, 7, 7, 7, 7, 7
};
inline int best_index_iq3nl(const int8_t * values, float x) {
int ix = (int)x - values[0];
if (ix < 0 || ix >= 111) return ix < 0 ? 0 : 7;
ix = iq3nl_index[ix];
return ix < 8 ? ix : x - values[ix-8] < values[ix-7] - x ? ix-8 : ix-7;
}
static void quantize_row_iq3_k_impl(const float * x, void * vy, int n_per_row, const float * quant_weights) {
const int ntry = 5;
block_iq3_k * y = (block_iq3_k *)vy;
float scales[QK_K/16];
float weight[16];
const int8_t * shifted_values = iq3nl_values + 8;
for (int ibl = 0; ibl < n_per_row/QK_K; ++ibl) {
memset(&y[ibl], 0, sizeof(block_iq3_k));
y[ibl].d = GGML_FP32_TO_FP16(0.f);
const float * xbl = x + ibl*QK_K;
float sumx2 = 0;
for (int j = 0; j < QK_K; ++j) sumx2 += xbl[j]*xbl[j];
const float sigma2 = sumx2/QK_K;
uint16_t extra = 0;
float max_abs_scale = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float amax = 0, max = 0;
for (int j = 0; j < 16; ++j) {
float ax = fabsf(xb[j]);
if (ax > amax) {
amax = ax; max = xb[j];
}
}
if (!amax) {
scales[ib] = 0;
continue;
}
float d = ntry > 0 ? -max/iq3nl_values[0] : max/iq3nl_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 < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq3nl(iq3nl_values, al);
float q = iq3nl_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq3nl(iq3nl_values, -al);
q = iq3nl_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 = -ntry; itry <= ntry; ++itry) {
id = (itry + iq3nl_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq3nl(iq3nl_values, al);
float q = iq3nl_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq3nl(iq3nl_values, -al);
q = iq3nl_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 = (itry + shifted_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq3nl(shifted_values, al);
float q = shifted_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq3nl(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;
}
}
if (d) {
const int8_t * block_values = is_shifted ? shifted_values : iq3nl_values;
float sumqx = 0, sumq2 = 0;
id = 1/d;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq3nl(block_values, al);
float q = block_values[l];
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
if (sumq2 > 0) d = sumqx/sumq2;
}
scales[ib] = d;
if (is_shifted) extra |= (1 << ib);
float abs_scale = fabsf(scales[ib]);
max_abs_scale = MAX(max_abs_scale, abs_scale);
}
if (!max_abs_scale) continue;
float d = max_abs_scale/31;
y[ibl].d = GGML_FP32_TO_FP16(d);
y[ibl].extra = extra;
float id = 1/d;
float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
int ls = nearest_int(0.5f*(id*fabsf(scales[ib])-1));
ls = MAX(0, MIN(15, ls));
y[ibl].scales_l[ib/2] |= (ls << 4*(ib%2));
if (scales[ib] < 0) y[ibl].scales_h |= (1 << ib);
ls = (2*ls + 1) * (scales[ib] < 0 ? -1 : 1);
float dl = d * ls;
if (dl) {
const int8_t * block_values = y[ibl].extra & (1 << ib) ? shifted_values : iq3nl_values;
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float idl = 1/dl;
int ib32 = ib/2;
int offset = 16*(ib%2);
uint8_t * qs = y[ibl].qs + 32*(ib32/4) + offset;
uint8_t * qh = y[ibl].qh + 32*(ib32/8) + offset;
for (int j = 0; j < 16; ++j) {
const float al = idl*xb[j];
int ibest = best_index_iq3nl(block_values, al);
qs[j] |= ((ibest & 3) << 2*(ib32%4));
qh[j] |= ((ibest >> 2) << (ib32%8));
float w = weight[j];
float q = block_values[ibest]*ls;
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
}
}
if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(sumqx/sumq2);
}
}
}
void quantize_row_iq3_k_ref(const float * x, block_iq3_k * y, int64_t k) {
assert(k % QK_K == 0);
quantize_iq3_k(x, (void *)y, 1, k, nullptr);
}
void quantize_row_iq3_k(const float * x, void * vy, int64_t k) {
assert(k % QK_K == 0);
block_iq3_k * y = (block_iq3_k *)vy;
quantize_row_iq3_k_ref(x, y, k);
}
size_t quantize_iq3_k(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
GGML_ASSERT(n_per_row%QK_K == 0);
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
for (int64_t row = 0; row < nrows; ++row) {
quantize_row_iq3_k_impl(src, (void *)qrow, n_per_row, imatrix);
src += n_per_row;
qrow += nblock*sizeof(block_iq3_k);
}
return nrows * nblock * sizeof(block_iq3_k);
}
void dequantize_row_iq3_k(const block_iq3_k * x, float * y, int64_t k) {
assert(k % QK_K == 0);
const int nb = k / QK_K;
for (int i = 0; i < nb; i++) {
const float d = GGML_FP16_TO_FP32(x[i].d);
const uint8_t * qs = x[i].qs;
const uint8_t * qh = x[i].qh;
uint16_t sh = x[i].scales_h;
uint16_t extra = x[i].extra;
for (int ib32 = 0; ib32 < QK_K/32; ++ib32) {
float dl1 = d * ((2*(x[i].scales_l[ib32] & 0xf) + 1) * ((sh & 1) ? -1 : 1));
float dl2 = d * ((2*(x[i].scales_l[ib32] >> 4) + 1) * ((sh & 2) ? -1 : 1));
sh >>= 2;
const int8_t * values1 = extra & 1 ? iq3nl_values + 8 : iq3nl_values;
const int8_t * values2 = extra & 2 ? iq3nl_values + 8 : iq3nl_values;
extra >>= 2;
int shift_l = 2*(ib32%4);
int shift_h = ib32%8;
for (int j = 0; j < 16; ++j) {
y[j+ 0] = dl1 * values1[((qs[j+ 0] >> shift_l) & 3) | (((qh[j+ 0] >> shift_h) & 1) << 2)];
y[j+16] = dl2 * values2[((qs[j+16] >> shift_l) & 3) | (((qh[j+16] >> shift_h) & 1) << 2)];
}
y += 32;
if (shift_l == 6) qs += 32;
}
}
}
void vec_dot_iq3_k_q8_k(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
assert(n % QK_K == 0);
assert(nrc == 1);
GGML_UNUSED(nrc);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(bs);
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ3_K, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
const int nb = n / QK_K;
const block_iq2_k * x = (const block_iq2_k *)vx;
const block_q8_K * y = (const block_q8_K *)vy;
}
//
// ============================================== iq4_K
//
void dequantize_row_iq4_k(const block_iq4_k * x, float * y, int64_t k) {
assert(k % QK_K == 0);
const int nb = k / QK_K;
for (int i = 0; i < nb; i++) {
const uint8_t * qs = x[i].qs;
const float d = GGML_FP16_TO_FP32(x[i].d);
uint16_t extra = x[i].extra;
for (int ib = 0; ib < QK_K/32; ++ib) {
const uint8_t sh = x[i].scales_h[ib/2] >> 4*(ib%2);
const float dl1 = d * (((x[i].scales_l[ib] & 0xf) | ((sh << 4) & 0x30)) - 32);
const float dl2 = d * (((x[i].scales_l[ib] >> 4) | ((sh << 2) & 0x30)) - 32);
const int8_t * values1 = extra & 1 ? iq4k_values + 16 : iq4k_values;
const int8_t * values2 = extra & 2 ? iq4k_values + 16 : iq4k_values;
extra >>= 2;
for (int j = 0; j < 16; ++j) {
y[j+ 0] = dl1 * values1[qs[j] & 0xf];
y[j+16] = dl2 * values2[qs[j] >> 4];
}
y += 32;
qs += 16;
}
}
}
void vec_dot_iq4_k_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
assert(n % QK_K == 0);
assert(nrc == 1);
GGML_UNUSED(nrc);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(bs);
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ4_K, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
const int nb = n / QK_K;
const block_iq4_k * x = (const block_iq4_k *)vx;
const block_q8_K * y = (const block_q8_K *)vy;
float sumf = 0;
for (int ibl = 0; ibl < nb; ++ibl) {
const float d4d8 = GGML_FP16_TO_FP32(x[ibl].d) * y[ibl].d;
uint16_t extra = x[ibl].extra;
uint32_t h = *((const uint32_t *)x[ibl].scales_h);
const uint8_t * qs = x[ibl].qs;
const int8_t * q8 = y[ibl].qs;
int32_t sum = 0;
for (int ib = 0; ib < QK_K/32; ++ib) {
const int ls1 = ((x[ibl].scales_l[ib] & 0xf) | ((h << 4) & 0x30)) - 32;
const int ls2 = ((x[ibl].scales_l[ib] >> 4) | ((h << 2) & 0x30)) - 32;
h >>= 4;
const int8_t * values1 = iq4k_values + 16*(extra & 1);
const int8_t * values2 = iq4k_values + 8*(extra & 2);
extra >>= 2;
int sumi1 = 0, sumi2 = 0;
for (int j = 0; j < 16; ++j) {
sumi1 += q8[j+ 0] * values1[qs[j] & 0xf];
sumi2 += q8[j+16] * values2[qs[j] >> 4];
}
sum += ls1*sumi1 + ls2*sumi2;
qs += 16;
q8 += 32;
}
sumf += d4d8 * sum;
}
*s = sumf;
}
namespace {
const int8_t iq4nl_index[241] = {
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 16, 16, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 17, 17, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 18, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3,
3, 3, 3, 3, 3, 3, 19, 4, 4, 4, 4, 4, 4, 4, 4, 4, 4, 4, 4, 4, 4, 20, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5,
5, 5, 21, 21, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 22, 7, 7, 7, 7, 7, 7, 7, 7, 7, 7, 23, 23, 8, 8, 8, 8,
8, 8, 8, 8, 8, 8, 24, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 25, 10, 10, 10, 10, 10, 10, 10, 10, 10, 10, 10, 26, 26,
11, 11, 11, 11, 11, 11, 11, 11, 11, 11, 11, 11, 27, 27, 12, 12, 12, 12, 12, 12, 12, 12, 12, 12, 12, 12, 12, 12, 28, 13, 13, 13,
13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 29, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14, 14,
14, 14, 14, 14, 30, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15
};
inline int best_index_iq4nl(const int8_t * values, float x) {
int ix = (int)x - values[0];
if (ix < 0 || ix >= 241) return ix < 0 ? 0 : 15;
ix = iq4nl_index[ix];
return ix < 16 ? ix : x - values[ix-16] < values[ix-15] - x ? ix-16 : ix-15;
}
static void quantize_row_iq4_k_impl_bs16(const int super_block_size, const int block_size, const float * x,
block_iq4_k * y,
float * scales, float * weight, uint8_t * L,
const int8_t * values,
const float * quant_weights,
const int ntry) {
GGML_ASSERT(super_block_size == 256 && block_size == 16);
float sigma2 = 0;
for (int j = 0; j < super_block_size; ++j) sigma2 += x[j]*x[j];
sigma2 *= 2.f/super_block_size;
memset(y, 0, sizeof(block_iq4_k));
y->d = GGML_FP32_TO_FP16(0.f);
uint16_t * scales_h = (uint16_t *)y->scales_h;
const int8_t * shifted_values = values + 16;
float max_scale = 0, amax_scale = 0;
uint16_t extra = 0;
for (int ib = 0; ib < super_block_size/block_size; ++ib) {
const float * xb = x + ib*block_size;
if (quant_weights) {
const float * qw = quant_weights + ib*block_size;
for (int j = 0; j < block_size; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < block_size; ++j) weight[j] = xb[j]*xb[j];
}
float amax = 0, max = 0;
for (int j = 0; j < block_size; ++j) {
float ax = fabsf(xb[j]);
if (ax > amax) {
amax = ax; max = xb[j];
}
}
if (!amax) {
scales[ib] = 0;
continue;
}
float d = ntry > 0 ? -max/values[0] : max/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 < block_size; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq4nl(values, al);
float q = values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq4nl(values, -al);
q = values[l];
sumqx_m += w*q*xb[j];
sumq2_m += w*q*q;
}
d = sumqx_p/sumq2_p;
bool is_shifted = false;
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;
}
for (int itry = -ntry; itry <= ntry; ++itry) {
id = (itry + values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < block_size; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq4nl(values, al);
float q = values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq4nl(values, -al);
q = 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 = (itry + shifted_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < block_size; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq4nl(shifted_values, al);
float q = shifted_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq4nl(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;
}
}
if (is_shifted) extra |= (1 << ib);
scales[ib] = d;
float abs_d = fabsf(d);
if (abs_d > amax_scale) {
amax_scale = abs_d; max_scale = d;
}
}
float d = -max_scale/32;
y->d = GGML_FP32_TO_FP16(d);
y->extra = extra;
float id = d ? 1/d : 0.f;
float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < super_block_size/block_size; ++ib) {
const int8_t * block_values = extra & (1 << ib) ? shifted_values : values;
int l = nearest_int(id*scales[ib]);
l = MAX(-32, MIN(31, l));
float dl = d * l;
float idl = dl ? 1/dl : 0.f;
uint8_t * Lb = L + ib*block_size;
const float * xb = x + ib*block_size;
if (quant_weights) {
const float * qw = quant_weights + ib*block_size;
for (int j = 0; j < block_size; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < block_size; ++j) weight[j] = xb[j]*xb[j];
}
for (int j = 0; j < block_size; ++j) {
Lb[j] = best_index_iq4nl(block_values, idl*xb[j]);
float w = weight[j];
float q = block_values[Lb[j]]*l;
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
l += 32;
uint8_t l_l = l & 0xf;
uint8_t l_h = l >> 4;
if (ib%2 == 0) y->scales_l[ib/2] = l_l;
else y->scales_l[ib/2] |= (l_l << 4);
scales_h[ib/8] |= (l_h << 2*(ib%8));
}
if (sumq2 > 0) y->d = GGML_FP32_TO_FP16(sumqx/sumq2);
for (int i = 0; i < super_block_size/32; ++i) {
for (int j = 0; j < 16; ++j) {
y->qs[16*i + j] = L[32*i + j] | (L[32*i + 16 + j] << 4);
}
}
}
}
void quantize_row_iq4_k_ref(const float * x, block_iq4_k * y, int64_t k) {
assert(k % QK_K == 0);
quantize_iq4_k(x, (void *)y, 1, k, nullptr);
}
void quantize_row_iq4_k(const float * x, void * vy, int64_t k) {
assert(k % QK_K == 0);
block_iq4_k * y = (block_iq4_k *)vy;
quantize_row_iq4_k_ref(x, y, k);
}
size_t quantize_iq4_k(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
GGML_ASSERT(n_per_row%QK_K == 0);
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
uint8_t L[QK_K];
float weight[16];
float scales[QK_K/16];
for (int64_t row = 0; row < nrows; ++row) {
block_iq4_k * iq4 = (block_iq4_k *)qrow;
for (int ibl = 0; ibl < nblock; ++ibl) {
const float * qw = imatrix ? imatrix + QK_K*ibl : NULL;
quantize_row_iq4_k_impl_bs16(QK_K, 16, src + QK_K*ibl, iq4 + ibl,
scales, weight, L, iq4k_values, qw, 7);
}
src += n_per_row;
qrow += nblock*sizeof(block_iq4_k);
}
return nrows * nblock * sizeof(block_iq4_k);
}
//
// ============================================== iq5_K
//
void dequantize_row_iq5_k(const block_iq5_k * x, float * y, int64_t k) {
assert(k % QK_K == 0);
const int nb = k / QK_K;
for (int i = 0; i < nb; i++) {
const float d = GGML_FP16_TO_FP32(x[i].d);
const uint8_t * qs = x[i].qs;
const uint8_t * qh = x[i].qh;
const uint8_t * sl = x[i].scales_l;
const uint8_t * sh = x[i].scales_h;
uint16_t extra = x[i].extra;
int shift = 0;
for (int ib64 = 0; ib64 < QK_K/64; ++ib64) {
float dl1 = d * (((sl[2*ib64+0] & 0xf) | ((sh[ib64] << 4) & 0x30)) - 32);
float dl2 = d * (((sl[2*ib64+0] >> 4) | ((sh[ib64] << 2) & 0x30)) - 32);
float dl3 = d * (((sl[2*ib64+1] & 0xf) | ((sh[ib64] >> 0) & 0x30)) - 32);
float dl4 = d * (((sl[2*ib64+1] >> 4) | ((sh[ib64] >> 2) & 0x30)) - 32);
const int8_t * values1 = iq5nl_values + ((extra & 1) << 5);
const int8_t * values2 = iq5nl_values + ((extra & 2) << 4);
const int8_t * values3 = iq5nl_values + ((extra & 4) << 3);
const int8_t * values4 = iq5nl_values + ((extra & 8) << 2);
for (int j = 0; j < 16; ++j) {
y[j+ 0] = dl1 * values1[(qs[j+ 0] & 0xf) | (((qh[j+ 0] >> shift) & 1) << 4)];
y[j+16] = dl2 * values2[(qs[j+16] & 0xf) | (((qh[j+16] >> shift) & 1) << 4)];
y[j+32] = dl3 * values3[(qs[j+ 0] >> 4) | (((qh[j+ 0] >> shift) & 2) << 3)];
y[j+48] = dl4 * values4[(qs[j+16] >> 4) | (((qh[j+16] >> shift) & 2) << 3)];
}
y += 64;
qs += 32;
extra >>= 4;
shift += 2;
if (shift == 8) { qh += 32; shift = 0; }
}
}
}
void vec_dot_iq5_k_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
assert(n % QK_K == 0);
assert(nrc == 1);
GGML_UNUSED(nrc);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(bs);
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ5_K, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
const int nb = n / QK_K;
const block_iq5_k * x = (const block_iq5_k *)vx;
const block_q8_K * y = (const block_q8_K *)vy;
float sumf = 0;
for (int i = 0; i < nb; i++) {
const float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
const uint8_t * qs = x[i].qs;
const uint8_t * qh = x[i].qh;
const uint8_t * sl = x[i].scales_l;
const uint8_t * sh = x[i].scales_h;
const int8_t * q8 = y[i].qs;
uint16_t extra = x[i].extra;
int shift = 0;
int sumb = 0;
for (int ib64 = 0; ib64 < QK_K/64; ++ib64) {
int dl1 = (((sl[2*ib64+0] & 0xf) | ((sh[ib64] << 4) & 0x30)) - 32);
int dl2 = (((sl[2*ib64+0] >> 4) | ((sh[ib64] << 2) & 0x30)) - 32);
int dl3 = (((sl[2*ib64+1] & 0xf) | ((sh[ib64] >> 0) & 0x30)) - 32);
int dl4 = (((sl[2*ib64+1] >> 4) | ((sh[ib64] >> 2) & 0x30)) - 32);
const int8_t * values1 = iq5nl_values + ((extra & 1) << 5);
const int8_t * values2 = iq5nl_values + ((extra & 2) << 4);
const int8_t * values3 = iq5nl_values + ((extra & 4) << 3);
const int8_t * values4 = iq5nl_values + ((extra & 8) << 2);
int sumi1 = 0, sumi2 = 0, sumi3 = 0, sumi4 = 0;
for (int j = 0; j < 16; ++j) {
sumi1 += q8[j+ 0] * values1[(qs[j+ 0] & 0xf) | (((qh[j+ 0] >> shift) & 1) << 4)];
sumi2 += q8[j+16] * values2[(qs[j+16] & 0xf) | (((qh[j+16] >> shift) & 1) << 4)];
sumi3 += q8[j+32] * values3[(qs[j+ 0] >> 4) | (((qh[j+ 0] >> shift) & 2) << 3)];
sumi4 += q8[j+48] * values4[(qs[j+16] >> 4) | (((qh[j+16] >> shift) & 2) << 3)];
}
sumb += dl1 * sumi1 + dl2 * sumi2 + dl3 * sumi3 + dl4 * sumi4;
q8 += 64;
qs += 32;
extra >>= 4;
shift += 2;
}
sumf += d * sumb;
}
*s = sumf;
}
namespace {
const int8_t iq5nl_index[248] = {
0, 0, 0, 0, 0, 0, 32, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 33, 33, 2, 2, 2, 2, 2, 2, 2, 2, 2, 34, 34, 3, 3,
3, 3, 3, 3, 3, 3, 35, 35, 4, 4, 4, 4, 4, 4, 4, 36, 36, 5, 5, 5, 5, 5, 5, 5, 37, 37, 6, 6, 6, 6, 6, 6,
6, 38, 7, 7, 7, 7, 7, 7, 39, 39, 8, 8, 8, 8, 8, 40, 40, 9, 9, 9, 9, 9, 41, 41, 10, 10, 10, 10, 10, 42, 11, 11,
11, 11, 11, 43, 12, 12, 12, 12, 12, 44, 13, 13, 13, 13, 13, 45, 14, 14, 14, 14, 14, 46, 15, 15, 15, 15, 47, 47, 16, 16, 16, 16,
48, 17, 17, 17, 17, 17, 49, 18, 18, 18, 18, 18, 50, 19, 19, 19, 19, 19, 51, 20, 20, 20, 20, 20, 52, 21, 21, 21, 21, 21, 53, 53,
22, 22, 22, 22, 22, 54, 54, 23, 23, 23, 23, 23, 23, 55, 24, 24, 24, 24, 24, 24, 24, 56, 25, 25, 25, 25, 25, 25, 25, 57, 57, 26,
26, 26, 26, 26, 26, 26, 58, 58, 27, 27, 27, 27, 27, 27, 27, 27, 59, 28, 28, 28, 28, 28, 28, 28, 28, 28, 60, 29, 29, 29, 29, 29,
29, 29, 29, 29, 29, 61, 30, 30, 30, 30, 30, 30, 30, 30, 30, 30, 30, 62, 31, 31, 31, 31, 31, 31
};
inline int best_index_iq5nl(const int8_t * values, float x) {
int ix = (int)x - values[0];
if (ix < 0 || ix >= 247) return ix < 0 ? 0 : 31;
ix = iq5nl_index[ix];
return ix < 32 ? ix : x - values[ix-32] < values[ix-31] - x ? ix-32 : ix-31;
}
void quantize_row_iq5_k_impl(const float * x, void * vy, int n_per_row, const float * quant_weights) {
const int ntry = 5;
const float step = 1.f;
block_iq5_k * y = (block_iq5_k *)vy;
float scales[QK_K/16];
float weight[16];
const int8_t * shifted_values = iq5nl_values + 32;
for (int ibl = 0; ibl < n_per_row/QK_K; ++ibl) {
memset(&y[ibl], 0, sizeof(block_iq5_k));
y[ibl].d = GGML_FP32_TO_FP16(0.f);
const float * xbl = x + ibl*QK_K;
float sumx2 = 0;
for (int j = 0; j < QK_K; ++j) sumx2 += xbl[j]*xbl[j];
const float sigma2 = 2*sumx2/QK_K;
float max_scale = 0, max_abs_scale = 0;
uint16_t extra = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float amax = 0, max = 0;
for (int j = 0; j < 16; ++j) {
float ax = fabsf(xb[j]);
if (ax > amax) {
amax = ax; max = xb[j];
}
}
if (!amax) {
scales[ib] = 0;
continue;
}
float d = ntry > 0 ? -max/iq5nl_values[0] : max/iq5nl_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 < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq5nl(iq5nl_values, al);
float q = iq5nl_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq5nl(iq5nl_values, -al);
q = iq5nl_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 = -ntry; itry <= ntry; ++itry) {
id = (itry*step + iq5nl_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq5nl(iq5nl_values, al);
float q = iq5nl_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq5nl(iq5nl_values, -al);
q = iq5nl_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 = (itry*step + shifted_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq5nl(shifted_values, al);
float q = shifted_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
l = best_index_iq5nl(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;
}
}
if (d) {
const int8_t * block_values = is_shifted ? shifted_values : iq5nl_values;
float sumqx = 0, sumq2 = 0;
id = 1/d;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
int l = best_index_iq5nl(block_values, al);
float q = block_values[l];
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
if (sumq2 > 0) d = sumqx/sumq2;
}
scales[ib] = d;
if (is_shifted) extra |= (1 << ib);
float abs_scale = fabsf(scales[ib]);
if (abs_scale > max_abs_scale) {
max_abs_scale = abs_scale; max_scale = scales[ib];
}
}
if (!max_abs_scale) continue;
float d = -max_scale/32;
y[ibl].d = GGML_FP32_TO_FP16(d);
y[ibl].extra = extra;
float id = 1/d;
float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
int ls = nearest_int(id*scales[ib]);
ls = MAX(-32, MIN(31, ls));
int uls = ls + 32;
y[ibl].scales_l[ib/2] |= ((uls & 0xf) << 4*(ib%2));
y[ibl].scales_h[ib/4] |= ((uls >> 4) << 2*(ib%4));
float dl = d * ls;
if (dl) {
const int8_t * block_values = y[ibl].extra & (1 << ib) ? shifted_values : iq5nl_values;
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float idl = 1/dl;
int ib32 = ib/2;
int offset = 16*(ib%2);
uint8_t * qs = y[ibl].qs + 32*(ib32/2) + offset;
uint8_t * qh = y[ibl].qh + 32*(ib32/8) + offset;
for (int j = 0; j < 16; ++j) {
const float al = idl*xb[j];
int ibest = best_index_iq5nl(block_values, al);
qs[j] |= ((ibest & 0xf) << 4*(ib32%2));
qh[j] |= ((ibest >> 4) << (ib32%8));
float w = weight[j];
float q = block_values[ibest]*ls;
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
}
}
if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(sumqx/sumq2);
}
}
}
void quantize_row_iq5_k_ref(const float * x, block_iq5_k * y, int64_t k) {
assert(k % QK_K == 0);
quantize_iq5_k(x, (void *)y, 1, k, nullptr);
}
void quantize_row_iq5_k(const float * x, void * vy, int64_t k) {
assert(k % QK_K == 0);
block_iq5_k * y = (block_iq5_k *)vy;
quantize_row_iq5_k_ref(x, y, k);
}
size_t quantize_iq5_k(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
GGML_ASSERT(n_per_row%QK_K == 0);
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
for (int64_t row = 0; row < nrows; ++row) {
quantize_row_iq5_k_impl(src, (void *)qrow, n_per_row, imatrix);
src += n_per_row;
qrow += nblock*sizeof(block_iq5_k);
}
return nrows * nblock * sizeof(block_iq5_k);
}
//
// ============================================== iq6_K
//
#define A_IQ6K -127.f
#define B_IQ6K 6.2568f
#define C_IQ6K 0.11218f
#define D_IQ6K 0.0011972f
#define S_IQ6K 1.f
void dequantize_row_iq6_k(const block_iq6_k * x, float * y, int64_t k) {
assert(k % QK_K == 0);
const int nb = k / QK_K;
for (int i = 0; i < nb; i++) {
const float d = GGML_FP16_TO_FP32(x[i].d);
const uint8_t * qs = x[i].qs;
const uint8_t * qh = x[i].qh;
const int8_t * sl = x[i].scales;
uint16_t extra = x[i].extra;
int shift = 0;
for (int ib64 = 0; ib64 < QK_K/64; ++ib64) {
float dl1 = d * sl[4*ib64 + 0];
float dl2 = d * sl[4*ib64 + 1];
float dl3 = d * sl[4*ib64 + 2];
float dl4 = d * sl[4*ib64 + 3];
float m1 = extra & 1 ? S_IQ6K : 0;
float m2 = extra & 2 ? S_IQ6K : 0;
float m3 = extra & 4 ? S_IQ6K : 0;
float m4 = extra & 8 ? S_IQ6K : 0;
for (int j = 0; j < 16; ++j) {
float q1 = ((qs[j+ 0] & 0xf) | (((qh[j+ 0] >> shift) & 0x03) << 4));
float q2 = ((qs[j+16] & 0xf) | (((qh[j+16] >> shift) & 0x03) << 4));
float q3 = ((qs[j+ 0] >> 4) | (((qh[j+ 0] >> shift) & 0x0c) << 2));
float q4 = ((qs[j+16] >> 4) | (((qh[j+16] >> shift) & 0x0c) << 2));
y[j+ 0] = dl1 * (A_IQ6K + q1*(B_IQ6K + q1*(-C_IQ6K + q1*D_IQ6K)) + m1);
y[j+16] = dl2 * (A_IQ6K + q2*(B_IQ6K + q2*(-C_IQ6K + q2*D_IQ6K)) + m2);
y[j+32] = dl3 * (A_IQ6K + q3*(B_IQ6K + q3*(-C_IQ6K + q3*D_IQ6K)) + m3);
y[j+48] = dl4 * (A_IQ6K + q4*(B_IQ6K + q4*(-C_IQ6K + q4*D_IQ6K)) + m4);
}
y += 64;
qs += 32;
extra >>= 4;
shift += 4;
if (shift == 8) { qh += 32; shift = 0; }
}
}
}
void vec_dot_iq6_k_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
assert(n % QK_K == 0);
assert(nrc == 1);
GGML_UNUSED(nrc);
GGML_UNUSED(bx);
GGML_UNUSED(by);
GGML_UNUSED(bs);
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ6_K, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
// TODO
//const int nb = n / QK_K;
//const block_iq5_k * x = (const block_iq5_k *)vx;
//const block_q8_K * y = (const block_q8_K *)vy;
//float sumf = 0;
//for (int i = 0; i < nb; i++) {
// const float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
// const uint8_t * qs = x[i].qs;
// const uint8_t * qh = x[i].qh;
// const uint8_t * sl = x[i].scales_l;
// const uint8_t * sh = x[i].scales_h;
// const int8_t * q8 = y[i].qs;
// uint16_t extra = x[i].extra;
// int shift = 0;
// int sumb = 0;
// for (int ib64 = 0; ib64 < QK_K/64; ++ib64) {
// int dl1 = (((sl[2*ib64+0] & 0xf) | ((sh[ib64] << 4) & 0x30)) - 32);
// int dl2 = (((sl[2*ib64+0] >> 4) | ((sh[ib64] << 2) & 0x30)) - 32);
// int dl3 = (((sl[2*ib64+1] & 0xf) | ((sh[ib64] >> 0) & 0x30)) - 32);
// int dl4 = (((sl[2*ib64+1] >> 4) | ((sh[ib64] >> 2) & 0x30)) - 32);
// const int8_t * values1 = iq5nl_values + ((extra & 1) << 5);
// const int8_t * values2 = iq5nl_values + ((extra & 2) << 4);
// const int8_t * values3 = iq5nl_values + ((extra & 4) << 3);
// const int8_t * values4 = iq5nl_values + ((extra & 8) << 2);
// int sumi1 = 0, sumi2 = 0, sumi3 = 0, sumi4 = 0;
// for (int j = 0; j < 16; ++j) {
// sumi1 += q8[j+ 0] * values1[(qs[j+ 0] & 0xf) | (((qh[j+ 0] >> shift) & 1) << 4)];
// sumi2 += q8[j+16] * values2[(qs[j+16] & 0xf) | (((qh[j+16] >> shift) & 1) << 4)];
// sumi3 += q8[j+32] * values3[(qs[j+ 0] >> 4) | (((qh[j+ 0] >> shift) & 2) << 3)];
// sumi4 += q8[j+48] * values4[(qs[j+16] >> 4) | (((qh[j+16] >> shift) & 2) << 3)];
// }
// sumb += dl1 * sumi1 + dl2 * sumi2 + dl3 * sumi3 + dl4 * sumi4;
// q8 += 64;
// qs += 32;
// extra >>= 4;
// shift += 2;
// }
// sumf += d * sumb;
//}
//*s = sumf;
}
namespace {
inline int best_index(int n, const float * val, float x) {
if (x <= val[0]) return 0;
if (x >= val[n-1]) return n-1;
int ml = 0, mu = n-1;
while (mu-ml > 1) {
int mav = (ml+mu)/2;
if (x < val[mav]) mu = mav; else ml = mav;
}
return x - val[mu-1] < val[mu] - x ? mu-1 : mu;
}
uint8_t iq6nl_index[249] = {
0, 0, 0, 64, 1, 1, 1, 1, 1, 65, 2, 2, 2, 2, 2, 66, 3, 3, 3, 3, 67, 67, 4, 4, 4, 4, 68, 5, 5, 5, 5, 69,
69, 6, 6, 6, 70, 70, 7, 7, 7, 71, 8, 8, 8, 72, 72, 9, 9, 9, 73, 73, 10, 10, 10, 74, 11, 11, 11, 75, 12, 12, 12, 76,
13, 13, 13, 77, 14, 14, 14, 78, 15, 15, 79, 79, 16, 16, 80, 17, 17, 81, 81, 18, 18, 82, 19, 19, 83, 83, 20, 84, 84, 21, 85, 85,
22, 86, 86, 23, 87, 87, 24, 88, 88, 25, 89, 89, 26, 90, 90, 27, 91, 91, 28, 92, 29, 93, 93, 30, 94, 94, 31, 95, 95, 32, 96, 33,
97, 97, 34, 98, 98, 35, 99, 99, 36, 100, 100, 37, 101, 38, 102, 102, 39, 103, 103, 40, 104, 104, 41, 41, 105, 42, 42, 106, 106, 43, 107, 107,
44, 108, 108, 45, 45, 109, 46, 46, 46, 110, 47, 47, 111, 111, 48, 48, 112, 49, 49, 49, 113, 50, 50, 50, 114, 51, 51, 51, 115, 52, 52, 52,
116, 116, 53, 53, 53, 117, 54, 54, 54, 118, 118, 55, 55, 55, 119, 119, 56, 56, 56, 120, 120, 57, 57, 57, 121, 121, 58, 58, 58, 58, 122, 59,
59, 59, 59, 123, 123, 60, 60, 60, 60, 124, 61, 61, 61, 61, 61, 125, 62, 62, 62, 62, 62, 126, 63, 63, 63,
};
inline int best_index_iq6nl(const float * values, float x) {
int ix = (int)(x - values[0]);
if (ix < 0 || ix >= 249) return ix < 0 ? 0 : 63;
ix = iq6nl_index[ix];
return ix < 64 ? ix : x - values[ix-64] < values[ix-63] - x ? ix-64 : ix-63;
//if (x <= val[0]) return 0;
//if (x >= val[63]) return 63;
//int index = iq6nl_index[int(x - val[0])];
//return index < 64 ? index : x - val[index-64] < val[index-63] - x ? index - 64 : index - 63;
}
void quantize_row_iq6_k_impl(const float * x, void * vy, int n_per_row, const float * quant_weights, const float * values, const float * shifted_values) {
const int ntry = 5;
const float step = 1.f;
block_iq6_k * y = (block_iq6_k *)vy;
float scales[QK_K/16];
float weight[16];
for (int ibl = 0; ibl < n_per_row/QK_K; ++ibl) {
memset(&y[ibl], 0, sizeof(block_iq6_k));
y[ibl].d = GGML_FP32_TO_FP16(0.f);
const float * xbl = x + ibl*QK_K;
float sumx2 = 0;
for (int j = 0; j < QK_K; ++j) sumx2 += xbl[j]*xbl[j];
const float sigma2 = 2*sumx2/QK_K;
float max_scale = 0, max_abs_scale = 0;
uint16_t extra = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float amax = 0, max = 0;
for (int j = 0; j < 16; ++j) {
float ax = fabsf(xb[j]);
if (ax > amax) {
amax = ax; max = xb[j];
}
}
if (!amax) {
scales[ib] = 0;
continue;
}
float d = ntry > 0 ? -max/values[0] : max/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 < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
//int l = best_index(64, values, al);
int l = best_index_iq6nl(values, al);
float q = values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
//l = best_index(64, values, -al);
l = best_index_iq6nl(values, -al);
q = 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 = -ntry; itry <= ntry; ++itry) {
id = (itry*step + values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
//int l = best_index(64, values, al);
int l = best_index_iq6nl(values, al);
float q = values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
//l = best_index(64, values, -al);
l = best_index_iq6nl(values, -al);
q = 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 = (itry*step + shifted_values[0])/max;
sumqx_p = sumq2_p = 0;
sumqx_m = sumq2_m = 0;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
//int l = best_index(64, shifted_values, al);
int l = best_index_iq6nl(shifted_values, al);
float q = shifted_values[l];
sumqx_p += w*q*xb[j];
sumq2_p += w*q*q;
//l = best_index(64, shifted_values, -al);
l = best_index_iq6nl(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;
}
}
if (d) {
const float * block_values = is_shifted ? shifted_values : values;
float sumqx = 0, sumq2 = 0;
id = 1/d;
for (int j = 0; j < 16; ++j) {
float w = weight[j];
float al = id*xb[j];
//int l = best_index(64, block_values, al);
int l = best_index_iq6nl(block_values, al);
float q = block_values[l];
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
if (sumq2 > 0) d = sumqx/sumq2;
}
scales[ib] = d;
if (is_shifted) extra |= (1 << ib);
float abs_scale = fabsf(scales[ib]);
if (abs_scale > max_abs_scale) {
max_abs_scale = abs_scale; max_scale = scales[ib];
}
}
if (!max_abs_scale) continue;
float d = -max_scale/127;
y[ibl].d = GGML_FP32_TO_FP16(d);
y[ibl].extra = extra;
float id = 1/d;
float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < QK_K/16; ++ib) {
int ls = nearest_int(id*scales[ib]);
ls = MAX(-127, MIN(127, ls));
y[ibl].scales[ib] |= ls;
float dl = d * ls;
if (dl) {
const float * block_values = y[ibl].extra & (1 << ib) ? shifted_values : values;
const float * xb = xbl + 16*ib;
if (quant_weights) {
const float * qw = quant_weights + ibl*QK_K + ib*16;
for (int j = 0; j < 16; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
} else {
for (int j = 0; j < 16; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
}
float idl = 1/dl;
int ib32 = ib/2;
int offset = 16*(ib%2);
uint8_t * qs = y[ibl].qs + 32*(ib32/2) + offset;
uint8_t * qh = y[ibl].qh + 32*(ib32/4) + offset;
for (int j = 0; j < 16; ++j) {
const float al = idl*xb[j];
//int ibest = best_index(64, block_values, al);
int ibest = best_index_iq6nl(block_values, al);
qs[j] |= ((ibest & 0xf) << 4*(ib32%2));
qh[j] |= ((ibest >> 4) << 2*(ib32%4));
float w = weight[j];
float q = block_values[ibest]*ls;
sumqx += w*q*xb[j];
sumq2 += w*q*q;
}
}
}
if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(sumqx/sumq2);
}
}
}
void quantize_row_iq6_k_ref(const float * x, block_iq6_k * y, int64_t k) {
assert(k % QK_K == 0);
quantize_iq6_k(x, (void *)y, 1, k, nullptr);
}
void quantize_row_iq6_k(const float * x, void * vy, int64_t k) {
assert(k % QK_K == 0);
block_iq6_k * y = (block_iq6_k *)vy;
quantize_row_iq6_k_ref(x, y, k);
}
size_t quantize_iq6_k(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) {
GGML_ASSERT(n_per_row%QK_K == 0);
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
float values[128];
for (int i = 0; i < 64; ++i) {
values[i] = iq6nl_values[i];
values[i+64] = values[i] + S_IQ6K;
}
for (int64_t row = 0; row < nrows; ++row) {
quantize_row_iq6_k_impl(src, (void *)qrow, n_per_row, imatrix, values, values + 64);
src += n_per_row;
qrow += nblock*sizeof(block_iq6_k);
}
return nrows * nblock * sizeof(block_iq6_k);
}
//
// ========================== IQ2_TN
//
void quantize_row_iq2_tn_ref(const float * x, block_iq2_tn * y, int64_t k) {
GGML_ASSERT(k%QK_K == 0);
int nb = k/QK_K;
auto quantize = [] (float xmax, float x) {
return x < -0.5f*xmax ? 0 : x < 0.5f*xmax ? 1 : 2;
};
for (int ibl = 0; ibl < nb; ++ibl) {
auto xb = x + QK_K*ibl;
float max = xb[0];
for (int j = 0; j < QK_K; ++j) {
float ax = fabsf(xb[j]);
max = std::max(ax, max);
}
y[ibl].d = GGML_FP32_TO_FP16(max);
auto qs = y[ibl].qs;
for (int l = 0; l < QK_K/128; ++l) {
for (int j = 0; j < 32; ++j) {
qs[j] = quantize(max, xb[j]) | (quantize(max, xb[j+32]) << 2) | (quantize(max, xb[j+64]) << 4) | (quantize(max, xb[j+96]) << 6);
}
xb += 128;
qs += 32;
}
}
}
void quantize_row_iq2_tn(const float * x, void * y, int64_t k) {
quantize_row_iq2_tn_ref(x, (block_iq2_tn *)y, k);
}
size_t quantize_iq2_tn(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * /*imatrix*/) {
auto row_size = ggml_row_size(GGML_TYPE_IQ2_TN, n_per_row);
char * qrow = (char *)dst;
for (int row = 0; row < nrows; ++row) {
quantize_row_iq2_tn_ref(src, (block_iq2_tn *)qrow, n_per_row);
qrow += row_size;
src += n_per_row;
}
return row_size*nrows;
}
void dequantize_row_iq2_tn(const block_iq2_tn * x, float * y, int64_t k) {
GGML_ASSERT(k%QK_K == 0);
int nb = k/QK_K;
for (int ibl = 0; ibl < nb; ++ibl) {
float d = GGML_FP16_TO_FP32(x[ibl].d);
auto qs = x[ibl].qs;
for (int l = 0; l < QK_K/128; ++l) {
for (int j = 0; j < 32; ++j) {
y[j+ 0] = d*((qs[j] >> 0) & 3) - d;
y[j+32] = d*((qs[j] >> 2) & 3) - d;
y[j+64] = d*((qs[j] >> 4) & 3) - d;
y[j+96] = d*((qs[j] >> 6) & 3) - d;
}
y += 128;
qs += 32;
}
}
}
void vec_dot_iq2_tn_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ2_TN, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) {
return;
}
const int nb = n / QK_K;
const block_iq2_tn * x = (const block_iq2_tn *)vx;
const block_q8_K * y = (const block_q8_K *)vy;
float sumf = 0;
for (int i = 0; i < nb; i++) {
float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
auto qs = x[i].qs;
auto q8 = y[i].qs;
int sumi1 = 0, sumi2 = 0, sumi3 = 0,sumi4 = 0;
for (int j = 0; j < QK_K/16; ++j) sumi1 -= y[i].bsums[j];
for (int l = 0; l < QK_K/128; ++l) {
for (int j = 0; j < 32; ++j) {
sumi1 += q8[j+ 0] * (qs[j] & 0x03);
sumi2 += q8[j+32] * (qs[j] & 0x0c);
sumi3 += q8[j+64] * (qs[j] & 0x30);
sumi4 += q8[j+96] * (qs[j] & 0xc0);
}
q8 += 128;
qs += 32;
}
sumf += d * (sumi1 + 0.25f*sumi2 + 0.0625f*sumi3 + 0.015625f*sumi4);
}
*s = sumf;
}
Reference in New Issue View Git Blame Copy Permalink