New IQ2_KT, IQ3_KT and IQ4_KT, V2 (#529)
* New iq4_kt trellis The new trellis generates int8_t values via sum_as_uint8_t[(ka * idx + kb) & 0x3f33f3f3f] - 126. CUDA dequantize works. AVX2 case Ny > 32 works, and we get 273 t/s for L3-8B. PPL is on par or even slightly lower than original QTIP trellis. * Something is not working with the AVX2 dot product * New iq4_kt: CUDA MMVQ * New iq4_kt: CUDA MMQ * For now have only iq4_kt use the new trellis * Fix iq2_kt that got broken along the way * New iq4_kt: AVX2 dot product finally works We get 13.6 t/s vs 8.4 t/s with the f16 trellis and f32 arithmetic. Still somewhat slower than other quants, but no longer pathetic. * New iq4_kt: fix vanilla AVX2 * New iq4_kt: NEON implementation We get very respectable PP-512 = 120 t/s. TG-128 is pathetic at 5.3 t/s, so 20+% slower than the f16 variant. * New iq4_kt: slightly faster NEON * New iq4_kt: slightly faster NEON * New iq4_kt: faster NEON We are now at 9.4 t/s, up from 6.6 t/s for the f16 trellis. * Minor * New iq4_kt trellis: not working Metal implementation * Remove the extra 4 bytes of row meta data that is no longer used * Cleanup * Adding forgottent file * Switching iq2_kt to new trellis - CUDA MMQ * New iq2_kt: CUDA GEMV * New iq2_kt: AVX2 dequantize * New iq2_kt: AVX2 GEMM/GEMV * Adding forgotten file * New iq2_kt: NEON GEMM/GEMV * New iq2_kt: slightly faster NEON GEMM * New iq2_kt: Metal - very slow. It seems Apple Silicon cannot quickly add 4 8-bit ints. Or I don't know how to do it - but I didn't find anything in the Metal Shading Language Specification. So, performance is quite a bit worse than the original trellis. * Add missing break * Trying @louiehelm's multiplier * CPU * iq3_kt: use integer trellis + CUDA dequantize and MMVQ * iq3_kt: MMQ * iq3_kt: AVX2 GEMM * iq3_kt: AVX2 GEMV * The trellis quants now need super-blocks of 256, so we need a check --------- Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
This commit is contained in:
Kawrakow
GitHub
parent
129b58b150
commit
d345a15a84
|
|
@ -578,6 +578,20 @@ struct ggml_cuda_type_traits<GGML_TYPE_IQ2_KT> {
|
|||
static constexpr int qi = QI4_XS;
|
||||
};
|
||||
|
||||
template<>
|
||||
struct ggml_cuda_type_traits<GGML_TYPE_IQ3_KT> {
|
||||
static constexpr int qk = QK_K;
|
||||
static constexpr int qr = QR4_XS;
|
||||
static constexpr int qi = QI4_XS;
|
||||
};
|
||||
|
||||
template<>
|
||||
struct ggml_cuda_type_traits<GGML_TYPE_IQ4_KT> {
|
||||
static constexpr int qk = QK_K;
|
||||
static constexpr int qr = QR4_XS;
|
||||
static constexpr int qi = QI4_XS;
|
||||
};
|
||||
|
||||
template<>
|
||||
struct ggml_cuda_type_traits<GGML_TYPE_IQ3_K> {
|
||||
static constexpr int qk = QK_K;
|
||||
|
|
|
|||
|
|
@ -340,6 +340,12 @@ inline __device__ int nearest_int(float fval) {
|
|||
return (i & 0x007fffff) - 0x00400000;
|
||||
}
|
||||
|
||||
int __device__ __forceinline__ trellis_next_int(uint32_t& val) {
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
val = ka*val;
|
||||
return ggml_cuda_dp4a(val & 0x3f3f3f3f, 0x01010101, -126);
|
||||
}
|
||||
|
||||
float __device__ __forceinline__ trellis_next(uint32_t& val) {
|
||||
constexpr uint32_t ka = 89226354;
|
||||
constexpr uint32_t kb = 64248484;
|
||||
|
|
@ -367,9 +373,9 @@ static __global__ void dequantize_block_iq2_kt(const void * __restrict__ vx, dst
|
|||
dst_t * y = yy + ii*QK_K + 8*ib;
|
||||
const uint16_t * ql = (const uint16_t *)x[i].ql;
|
||||
uint32_t idx = ql[ib] + 4096;
|
||||
const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)] * 31.75f * 1.05f;
|
||||
const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)] * 1.05f;
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
y[j] = dl * trellis_next(idx);
|
||||
y[j] = dl * trellis_next_int(idx);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -388,10 +394,10 @@ static __global__ void dequantize_block_iq3_kt(const void * __restrict__ vx, dst
|
|||
dst_t * y = yy + ii*QK_K + 8*ib;
|
||||
const uint16_t * ql = (const uint16_t *)x[i].ql;
|
||||
uint32_t idx = ql[ib] + 4096;
|
||||
const float dl = scale * ((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf) * 31.75f * 1.01f; //1.015f;
|
||||
const float dl = scale * ((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf) * 1.01f; //1.015f;
|
||||
uint8_t mask = 1 << (ib/4);
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
y[j] = dl * std::abs(trellis_next(idx)) * (x[i].qh[(8*ib+j)%32] & mask ? -1.f : 1.f);
|
||||
y[j] = dl * std::abs(trellis_next_int(idx)) * (x[i].qh[(8*ib+j)%32] & mask ? -1.f : 1.f);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -401,9 +407,8 @@ static __global__ void dequantize_block_iq4_kt(const void * __restrict__ vx, dst
|
|||
int64_t ii = blockIdx.x;
|
||||
int64_t row = (QK_K * ii) / n_per_row;
|
||||
const float * dptr = (const float *)((const char *)vx + row * row_size);
|
||||
float scale = dptr[0] * 31.75f * 1.01f;
|
||||
float row_av = dptr[1];
|
||||
const block_iq4_kt * x = (const block_iq4_kt *)(dptr + 2);
|
||||
float scale = dptr[0] * 1.00f;
|
||||
const block_iq4_kt * x = (const block_iq4_kt *)(dptr + 1);
|
||||
const int64_t i = ii - (row*n_per_row)/QK_K;
|
||||
|
||||
constexpr int kNumGroups = 64;
|
||||
|
|
@ -423,8 +428,8 @@ static __global__ void dequantize_block_iq4_kt(const void * __restrict__ vx, dst
|
|||
int ls = ((shb[ib32] & 0xff) >> 1) - 64;
|
||||
const float dl = scale * ls;
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
y[j+0] = dl * trellis_next(idx1) + row_av;
|
||||
y[j+4] = dl * trellis_next(idx2) + row_av;
|
||||
y[j+0] = dl * trellis_next_int(idx1);
|
||||
y[j+4] = dl * trellis_next_int(idx2);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -433,6 +433,119 @@ __device__ __forceinline__ void vec_dot_iq4_ks_q8_1(
|
|||
*result += dl * __low2float(bq8_1[ib32].ds) * sumi;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void vec_dot_iq4_kt_q8_1(
|
||||
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs, float * result) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
float scale = *(const float *)vbq;
|
||||
const block_iq4_kt * bq4 = (const block_iq4_kt *)((const char *)vbq + sizeof(float)) + kbx;
|
||||
|
||||
// iqs is 0...28
|
||||
const int ib32 = iqs/4; // Why iqs/4 ?
|
||||
const int32_t * q8 = (const int *)bq8_1[ib32].qs;
|
||||
//const int8_t * q8 = bq8_1[ib32].qs;
|
||||
const int ls = (bq4->qs[ib32] & 0xff) >> 1;
|
||||
const float dl = scale * (ls - 64);
|
||||
const uint32_t idx0 = ((bq4->qs[ib32] & 1) << 15) + 4096;
|
||||
auto ql = (const uint8_t *)(bq4->qs + 8);
|
||||
auto qh = ql + 64;
|
||||
ql += 8*ib32;
|
||||
qh += 8*(ib32%4);
|
||||
const int shift1 = 8 - 4*(ib32/4);
|
||||
int sumi = 0;
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
const uint32_t sh = bq4->qs[ib32] >> (8 + 3*j);
|
||||
uint32_t val = ql[j] + ((qh[j] << shift1) & 0xf00) + ((sh & 7) << 12) + idx0;
|
||||
int v4 = 0;
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
//int s = val & km;
|
||||
//sumi += q8[4*j+k] * ggml_cuda_dp4a(s, 0x01010101, -126);
|
||||
v4 |= (ggml_cuda_dp4a(val & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
sumi = ggml_cuda_dp4a(v4, q8[j], sumi);
|
||||
}
|
||||
*result += dl * __low2float(bq8_1[ib32].ds) * sumi;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void vec_dot_iq2_kt_q8_1(
|
||||
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs, float * result) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
float scale = *(const float *)vbq;
|
||||
const block_iq2_kt * bq2 = (const block_iq2_kt *)((const char *)vbq + sizeof(float)) + kbx;
|
||||
|
||||
// iqs is 0...28
|
||||
const int ib32 = iqs/4;
|
||||
const int32_t * q8 = (const int *)bq8_1[ib32].qs;
|
||||
const int ls = iq4k_values[(bq2->scales[ib32%4] >> 4*(ib32/4)) & 0xf];
|
||||
const float dl = scale * ls * 1.05f;
|
||||
auto ql = (const uint16_t *)bq2->ql;
|
||||
int sumi = 0;
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
uint32_t val = ql[4*ib32+j] + 4096;
|
||||
int v4 = 0;
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v4 |= (ggml_cuda_dp4a(val & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
sumi = ggml_cuda_dp4a(v4, q8[2*j+0], sumi);
|
||||
v4 = 0;
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v4 |= (ggml_cuda_dp4a(val & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
sumi = ggml_cuda_dp4a(v4, q8[2*j+1], sumi);
|
||||
}
|
||||
*result += dl * __low2float(bq8_1[ib32].ds) * sumi;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void vec_dot_iq3_kt_q8_1(
|
||||
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs, float * result) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
float scale = *(const float *)vbq;
|
||||
const block_iq3_kt * bq3 = (const block_iq3_kt *)((const char *)vbq + sizeof(float)) + kbx;
|
||||
|
||||
// iqs is 0...28
|
||||
const int ib32 = iqs/4;
|
||||
const int32_t * q8 = (const int *)bq8_1[ib32].qs;
|
||||
const int ls = (bq3->scales[ib32%4] >> 4*(ib32/4)) & 0xf;
|
||||
const float dl = scale * ls * 1.015f;
|
||||
auto ql = (const uint16_t *)bq3->ql;
|
||||
uint32_t mask = 0x01010101 << ib32;
|
||||
const uint32_t * qh = (const uint32_t *)bq3->qh;
|
||||
int sumi = 0;
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
uint32_t val = ql[4*ib32+j] + 4096;
|
||||
int v4 = 0;
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
int8_t q = std::abs(ggml_cuda_dp4a(val & km, 0x01010101, -126));
|
||||
v4 |= q << 8*k;
|
||||
}
|
||||
uint32_t signs = __vcmpne4(qh[2*j+0] & mask, 0);
|
||||
v4 = __vsub4(v4 ^ signs, signs);
|
||||
sumi = ggml_cuda_dp4a(v4, q8[2*j+0], sumi);
|
||||
v4 = 0;
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
int8_t q = std::abs(ggml_cuda_dp4a(val & km, 0x01010101, -126));
|
||||
v4 |= q << 8*k;
|
||||
}
|
||||
signs = __vcmpne4(qh[2*j+1] & mask, 0);
|
||||
v4 = __vsub4(v4 ^ signs, signs);
|
||||
sumi = ggml_cuda_dp4a(v4, q8[2*j+1], sumi);
|
||||
}
|
||||
*result += dl * __low2float(bq8_1[ib32].ds) * sumi;
|
||||
}
|
||||
|
||||
#define VDR_IQ4_KSS_Q8_1_MMVQ 4
|
||||
#define VDR_IQ4_KSS_Q8_1_MMQ 4
|
||||
|
||||
|
|
@ -1217,6 +1330,30 @@ void mul_mat_vec_iq4_ks_q8_1_cuda(
|
|||
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ4_KS, VDR_IQ4_KS_Q8_1_MMVQ, vec_dot_iq4_ks_q8_1>(vx, vy, dst, ids_data, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
}
|
||||
|
||||
void mul_mat_vec_iq4_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, int64_t ids_nb0, cudaStream_t stream) {
|
||||
|
||||
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ4_KT, VDR_IQ4_KS_Q8_1_MMVQ, vec_dot_iq4_kt_q8_1>(vx, vy, dst, ids_data, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
}
|
||||
|
||||
void mul_mat_vec_iq2_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, int64_t ids_nb0, cudaStream_t stream) {
|
||||
|
||||
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ2_KT, VDR_IQ4_KS_Q8_1_MMVQ, vec_dot_iq2_kt_q8_1>(vx, vy, dst, ids_data, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
}
|
||||
|
||||
void mul_mat_vec_iq3_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, int64_t ids_nb0, cudaStream_t stream) {
|
||||
|
||||
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ3_KT, VDR_IQ4_KS_Q8_1_MMVQ, vec_dot_iq3_kt_q8_1>(vx, vy, dst, ids_data, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
}
|
||||
|
||||
void mul_mat_vec_iq4_kss_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
|
|
|
|||
|
|
@ -100,3 +100,18 @@ void mul_mat_vec_iq1_m_r4_q8_1_cuda(
|
|||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, const int64_t ids_nb0, cudaStream_t stream);
|
||||
|
||||
void mul_mat_vec_iq2_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, const int64_t ids_nb0, cudaStream_t stream);
|
||||
|
||||
void mul_mat_vec_iq3_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, const int64_t ids_nb0, cudaStream_t stream);
|
||||
|
||||
void mul_mat_vec_iq4_kt_q8_1_cuda(
|
||||
const void * vx, const void * vy, float * dst, const char * ids_data,
|
||||
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst,
|
||||
const int ne2, const uint64_t nb02, const uint64_t nb12, const uint64_t nb2, const int64_t ids_nb0, cudaStream_t stream);
|
||||
|
|
|
|||
|
|
@ -100,6 +100,15 @@ void ggml_cuda_op_mul_mat_q(
|
|||
case GGML_TYPE_IQ4_KS_R4:
|
||||
mul_mat_q_case<GGML_TYPE_IQ4_KS_R4>(ctx, args, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
mul_mat_q_case<GGML_TYPE_IQ4_KT>(ctx, args, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ2_KT:
|
||||
mul_mat_q_case<GGML_TYPE_IQ2_KT>(ctx, args, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
mul_mat_q_case<GGML_TYPE_IQ3_KT>(ctx, args, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ5_KS:
|
||||
mul_mat_q_case<GGML_TYPE_IQ5_KS>(ctx, args, stream);
|
||||
break;
|
||||
|
|
@ -172,6 +181,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
|
|||
case GGML_TYPE_IQ4_K:
|
||||
case GGML_TYPE_IQ5_K:
|
||||
case GGML_TYPE_IQ6_K:
|
||||
case GGML_TYPE_IQ2_KT:
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
mmq_supported = true;
|
||||
break;
|
||||
default:
|
||||
|
|
|
|||
|
|
@ -93,6 +93,9 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
|
|||
case GGML_TYPE_IQ5_KS:
|
||||
case GGML_TYPE_IQ5_KS_R4:
|
||||
case GGML_TYPE_IQ6_K:
|
||||
case GGML_TYPE_IQ2_KT:
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
return MMQ_Q8_1_DS_LAYOUT_D4;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
|
|
@ -202,6 +205,9 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml
|
|||
case GGML_TYPE_IQ4_K : return MMQ_DP4A_TXS_Q8_0_16;
|
||||
case GGML_TYPE_IQ5_K : return MMQ_DP4A_TXS_Q8_0_16;
|
||||
case GGML_TYPE_IQ6_K : return MMQ_DP4A_TXS_Q8_0_16;
|
||||
case GGML_TYPE_IQ2_KT : return MMQ_DP4A_TXS_Q8_0;
|
||||
case GGML_TYPE_IQ3_KT : return MMQ_DP4A_TXS_Q8_0;
|
||||
case GGML_TYPE_IQ4_KT : return MMQ_DP4A_TXS_Q8_0;
|
||||
default : return tile_x_sizes{0, 0, 0};
|
||||
}
|
||||
}
|
||||
|
|
@ -250,6 +256,9 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) {
|
|||
case GGML_TYPE_IQ4_K : return MMQ_MMA_TILE_X_K_Q3_K;
|
||||
case GGML_TYPE_IQ5_K : return MMQ_MMA_TILE_X_K_Q3_K;
|
||||
case GGML_TYPE_IQ6_K : return MMQ_MMA_TILE_X_K_Q3_K;
|
||||
case GGML_TYPE_IQ2_KT : return MMQ_MMA_TILE_X_K_Q8_0;
|
||||
case GGML_TYPE_IQ3_KT : return MMQ_MMA_TILE_X_K_Q8_0;
|
||||
case GGML_TYPE_IQ4_KT : return MMQ_MMA_TILE_X_K_Q8_0;
|
||||
default : return 0;
|
||||
}
|
||||
}
|
||||
|
|
@ -2790,6 +2799,226 @@ template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinlin
|
|||
|
||||
}
|
||||
|
||||
template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinline__ void load_tiles_iq4_kt(
|
||||
const char * __restrict__ x, int * __restrict__ x_tile, const int & kbx0, const int & i_max, const int & stride) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + WARP_SIZE*2);
|
||||
#else
|
||||
constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ4_XS, mmq_y);
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + txs.qs);
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
|
||||
const int kqsx = threadIdx.x;
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps) {
|
||||
int i = i0 + threadIdx.y;
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const block_iq4_kt * bxi = (const block_iq4_kt *)(x + i*stride + sizeof(float)) + kbx0;
|
||||
|
||||
int ib32 = kqsx/4;
|
||||
int j = kqsx%4;
|
||||
const auto shb = bxi->qs;
|
||||
const auto ql = (const uint8_t *)(shb + 8);
|
||||
const auto qh = ql + 64;
|
||||
const uint32_t sh = shb[ib32] >> (8 + 6*j);
|
||||
uint32_t offset = 4096 + ((shb[ib32] & 1) << 15);
|
||||
uint32_t val1 = offset + ql[8*ib32+2*j+0] + ((qh[8*(ib32%4)+2*j+0] << (8 - 4*(ib32/4))) & 0xf00) + ((sh & 7) << 12);
|
||||
uint32_t val2 = offset + ql[8*ib32+2*j+1] + ((qh[8*(ib32%4)+2*j+1] << (8 - 4*(ib32/4))) & 0xf00) + ((sh & 56) << 9);
|
||||
int2 v = {0, 0};
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val1 *= ka;
|
||||
val2 *= ka;
|
||||
v.x |= (ggml_cuda_dp4a(val1 & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
v.y |= (ggml_cuda_dp4a(val2 & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 1] = v.y;
|
||||
#else
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 1] = v.y;
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps * 4) {
|
||||
int i = i0 + threadIdx.y * 4 + threadIdx.x / (WARP_SIZE/4);
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const float * dptr = (const float *)(x + i*stride);
|
||||
const block_iq4_kt * bxi = (const block_iq4_kt *)(dptr + 1) + kbx0;
|
||||
const int ls = (bxi->qs[threadIdx.x % 8] & 0xff) >> 1;
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + threadIdx.x % 8] = dptr[0] * (ls - 64);
|
||||
#else
|
||||
x_df[i*(WARP_SIZE/4) + i/4 + threadIdx.x % 8] = dptr[0] * (ls - 64);
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
}
|
||||
|
||||
template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinline__ void load_tiles_iq2_kt(
|
||||
const char * __restrict__ x, int * __restrict__ x_tile, const int & kbx0, const int & i_max, const int & stride) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + WARP_SIZE*2);
|
||||
#else
|
||||
constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ4_XS, mmq_y);
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + txs.qs);
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
|
||||
const int kqsx = threadIdx.x;
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps) {
|
||||
int i = i0 + threadIdx.y;
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const block_iq2_kt * bxi = (const block_iq2_kt *)(x + i*stride + sizeof(float)) + kbx0;
|
||||
|
||||
int ib32 = kqsx/4;
|
||||
int j = kqsx%4;
|
||||
const auto ql = (const uint16_t *)bxi->ql;
|
||||
uint32_t val = ql[4*ib32+j] + 4096;
|
||||
int2 v = {0, 0};
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v.x |= (ggml_cuda_dp4a(val & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v.y |= (ggml_cuda_dp4a(val & km, 0x01010101, -126) & 0xff) << 8*k;
|
||||
}
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 1] = v.y;
|
||||
#else
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 1] = v.y;
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps * 4) {
|
||||
int i = i0 + threadIdx.y * 4 + threadIdx.x / (WARP_SIZE/4);
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const float * dptr = (const float *)(x + i*stride);
|
||||
const float d = dptr[0] * 1.05f;
|
||||
const block_iq2_kt * bxi = (const block_iq2_kt *)(dptr + 1) + kbx0;
|
||||
int ib32 = threadIdx.x % 8;
|
||||
const int ls = iq4k_values[(bxi->scales[ib32%4] >> 4*(ib32/4)) & 0xf];
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + threadIdx.x % 8] = d * ls;
|
||||
#else
|
||||
x_df[i*(WARP_SIZE/4) + i/4 + threadIdx.x % 8] = d * ls;
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
}
|
||||
|
||||
template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinline__ void load_tiles_iq3_kt(
|
||||
const char * __restrict__ x, int * __restrict__ x_tile, const int & kbx0, const int & i_max, const int & stride) {
|
||||
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
constexpr uint32_t km = 0x3f3f3f3f;
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + WARP_SIZE*2);
|
||||
#else
|
||||
constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ4_XS, mmq_y);
|
||||
int * x_qs = (int *) x_tile;
|
||||
float * x_df = (float *) (x_qs + txs.qs);
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
|
||||
const int kqsx = threadIdx.x;
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps) {
|
||||
int i = i0 + threadIdx.y;
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const block_iq3_kt * bxi = (const block_iq3_kt *)(x + i*stride + sizeof(float)) + kbx0;
|
||||
|
||||
int ib32 = kqsx/4;
|
||||
int j = kqsx%4;
|
||||
const auto ql = (const uint16_t *)bxi->ql;
|
||||
const auto qh = (const uint32_t *)bxi->qh;
|
||||
uint32_t mask = 0x01010101 << ib32;
|
||||
uint32_t val = ql[4*ib32+j] + 4096;
|
||||
int2 v = {0, 0};
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v.x |= std::abs(ggml_cuda_dp4a(val & km, 0x01010101, -126)) << 8*k;
|
||||
}
|
||||
auto signs = __vcmpne4(qh[2*j+0] & mask, 0);
|
||||
v.x = __vsub4(v.x ^ signs, signs);
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
val *= ka;
|
||||
v.y |= std::abs(ggml_cuda_dp4a(val & km, 0x01010101, -126)) << 8*k;
|
||||
}
|
||||
signs = __vcmpne4(qh[2*j+1] & mask, 0);
|
||||
v.y = __vsub4(v.y ^ signs, signs);
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*ib32 + 2*j + 1] = v.y;
|
||||
#else
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 0] = v.x;
|
||||
x_qs[i*(2*WARP_SIZE + 1) + 8*ib32 + 2*j + 1] = v.y;
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int i0 = 0; i0 < mmq_y; i0 += nwarps * 4) {
|
||||
int i = i0 + threadIdx.y * 4 + threadIdx.x / (WARP_SIZE/4);
|
||||
|
||||
if (need_check) {
|
||||
i = min(i, i_max);
|
||||
}
|
||||
|
||||
const float * dptr = (const float *)(x + i*stride);
|
||||
const float d = dptr[0] * 1.01f;
|
||||
const block_iq3_kt * bxi = (const block_iq3_kt *)(dptr + 1) + kbx0;
|
||||
int ib32 = threadIdx.x % 8;
|
||||
const int ls = (bxi->scales[ib32%4] >> 4*(ib32/4)) & 0xf;
|
||||
|
||||
#ifdef INT8_MMA_AVAILABLE
|
||||
x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + threadIdx.x % 8] = d * ls;
|
||||
#else
|
||||
x_df[i*(WARP_SIZE/4) + i/4 + threadIdx.x % 8] = d * ls;
|
||||
#endif // INT8_MMA_AVAILABLE
|
||||
}
|
||||
}
|
||||
|
||||
template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinline__ void load_tiles_iq5_ks_r4(
|
||||
const char * __restrict__ x, int * __restrict__ x_tile, const int & kbx0, const int & i_max, const int & stride) {
|
||||
|
||||
|
|
@ -3382,6 +3611,27 @@ struct mmq_type_traits<mmq_x, mmq_y, nwarps, need_check, GGML_TYPE_IQ4_KS_R4> {
|
|||
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y, nwarps>;
|
||||
};
|
||||
|
||||
template <int mmq_x, int mmq_y, int nwarps, bool need_check>
|
||||
struct mmq_type_traits<mmq_x, mmq_y, nwarps, need_check, GGML_TYPE_IQ4_KT> {
|
||||
static constexpr load_tiles_mmq_t load_tiles = load_tiles_iq4_kt<mmq_y, nwarps, need_check>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, nwarps, MMQ_Q8_1_DS_LAYOUT_D4>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y, nwarps>;
|
||||
};
|
||||
|
||||
template <int mmq_x, int mmq_y, int nwarps, bool need_check>
|
||||
struct mmq_type_traits<mmq_x, mmq_y, nwarps, need_check, GGML_TYPE_IQ2_KT> {
|
||||
static constexpr load_tiles_mmq_t load_tiles = load_tiles_iq2_kt<mmq_y, nwarps, need_check>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, nwarps, MMQ_Q8_1_DS_LAYOUT_D4>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y, nwarps>;
|
||||
};
|
||||
|
||||
template <int mmq_x, int mmq_y, int nwarps, bool need_check>
|
||||
struct mmq_type_traits<mmq_x, mmq_y, nwarps, need_check, GGML_TYPE_IQ3_KT> {
|
||||
static constexpr load_tiles_mmq_t load_tiles = load_tiles_iq3_kt<mmq_y, nwarps, need_check>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, nwarps, MMQ_Q8_1_DS_LAYOUT_D4>;
|
||||
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y, nwarps>;
|
||||
};
|
||||
|
||||
template <int mmq_x, int mmq_y, int nwarps, bool need_check>
|
||||
struct mmq_type_traits<mmq_x, mmq_y, nwarps, need_check, GGML_TYPE_IQ5_KS> {
|
||||
static constexpr load_tiles_mmq_t load_tiles = load_tiles_iq5_ks<mmq_y, nwarps, need_check>;
|
||||
|
|
@ -3843,6 +4093,9 @@ extern DECL_MMQ_CASE(GGML_TYPE_IQ5_K);
|
|||
extern DECL_MMQ_CASE(GGML_TYPE_IQ5_KS);
|
||||
extern DECL_MMQ_CASE(GGML_TYPE_IQ6_K);
|
||||
extern DECL_MMQ_CASE(GGML_TYPE_IQ1_S_R4);
|
||||
extern DECL_MMQ_CASE(GGML_TYPE_IQ4_KT);
|
||||
extern DECL_MMQ_CASE(GGML_TYPE_IQ2_KT);
|
||||
extern DECL_MMQ_CASE(GGML_TYPE_IQ3_KT);
|
||||
|
||||
// -------------------------------------------------------------------------------------------------------------------------
|
||||
|
||||
|
|
|
|||
|
|
@ -527,6 +527,15 @@ static void ggml_cuda_op_mul_mat_vec_q_impl(ggml_backend_cuda_context & ctx, ggm
|
|||
case GGML_TYPE_IQ4_KSS:
|
||||
mul_mat_vec_iq4_kss_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ids_data, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ2_KT:
|
||||
mul_mat_vec_iq2_kt_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ids_data, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
mul_mat_vec_iq3_kt_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ids_data, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
mul_mat_vec_iq4_kt_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ids_data, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
break;
|
||||
case GGML_TYPE_IQ2_KS:
|
||||
mul_mat_vec_iq2_ks_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ids_data, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, ne2, nb02, nb12, nb2, ids_nb0, stream);
|
||||
break;
|
||||
|
|
@ -687,6 +696,9 @@ bool ggml_cuda_mmvq_type_supported(ggml_type src0_type) {
|
|||
case GGML_TYPE_IQ5_KS_R4:
|
||||
case GGML_TYPE_IQ1_S_R4:
|
||||
case GGML_TYPE_IQ1_M_R4:
|
||||
case GGML_TYPE_IQ2_KT:
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
return true;
|
||||
default:
|
||||
return false;
|
||||
|
|
|
|||
|
|
@ -0,0 +1,5 @@
|
|||
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
|
||||
|
||||
#include "../mmq.cuh"
|
||||
|
||||
DECL_MMQ_CASE(GGML_TYPE_IQ2_KT);
|
||||
|
|
@ -0,0 +1,5 @@
|
|||
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
|
||||
|
||||
#include "../mmq.cuh"
|
||||
|
||||
DECL_MMQ_CASE(GGML_TYPE_IQ3_KT);
|
||||
|
|
@ -0,0 +1,5 @@
|
|||
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
|
||||
|
||||
#include "../mmq.cuh"
|
||||
|
||||
DECL_MMQ_CASE(GGML_TYPE_IQ4_KT);
|
||||
|
|
@ -6596,6 +6596,37 @@ void kernel_mul_mv_iq2_k_f32_impl(
|
|||
}
|
||||
}
|
||||
|
||||
struct Trellis3 {
|
||||
constexpr constant static uint32_t kmask = 0x3f3f3f3f;
|
||||
constexpr constant static uint32_t ka = 89226354;
|
||||
constexpr constant static uint32_t kb = 64248484;
|
||||
constexpr constant static uint32_t ka1 = ka*ka;
|
||||
constexpr constant static uint32_t kb1 = kb*ka+kb;
|
||||
constexpr constant static uint32_t ka2 = ka1*ka;
|
||||
constexpr constant static uint32_t kb2 = kb1*ka+kb;
|
||||
constexpr constant static uint32_t ka3 = ka2*ka;
|
||||
constexpr constant static uint32_t kb3 = kb2*ka+kb;
|
||||
static inline char4 gen4(uint32_t val) {
|
||||
thread uint32_t aux[4] = {(ka*val + kb) & kmask, (ka1*val + kb1) & kmask, (ka2*val + kb2) & kmask, (ka3*val + kb3) & kmask};
|
||||
thread const int8_t * a8 = (thread const int8_t *)aux;
|
||||
char4 result;
|
||||
for (int i = 0; i < 4; ++i) result[i] = -126 + a8[4*i+0] + a8[4*i+1] + a8[4*i+2] + a8[4*i+3];
|
||||
return result;
|
||||
}
|
||||
template <typename T4>
|
||||
static inline void gen8(uint32_t val, thread T4& v1, thread T4& v2) {
|
||||
thread uint32_t aux[4] = {ka*val + kb, ka1*val + kb1, ka2*val + kb2, ka3*val + kb3};
|
||||
uint32_t aux32[2];
|
||||
thread const int8_t * a8 = (thread const int8_t *)aux32;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
aux32[0] = aux[i] & kmask;
|
||||
aux32[1] = (ka3*aux[i] + kb3) & kmask;
|
||||
v1[i] = -126 + a8[0] + a8[1] + a8[2] + a8[3];
|
||||
v2[i] = -126 + a8[4] + a8[5] + a8[6] + a8[7];
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct Trellis {
|
||||
constexpr constant static uint32_t kmask1 = 0x8fff8fff;
|
||||
constexpr constant static uint32_t kmask2 = 0x3b603b60;
|
||||
|
|
@ -6691,7 +6722,7 @@ void kernel_mul_mv_iq2_kt_f32_impl(
|
|||
float drow[N_DST];
|
||||
for (int row = 0; row < N_DST; ++row) {
|
||||
device const float * dptr = (device const float *)(cx + row*row_size);
|
||||
drow[row] = dptr[0] * 31.75f * 1.05f;
|
||||
drow[row] = dptr[0] * 1.05f;
|
||||
}
|
||||
|
||||
device const block_iq2_kt * x = (device const block_iq2_kt *)(cx + sizeof(float));
|
||||
|
|
@ -6706,10 +6737,10 @@ void kernel_mul_mv_iq2_kt_f32_impl(
|
|||
|
||||
const float ls = drow[row] * iq4k_values[(sc[(it/2)%4] >> 4*(it/8)) & 0xf];
|
||||
|
||||
Trellis::gen8(q2[2*it+0]+4096, v1, v2);
|
||||
Trellis3::gen8(q2[2*it+0]+4096, v1, v2);
|
||||
auto sum = v1*y4[0] + v2*y4[1];
|
||||
|
||||
Trellis::gen8(q2[2*it+1]+4096, v1, v2);
|
||||
Trellis3::gen8(q2[2*it+1]+4096, v1, v2);
|
||||
sum += v1*y4[2] + v2*y4[3];
|
||||
|
||||
sum *= ls;
|
||||
|
|
@ -8542,19 +8573,18 @@ template <typename type4x4>
|
|||
void dequantize_iq2_kt(device const block_iq2_kt * x, short il, thread type4x4 & reg) {
|
||||
// il is 0...15 for QK_K = 256
|
||||
int ib32 = il/2;
|
||||
half scale = iq4k_values[((x->scales[ib32%4] >> 4*(ib32/4)) & 0xf)] * 31.75h * 1.05h;
|
||||
half scale = iq4k_values[((x->scales[ib32%4] >> 4*(ib32/4)) & 0xf)] * 1.05h;
|
||||
device const uint16_t * q2 = (device const uint16_t *)x->ql + 4*ib32 + 2*(il%2);
|
||||
|
||||
half4 v1, v2;
|
||||
char4 v1, v2;
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
Trellis::gen8(q2[i]+4096, v1, v2);
|
||||
v1 *= scale; v2 *= scale;
|
||||
Trellis3::gen8(q2[i]+4096, v1, v2);
|
||||
if constexpr (is_same_v<type4x4, half4x4>) {
|
||||
reg[2*i+0] = v1;
|
||||
reg[2*i+1] = v2;
|
||||
reg[2*i+0] = {scale*(half)v1[0], scale*(half)v1[1], scale*(half)v1[2], scale*(half)v1[3]};
|
||||
reg[2*i+1] = {scale*(half)v2[0], scale*(half)v2[1], scale*(half)v2[2], scale*(half)v2[3]};
|
||||
} else {
|
||||
reg[2*i+0] = {(float)v1[0], (float)v1[1], (float)v1[2], (float)v1[3]};
|
||||
reg[2*i+1] = {(float)v2[0], (float)v2[1], (float)v2[2], (float)v2[3]};
|
||||
reg[2*i+0] = {scale*(float)v1[0], scale*(float)v1[1], scale*(float)v1[2], scale*(float)v1[3]};
|
||||
reg[2*i+1] = {scale*(float)v2[0], scale*(float)v2[1], scale*(float)v2[2], scale*(float)v2[3]};
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -8586,20 +8616,20 @@ void dequantize_iq4_kt(device const block_iq4_kt * x, short il, float d, thread
|
|||
device const uint32_t * shb = x->qs;
|
||||
device const uint8_t * ql = (device const uint8_t *)(shb + 8);
|
||||
device const uint8_t * qh = ql + 64;
|
||||
float scale = d * (((shb[ib32] & 0xff) >> 1) - 64);
|
||||
const int ls = (shb[ib32] & 0xff) >> 1;
|
||||
const float scale = d * (ls - 64);
|
||||
const uint32_t offset = 4096 + ((shb[ib32] & 1) << 15);
|
||||
|
||||
const int jj = ib32*8 + 4*(il%2);
|
||||
ql += jj;
|
||||
qh += jj%32;
|
||||
ql += 8*ib32;
|
||||
qh += 8*(ib32%4);
|
||||
|
||||
uint32_t sh = (shb[ib32] >> (8 + 12*(il%2))) << 12;
|
||||
const int shift = 8 - 4*(jj/32);
|
||||
const int shift = 8 - 4*(ib32/4);
|
||||
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
uint32_t idx = ql[i] + ((qh[i] << shift) & 0xf00) + ((sh >> 3*i) & 0x7000) + offset;
|
||||
auto v = (float4)Trellis::gen4(idx);
|
||||
reg[i] = v * scale;
|
||||
auto c4 = Trellis3::gen4(idx);
|
||||
reg[i] = {scale*c4[0], scale*c4[1], scale*c4[2], scale*c4[3]};
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -8931,18 +8961,17 @@ struct DequantizerKT4 {
|
|||
using type4x4 = T4x4;
|
||||
DequantizerKT4(device const char * cx, short il = 0) : il(il) {
|
||||
device const float * dptr = (device const float *)cx;
|
||||
d[0] = dptr[0] * 31.75f * 1.01f;
|
||||
d[1] = dptr[1];
|
||||
x = (device const Block *)(dptr + 2);
|
||||
d = dptr[0] * 1.01f;
|
||||
x = (device const Block *)(dptr + 1);
|
||||
}
|
||||
inline void convert(thread T4x4& t) const {
|
||||
float4x4 tmp;
|
||||
dequantize_iq4_kt(x, il, d[0], tmp);
|
||||
dequantize_iq4_kt(x, il, d, tmp);
|
||||
for (int i = 0; i < 4; ++i) for (int j = 0; j < 4; ++j) t[i][j] = tmp[i][j];
|
||||
}
|
||||
inline void convert(int64_t ind, thread T4x4& t) {
|
||||
float4x4 tmp;
|
||||
dequantize_iq4_kt(x + ind/nl, ind%nl, d[0], tmp);
|
||||
dequantize_iq4_kt(x + ind/nl, ind%nl, d, tmp);
|
||||
for (int i = 0; i < 4; ++i) for (int j = 0; j < 4; ++j) t[i][j] = tmp[i][j];
|
||||
}
|
||||
inline void next() {
|
||||
|
|
@ -8951,7 +8980,7 @@ struct DequantizerKT4 {
|
|||
}
|
||||
device const Block * x;
|
||||
short il;
|
||||
float d[2];
|
||||
float d;
|
||||
};
|
||||
|
||||
template <typename T4x4, typename Block, typename Scale, int nl, void (*dequantize)(half d, device const Block *, short, thread T4x4&), bool may_not_be_aligned = false>
|
||||
|
|
|
|||
|
|
@ -1596,10 +1596,10 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
|
|||
.from_float = quantize_row_iq2_kt,
|
||||
.from_float_ref = (ggml_from_float_t)quantize_row_iq2_kt_ref,
|
||||
.vec_dot = vec_dot_iq2_kt_q8_k,
|
||||
#ifdef __ARM_NEON
|
||||
.vec_dot_type = GGML_TYPE_F16,
|
||||
#if defined __AVX2__
|
||||
.vec_dot_type = GGML_TYPE_Q8_2_X4,
|
||||
#else
|
||||
.vec_dot_type = GGML_TYPE_F32,
|
||||
.vec_dot_type = GGML_TYPE_Q8_0_X4,
|
||||
#endif
|
||||
.nrows = 1,
|
||||
.row_meta_size = 4,
|
||||
|
|
@ -1613,11 +1613,16 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
|
|||
.from_float = quantize_row_iq3_kt,
|
||||
.from_float_ref = (ggml_from_float_t)quantize_row_iq3_kt_ref,
|
||||
.vec_dot = vec_dot_iq3_kt_q8_k,
|
||||
#ifdef __ARM_NEON
|
||||
.vec_dot_type = GGML_TYPE_F16,
|
||||
#if defined __AVX2__
|
||||
.vec_dot_type = GGML_TYPE_Q8_2_X4,
|
||||
#else
|
||||
.vec_dot_type = GGML_TYPE_F32,
|
||||
.vec_dot_type = GGML_TYPE_Q8_0_X4,
|
||||
#endif
|
||||
//#ifdef __ARM_NEON
|
||||
// .vec_dot_type = GGML_TYPE_F16,
|
||||
//#else
|
||||
// .vec_dot_type = GGML_TYPE_F32,
|
||||
//#endif
|
||||
.nrows = 1,
|
||||
.row_meta_size = 4,
|
||||
},
|
||||
|
|
@ -1630,13 +1635,13 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
|
|||
.from_float = quantize_row_iq4_kt,
|
||||
.from_float_ref = (ggml_from_float_t)quantize_row_iq4_kt_ref,
|
||||
.vec_dot = vec_dot_iq4_kt_q8_k,
|
||||
#ifdef __ARM_NEON
|
||||
.vec_dot_type = GGML_TYPE_F16,
|
||||
#if defined __AVX2__
|
||||
.vec_dot_type = GGML_TYPE_Q8_2_X4,
|
||||
#else
|
||||
.vec_dot_type = GGML_TYPE_F32,
|
||||
.vec_dot_type = GGML_TYPE_Q8_0_X4,
|
||||
#endif
|
||||
.nrows = 1,
|
||||
.row_meta_size = 8,
|
||||
.row_meta_size = 4,
|
||||
},
|
||||
[GGML_TYPE_IQ3_K] = {
|
||||
.type_name = "iq3_k",
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -236,9 +236,6 @@ struct MulMat {
|
|||
static inline ggml_type is_dequant_better(ggml_type type, int nrc_y) {
|
||||
#ifdef __AVX2__
|
||||
switch (type) {
|
||||
case GGML_TYPE_IQ2_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type;
|
||||
case GGML_TYPE_IQ3_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type;
|
||||
case GGML_TYPE_IQ4_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type;
|
||||
case GGML_TYPE_IQ2_XXS: return nrc_y >= 32 ? GGML_TYPE_Q8_K_R8 : type;
|
||||
case GGML_TYPE_IQ2_XS : return nrc_y >= 32 ? GGML_TYPE_Q8_K_R8 : type;
|
||||
case GGML_TYPE_IQ2_S : return nrc_y >= 16 ? GGML_TYPE_Q8_K_R8 : type;
|
||||
|
|
@ -267,13 +264,16 @@ struct MulMat {
|
|||
case GGML_TYPE_Q6_0 : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_IQ4_NL : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_Q8_0 : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_IQ2_KT : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_IQ3_KT : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_IQ4_KT : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
default: break;
|
||||
}
|
||||
#else
|
||||
switch (type) {
|
||||
case GGML_TYPE_IQ2_KT: return nrc_y >= 32 ? GGML_TYPE_F16 : type;
|
||||
case GGML_TYPE_IQ2_KT: return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
case GGML_TYPE_IQ3_KT: return nrc_y >= 32 ? GGML_TYPE_F16 : type;
|
||||
case GGML_TYPE_IQ4_KT: return nrc_y >= 32 ? GGML_TYPE_F16 : type;
|
||||
case GGML_TYPE_IQ4_KT: return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type;
|
||||
default: break;
|
||||
}
|
||||
#endif
|
||||
|
|
@ -815,7 +815,7 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) {
|
|||
case GGML_TYPE_IQ2_KT:
|
||||
case GGML_TYPE_IQ3_KT:
|
||||
case GGML_TYPE_IQ4_KT:
|
||||
return ggml_type(typeB) == GGML_TYPE_F32 ? iqk_set_kernels_ktquants(ne00, typeA, typeB, mm.funcs, mm.func16) : false;
|
||||
return iqk_set_kernels_ktquants(ne00, typeA, typeB, mm.funcs, mm.func16);
|
||||
case GGML_TYPE_Q4_0:
|
||||
case GGML_TYPE_Q4_1:
|
||||
case GGML_TYPE_Q5_0:
|
||||
|
|
|
|||
|
|
@ -7397,7 +7397,7 @@ void dequantize_row_ms_i2s(const void * vx, float * y, int64_t k) {
|
|||
}
|
||||
|
||||
namespace {
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs = false>
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs = false, bool is_int = false>
|
||||
class QuantizerIQKT {
|
||||
static_assert(group_size == 8 || group_size == 4);
|
||||
static_assert(block_size >= 8 && block_size%8 == 0);
|
||||
|
|
@ -7408,7 +7408,7 @@ public:
|
|||
constexpr static int kNg = kBlockSize/kGroupSize;
|
||||
constexpr static int kNblock = kSuperBlockSize/kBlockSize;
|
||||
constexpr static int kNumVal = 1 << num_bits; // i.e, 16 bits per group of 8
|
||||
constexpr static float kScale = 31.75f;
|
||||
constexpr static float kScale = is_int ? 1.f : 31.75f;
|
||||
constexpr static bool kVerbose = false;
|
||||
|
||||
QuantizerIQKT(int num_clusters, int num_neighbours, int offset = 4096);
|
||||
|
|
@ -7419,17 +7419,32 @@ public:
|
|||
inline float find_best_inverse_scale(const float * xb, const float * weight, const int * best_idx) const;
|
||||
|
||||
static inline void set_values(uint32_t i, float * result, float scale, int offset = 4096) {
|
||||
constexpr uint32_t ka = 89226354;
|
||||
constexpr uint32_t kb = 64248484;
|
||||
constexpr uint32_t kmask = 0x8fff8fff;
|
||||
constexpr uint32_t km32 = 0x3b603b60;
|
||||
uint32_t x = i + offset;
|
||||
for (int k = 0; k < kGroupSize; ++k) {
|
||||
x = ka*x + kb;
|
||||
uint32_t s = (x & kmask) ^ km32;
|
||||
float val = GGML_FP16_TO_FP32(s & 65535) + GGML_FP16_TO_FP32(s >> 16);
|
||||
if constexpr (is_abs) result[k] = scale*std::abs(val);
|
||||
else result[k] = scale*val;
|
||||
if constexpr (is_int) {
|
||||
constexpr uint32_t ka = 0xCBAC1FED;
|
||||
uint32_t s;
|
||||
auto i8 = (const int8_t *)&s;
|
||||
for (int k = 0; k < kGroupSize; ++k) {
|
||||
x = ka*x;
|
||||
s = x & 0x3f3f3f3f;
|
||||
if constexpr (is_abs) {
|
||||
result[k] = scale*std::abs(i8[0] + i8[1] + i8[2] + i8[3] - 126.f);
|
||||
} else {
|
||||
result[k] = scale*(i8[0] + i8[1] + i8[2] + i8[3] - 126.f);
|
||||
}
|
||||
}
|
||||
} else {
|
||||
constexpr uint32_t ka = 89226354;
|
||||
constexpr uint32_t kb = 64248484;
|
||||
constexpr uint32_t kmask = 0x8fff8fff;
|
||||
constexpr uint32_t km32 = 0x3b603b60;
|
||||
for (int k = 0; k < kGroupSize; ++k) {
|
||||
x = ka*x + kb;
|
||||
uint32_t s = (x & kmask) ^ km32;
|
||||
float val = GGML_FP16_TO_FP32(s & 65535) + GGML_FP16_TO_FP32(s >> 16);
|
||||
if constexpr (is_abs) result[k] = scale*std::abs(val);
|
||||
else result[k] = scale*val;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -7478,14 +7493,15 @@ private:
|
|||
float m_mid[4*kGroupSize];
|
||||
};
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
QuantizerIQKT<block_size, group_size, num_bits, is_abs>::QuantizerIQKT(int num_clusters, int num_neighbours, int offset) {
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::QuantizerIQKT(int num_clusters, int num_neighbours, int offset) {
|
||||
m_values.resize(kNumVal*kGroupSize);
|
||||
float * data = m_values.data();
|
||||
for (int i = 0; i < kNumVal; ++i) {
|
||||
set_values(i, data, kScale, offset);
|
||||
data += kGroupSize;
|
||||
}
|
||||
if (num_clusters == 0) return;
|
||||
// Make 128 clusters.
|
||||
// Note: we get a slightly better result by using 64 clusters
|
||||
// at the expense of almost doubling the quantization time.
|
||||
|
|
@ -7494,8 +7510,8 @@ QuantizerIQKT<block_size, group_size, num_bits, is_abs>::QuantizerIQKT(int num_c
|
|||
m_in_cluster = finalize_clusters(num_neighbours, m_values, m_clusters, m_c_values);
|
||||
}
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
std::pair<float, float> QuantizerIQKT<block_size, group_size, num_bits, is_abs>::find_best_scale(
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
std::pair<float, float> QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::find_best_scale(
|
||||
const float * xb, const float * weight, const int * best_idx) const {
|
||||
float sumqx = 0, sumq2 = 0;
|
||||
#ifdef __AVX2__
|
||||
|
|
@ -7527,8 +7543,8 @@ std::pair<float, float> QuantizerIQKT<block_size, group_size, num_bits, is_abs>:
|
|||
return sumq2 > 0 ? std::make_pair(sumqx/sumq2, sumqx*sumqx/sumq2) : std::make_pair(0.f, 0.f);
|
||||
}
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
float QuantizerIQKT<block_size, group_size, num_bits, is_abs>::find_best_inverse_scale(
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
float QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::find_best_inverse_scale(
|
||||
const float * xb, const float * weight, const int * best_idx) const {
|
||||
float sumqx = 0, sumx2 = 0;
|
||||
#ifdef __AVX2__
|
||||
|
|
@ -7560,8 +7576,8 @@ float QuantizerIQKT<block_size, group_size, num_bits, is_abs>::find_best_inverse
|
|||
return sumx2 > 0 ? sumqx/sumx2 : 0.f;
|
||||
}
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
void QuantizerIQKT<block_size, group_size, num_bits, is_abs>::find_best_match(float d, const float * xb, const float * weight, int * best_idx) const {
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
void QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::find_best_match(float d, const float * xb, const float * weight, int * best_idx) const {
|
||||
if (!d) {
|
||||
std::memset(best_idx, 0, kNg*sizeof(int));
|
||||
return;
|
||||
|
|
@ -7739,8 +7755,8 @@ void QuantizerIQKT<block_size, group_size, num_bits, is_abs>::find_best_match(fl
|
|||
#endif
|
||||
}
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
std::vector<std::vector<int>> QuantizerIQKT<block_size, group_size, num_bits, is_abs>::finalize_clusters(int num_neighbours,
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
std::vector<std::vector<int>> QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::finalize_clusters(int num_neighbours,
|
||||
const std::vector<float>& values, const std::vector<float>& clusters, std::vector<std::vector<float>>& c_values) {
|
||||
int ncluster = clusters.size()/kGroupSize;
|
||||
std::vector<std::vector<int>> p_in_cluster(ncluster);
|
||||
|
|
@ -7826,8 +7842,8 @@ std::vector<std::vector<int>> QuantizerIQKT<block_size, group_size, num_bits, is
|
|||
return p_in_cluster;
|
||||
}
|
||||
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs>
|
||||
std::vector<float> QuantizerIQKT<block_size, group_size, num_bits, is_abs>::cluster_points(const std::vector<float>& points, int ncluster, int niter, float * mid) {
|
||||
template <int block_size, int group_size, int num_bits, bool is_abs, bool is_int>
|
||||
std::vector<float> QuantizerIQKT<block_size, group_size, num_bits, is_abs, is_int>::cluster_points(const std::vector<float>& points, int ncluster, int niter, float * mid) {
|
||||
constexpr int ndim = kGroupSize;
|
||||
GGML_ASSERT(points.size() % ndim == 0);
|
||||
int npoint = points.size() / ndim;
|
||||
|
|
@ -7995,7 +8011,7 @@ std::vector<float> QuantizerIQKT<block_size, group_size, num_bits, is_abs>::clus
|
|||
|
||||
// ========================================== iq2_kt ====================================================
|
||||
|
||||
using QuantizerIQ2KT = QuantizerIQKT<32, 8, 16>;
|
||||
using QuantizerIQ2KT = QuantizerIQKT<32, 8, 16, false, true>;
|
||||
|
||||
const QuantizerIQ2KT& iq2kt_quantizer() {
|
||||
static std::mutex mutex;
|
||||
|
|
@ -8006,7 +8022,7 @@ const QuantizerIQ2KT& iq2kt_quantizer() {
|
|||
}
|
||||
|
||||
void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const float * quant_weights, float * all_scales, float * all_weights,
|
||||
float * qtmp) {
|
||||
int * all_idx) {
|
||||
|
||||
constexpr float kSigmaScale = 2.0f;
|
||||
using Q = QuantizerIQ2KT;
|
||||
|
|
@ -8025,6 +8041,11 @@ void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
|
||||
Q::set_weights(kSigmaScale, nblock, x, quant_weights, all_weights);
|
||||
|
||||
float amax_row = 0;
|
||||
for (int j = 0; j < n_per_row; ++j) {
|
||||
amax_row = std::max(amax_row, std::abs(x[j]));
|
||||
}
|
||||
|
||||
float amax_scale = 0, max_scale = 0;
|
||||
|
||||
for (int ibl = 0; ibl < nblock; ++ibl) {
|
||||
|
|
@ -8042,9 +8063,10 @@ void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
float ax = std::abs(xb[j]);
|
||||
amax = std::max(amax, ax);
|
||||
}
|
||||
quantizer.find_best_match( amax/96.f, xb, weight, best_idx);
|
||||
float scale_0 = std::max(90.f, 124.f*amax/amax_row);
|
||||
quantizer.find_best_match( amax/scale_0, xb, weight, best_idx);
|
||||
auto [dp, score_p] = quantizer.find_best_scale(xb, weight, best_idx);
|
||||
quantizer.find_best_match(-amax/96.f, xb, weight, best_idx + Q::kNg);
|
||||
quantizer.find_best_match(-amax/scale_0, xb, weight, best_idx + Q::kNg);
|
||||
auto [dm, score_m] = quantizer.find_best_scale(xb, weight, best_idx + Q::kNg);
|
||||
|
||||
auto idx = best_idx;
|
||||
|
|
@ -8052,12 +8074,7 @@ void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
else {
|
||||
scales[ib] = dm; idx += Q::kNg;
|
||||
}
|
||||
auto qt = qtmp + ibl*Q::kSuperBlockSize + ib*Q::kBlockSize;
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) {
|
||||
auto q = quantizer.values() + idx[ig]*Q::kGroupSize;
|
||||
for (int j = 0; j < Q::kGroupSize; ++j) qt[j] = q[j];
|
||||
qt += Q::kGroupSize;
|
||||
}
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) all_idx[(ibl*Q::kSuperBlockSize + ib*Q::kBlockSize)/Q::kGroupSize + ig] = idx[ig];
|
||||
|
||||
float abs_scale = std::abs(scales[ib]);
|
||||
if (abs_scale > amax_scale) {
|
||||
|
|
@ -8080,20 +8097,22 @@ void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
float sumqx = 0, sumq2 = 0;
|
||||
for (int ibl = 0; ibl < nblock; ++ibl) {
|
||||
const float * xb = x + ibl*Q::kSuperBlockSize;
|
||||
const float * qb = qtmp + ibl*Q::kSuperBlockSize;
|
||||
const float * wb = all_weights + ibl*Q::kSuperBlockSize;
|
||||
auto scales = all_scales + ibl*Q::kNblock;
|
||||
for (int ib = 0; ib < Q::kNblock; ++ib) {
|
||||
int ls = best_index_iq4nl(iq4k_values, id*scales[ib]);
|
||||
float dl = iq4k_values[ls];
|
||||
for (int j = 0; j < Q::kBlockSize; ++j) {
|
||||
float q = dl*qb[j];
|
||||
sumqx += wb[j]*xb[j]*q;
|
||||
sumq2 += wb[j]*q*q;
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) {
|
||||
auto qb = quantizer.values() + Q::kGroupSize*all_idx[(ibl*Q::kSuperBlockSize + ib*Q::kBlockSize)/Q::kGroupSize + ig];
|
||||
for (int j = 0; j < Q::kGroupSize; ++j) {
|
||||
int jj = ig*Q::kGroupSize + j;
|
||||
float q = dl*qb[j];
|
||||
sumqx += wb[jj]*xb[jj]*q;
|
||||
sumq2 += wb[jj]*q*q;
|
||||
}
|
||||
}
|
||||
xb += Q::kBlockSize;
|
||||
wb += Q::kBlockSize;
|
||||
qb += Q::kBlockSize;
|
||||
}
|
||||
}
|
||||
if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
|
||||
|
|
@ -8129,6 +8148,26 @@ void quantize_row_iq2_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
float dl = d*ls;
|
||||
quantizer.find_best_match(dl, xb, weight, best_idx);
|
||||
|
||||
auto prev_idx = all_idx + (ibl*Q::kSuperBlockSize + ib*Q::kBlockSize)/Q::kGroupSize;
|
||||
|
||||
float mse1 = 0, mse2 = 0;
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) {
|
||||
auto q1 = quantizer.values() + Q::kGroupSize*prev_idx[ig];
|
||||
auto q2 = quantizer.values() + Q::kGroupSize*best_idx[ig];
|
||||
for (int j = 0; j < Q::kGroupSize; ++j) {
|
||||
int jj = ig*Q::kGroupSize + j;
|
||||
float diff1 = xb[jj] - dl*q1[j];
|
||||
float diff2 = xb[jj] - dl*q2[j];
|
||||
mse1 += weight[jj]*diff1*diff1;
|
||||
mse2 += weight[jj]*diff2*diff2;
|
||||
}
|
||||
}
|
||||
if (mse1 < mse2) {
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) best_idx[ig] = prev_idx[ig];
|
||||
} else {
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) prev_idx[ig] = best_idx[ig];
|
||||
}
|
||||
|
||||
for (int j = 0; j < Q::kNg; ++j) {
|
||||
qs[j] = best_idx[j];
|
||||
auto xl = xb + Q::kGroupSize*j;
|
||||
|
|
@ -8196,10 +8235,10 @@ size_t quantize_iq2_kt(const float * src, void * dst, int64_t nrows, int64_t n_p
|
|||
auto row_size = ggml_row_size(GGML_TYPE_IQ2_KT, n_per_row);
|
||||
std::vector<float> scales(n_per_row/QuantizerIQ2KT::kBlockSize);
|
||||
std::vector<float> weights(n_per_row);
|
||||
std::vector<float> xtmp(n_per_row);
|
||||
std::vector<int> idx(n_per_row/QuantizerIQ2KT::kGroupSize);
|
||||
char * qrow = (char *)dst;
|
||||
for (int64_t row = 0; row < nrows; ++row) {
|
||||
quantize_row_iq2_kt_impl(src, (void *)qrow, n_per_row, imatrix, scales.data(), weights.data(), xtmp.data());
|
||||
quantize_row_iq2_kt_impl(src, (void *)qrow, n_per_row, imatrix, scales.data(), weights.data(), idx.data());
|
||||
src += n_per_row;
|
||||
qrow += row_size;
|
||||
}
|
||||
|
|
@ -8209,7 +8248,7 @@ size_t quantize_iq2_kt(const float * src, void * dst, int64_t nrows, int64_t n_p
|
|||
void dequantize_row_iq2_kt(const block_iq2_kt * x, float * y, int64_t k) {
|
||||
assert(k % QuantizerIQ2KT::kSuperBlockSize == 0);
|
||||
#ifdef __AVX2__
|
||||
if (iqk_dequantize_ktquants(GGML_TYPE_IQ2_KT, k, x, 0, y, 0, 1)) return;
|
||||
//if (iqk_dequantize_ktquants(GGML_TYPE_IQ2_KT, k, x, 0, y, 0, 1)) return;
|
||||
#endif
|
||||
const int nb = k / QuantizerIQ2KT::kSuperBlockSize;
|
||||
const float * dptr = (const float *)x;
|
||||
|
|
@ -8254,7 +8293,7 @@ void vec_dot_iq2_kt_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx
|
|||
|
||||
namespace {
|
||||
|
||||
using QuantizerIQ3KT = QuantizerIQKT<32, 8, 16, true>;
|
||||
using QuantizerIQ3KT = QuantizerIQKT<32, 8, 16, true, true>;
|
||||
const QuantizerIQ3KT& iq3kt_quantizer() {
|
||||
static std::mutex mutex;
|
||||
std::lock_guard<std::mutex> lock(mutex);
|
||||
|
|
@ -8465,7 +8504,7 @@ size_t quantize_iq3_kt(const float * src, void * dst, int64_t nrows, int64_t n_p
|
|||
|
||||
void dequantize_row_iq3_kt(const block_iq3_kt * x, float * y, int64_t k) {
|
||||
#ifdef __AVX2__
|
||||
if (iqk_dequantize_ktquants(GGML_TYPE_IQ3_KT, k, x, 0, y, 0, 1)) return;
|
||||
//if (iqk_dequantize_ktquants(GGML_TYPE_IQ3_KT, k, x, 0, y, 0, 1)) return;
|
||||
#endif
|
||||
using Q = QuantizerIQ3KT;
|
||||
constexpr int kNumGroups = Q::kSuperBlockSize/Q::kGroupSize;
|
||||
|
|
@ -8521,7 +8560,7 @@ void vec_dot_iq3_kt_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx
|
|||
|
||||
namespace{
|
||||
|
||||
using QuantizerIQ4KT = QuantizerIQKT<32, 4, 15>;
|
||||
using QuantizerIQ4KT = QuantizerIQKT<32, 4, 15, false, true>;
|
||||
|
||||
const QuantizerIQ4KT& iq4kt_quantizer(bool with_offset = false) {
|
||||
static std::mutex mutex;
|
||||
|
|
@ -8536,6 +8575,14 @@ const QuantizerIQ4KT& iq4kt_quantizer(bool with_offset = false) {
|
|||
return *quantizer1;
|
||||
}
|
||||
|
||||
const QuantizerIQ4KT& iq4kt_dequantizer() {
|
||||
static std::mutex mutex;
|
||||
std::lock_guard<std::mutex> lock(mutex);
|
||||
static std::unique_ptr<QuantizerIQ4KT> dequantizer;
|
||||
if (!dequantizer) dequantizer = std::make_unique<QuantizerIQ4KT>(0, 0, 4096);
|
||||
return *dequantizer;
|
||||
}
|
||||
|
||||
void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const float * quant_weights, float * all_scales, float * all_weights) {
|
||||
|
||||
constexpr float kSigmaScale = 2.0f;
|
||||
|
|
@ -8546,7 +8593,7 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
|
||||
float * dptr = (float *)vy;
|
||||
|
||||
block_iq4_kt * y = (block_iq4_kt *)(dptr + 2);
|
||||
block_iq4_kt * y = (block_iq4_kt *)(dptr + 1);
|
||||
|
||||
auto& quantizer1 = iq4kt_quantizer();
|
||||
auto& quantizer2 = iq4kt_quantizer(true);
|
||||
|
|
@ -8555,13 +8602,10 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
|
||||
Q::set_weights(kSigmaScale, nblock, x, quant_weights, all_weights);
|
||||
|
||||
float amax_row = 0, row_av = 0;
|
||||
float amax_row = 0;
|
||||
for (int j = 0; j < n_per_row; ++j) {
|
||||
row_av += x[j];
|
||||
amax_row = std::max(amax_row, std::abs(x[j]));
|
||||
}
|
||||
row_av /= n_per_row;
|
||||
dptr[1] = row_av;
|
||||
if (!amax_row) {
|
||||
dptr[0] = 0.f;
|
||||
std::memset(y, 0, nblock*sizeof(block_iq4_kt));
|
||||
|
|
@ -8584,7 +8628,7 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
const float * weight = all_weights + ibl*Q::kSuperBlockSize + ib*Q::kBlockSize;
|
||||
float amax = 0;
|
||||
for (int j = 0; j < Q::kBlockSize; ++j) {
|
||||
xaux[j] = xbl[ib*Q::kBlockSize+j] - row_av;
|
||||
xaux[j] = xbl[ib*Q::kBlockSize+j];
|
||||
float ax = std::abs(xaux[j]);
|
||||
amax = std::max(amax, ax);
|
||||
}
|
||||
|
|
@ -8593,7 +8637,7 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
continue;
|
||||
}
|
||||
float best = 0;
|
||||
float scale_0 = std::max(92.f, 127.f*amax/amax_row);
|
||||
float scale_0 = std::max(90.f, 124.f*amax/amax_row);
|
||||
for (int itry = -kNtry; itry <= kNtry; ++itry) {
|
||||
quantizer1.find_best_match( amax/(8.f*itry + scale_0), xaux, weight, best_idx);
|
||||
auto [dp, score_p] = quantizer1.find_best_scale(xaux, weight, best_idx);
|
||||
|
|
@ -8664,7 +8708,7 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f
|
|||
for (int ib = 0; ib < Q::kNblock; ++ib) {
|
||||
auto& quantizer = y[ibl].qs[ib] & 1 ? quantizer2 : quantizer1;
|
||||
const float * weight = all_weights + ibl*Q::kSuperBlockSize + ib*Q::kBlockSize;
|
||||
for (int j = 0; j < Q::kBlockSize; ++j) xaux[j] = xbl[ib*Q::kBlockSize+j] - row_av;
|
||||
for (int j = 0; j < Q::kBlockSize; ++j) xaux[j] = xbl[ib*Q::kBlockSize+j];
|
||||
int ls = nearest_int(id*scales[ib]);
|
||||
ls = std::min(ls, 63);
|
||||
*(uint8_t *)(shb + ib) = ((ls + 64) << 1) | (shb[ib] & 1);
|
||||
|
|
@ -8724,7 +8768,7 @@ size_t quantize_iq4_kt(const float * src, void * dst, int64_t nrows, int64_t n_p
|
|||
|
||||
void dequantize_row_iq4_kt(const block_iq4_kt * x, float * y, int64_t k) {
|
||||
#ifdef __AVX2__
|
||||
if (iqk_dequantize_ktquants(GGML_TYPE_IQ4_KT, k, x, 0, y, 0, 1)) return;
|
||||
//if (iqk_dequantize_ktquants(GGML_TYPE_IQ4_KT, k, x, 0, y, 0, 1)) return;
|
||||
#endif
|
||||
using Q = QuantizerIQ4KT;
|
||||
assert(k % Q::kSuperBlockSize == 0);
|
||||
|
|
@ -8732,23 +8776,20 @@ void dequantize_row_iq4_kt(const block_iq4_kt * x, float * y, int64_t k) {
|
|||
const int nb = k / Q::kSuperBlockSize;
|
||||
const float * dptr = (const float *)x;
|
||||
const float d = dptr[0] * Q::kScale;
|
||||
const float row_av = dptr[1];
|
||||
x = (const block_iq4_kt *)(dptr + 2);
|
||||
auto& deq = iq4kt_quantizer();
|
||||
x = (const block_iq4_kt *)(dptr + 1);
|
||||
auto& deq = iq4kt_dequantizer();
|
||||
for (int ibl = 0; ibl < nb; ++ibl) {
|
||||
auto shb = x[ibl].qs;
|
||||
auto ql = (const uint8_t *)(shb + Q::kNblock);
|
||||
auto qh = ql + kNumGroups;
|
||||
for (int ib = 0; ib < Q::kNblock; ++ib) {
|
||||
int offset = shb[ib] & 1 ? 32768 + 4096 : 4096;
|
||||
//auto& deq = shb[ib] & 1 ? deq2 : deq1;
|
||||
int ls = int((shb[ib] & 0xff) >> 1) - 64;
|
||||
float sl = d * ls;
|
||||
for (int ig = 0; ig < Q::kNg; ++ig) {
|
||||
int jj = ib*Q::kNg+ig;
|
||||
uint16_t idx = ql[jj] | ((qh[jj%(kNumGroups/2)] << (8 - 4*(jj/(kNumGroups/2)))) & 0xf00) | (((shb[ib] >> (8 + 3*ig)) & 7) << 12);
|
||||
deq.set_values(idx, y, sl, offset);
|
||||
for (int j = 0; j < Q::kGroupSize; ++j) y[j] += row_av;
|
||||
y += Q::kGroupSize;
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -18627,6 +18627,7 @@ static ggml_type change_type_if_necessary(ggml_type new_type, int nx, int ny) {
|
|||
new_type == GGML_TYPE_IQ2_K_R4|| new_type == GGML_TYPE_IQ5_K_R4|| new_type == GGML_TYPE_IQ4_KS_R4 ||
|
||||
new_type == GGML_TYPE_IQ3_XXS_R4 || new_type == GGML_TYPE_IQ2_XXS_R4 || new_type == GGML_TYPE_IQ2_XS_R4 ||
|
||||
new_type == GGML_TYPE_IQ2_S_R4|| new_type == GGML_TYPE_IQ3_S_R4||
|
||||
new_type == GGML_TYPE_IQ2_KT || new_type == GGML_TYPE_IQ3_KT || new_type == GGML_TYPE_IQ4_KT ||
|
||||
new_type == GGML_TYPE_IQ5_KS || new_type == GGML_TYPE_IQ5_KS_R4) {
|
||||
if (nx % QK_K != 0) {
|
||||
LLAMA_LOG_WARN("\n\n%s : tensor cols %d x %d are not divisible by %d, required for %s", __func__, nx, ny, QK_K, ggml_type_name(new_type));
|
||||
|
|
|
|||
Loading…
Reference in New Issue