Skip to content

Commit a6a0ee5

Browse files
NexesenexNexesenex
committed
Full alignment of MMVQ on IK_Llama
1 parent 2e1d660 commit a6a0ee5

File tree

10 files changed

+3027
-3938
lines changed

10 files changed

+3027
-3938
lines changed

‎ggml/src/ggml-cpu/ggml-cpu.c

+17-17
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313
// #include "amx/amx.h"
1414
#include "ggml.h"
1515

16-
// #include "iqk/iqk_quantize.h"
16+
#include "iqk_croco/iqk_quantize_croco.h"
1717

1818
#if defined(_MSC_VER) || defined(__MINGW32__)
1919
#include <malloc.h> // using malloc.h with MSC/MINGW
@@ -398,13 +398,13 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
398398
},
399399
// [GGML_TYPE_IQ1_BN] = {
400400
// .from_float = quantize_row_iq1_bn,
401-
// .vec_dot = ggml_vec_dot_iq1_bn_q8_0,
401+
// .vec_dot = vec_dot_iq1_bn_q8_0,
402402
// .vec_dot_type = GGML_TYPE_IQ1_BN,
403403
// .nrows = 1,
404404
// },
405405
// [GGML_TYPE_IQ2_BN] = {
406406
// .from_float = quantize_row_iq2_bn,
407-
// .vec_dot = ggml_vec_dot_iq2_bn_q8_0,
407+
// .vec_dot = vec_dot_iq2_bn_q8_0,
408408
// .vec_dot_type = GGML_TYPE_IQ2_BN,
409409
// .nrows = 1,
410410
// },
@@ -422,67 +422,67 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
422422
},
423423
[GGML_TYPE_IQ4_KS] = {
424424
.from_float = quantize_row_iq4_ks,
425-
.vec_dot = ggml_vec_dot_iq4_ks_q8_K,
425+
.vec_dot = vec_dot_iq4_ks_q8_k,
426426
.vec_dot_type = GGML_TYPE_Q8_K,
427427
.nrows = 1,
428428
},
429429
[GGML_TYPE_IQ4_KSS] = {
430430
.from_float = quantize_row_iq4_kss,
431-
.vec_dot = ggml_vec_dot_iq4_kss_q8_K,
431+
.vec_dot = vec_dot_iq4_kss_q8_k,
432432
.vec_dot_type = GGML_TYPE_Q8_K,
433433
.nrows = 1,
434434
},
435435
[GGML_TYPE_IQ2_K] = {
436436
.from_float = quantize_row_iq2_k,
437-
.vec_dot = ggml_vec_dot_iq2_k_q8_K,
437+
.vec_dot = vec_dot_iq2_k_q8_k,
438438
.vec_dot_type = GGML_TYPE_Q8_K,
439439
.nrows = 1,
440440
},
441441
[GGML_TYPE_IQ2_KS] = {
442442
.from_float = quantize_row_iq2_ks,
443-
.vec_dot = ggml_vec_dot_iq2_ks_q8_K,
443+
.vec_dot = vec_dot_iq2_ks_q8_k,
444444
.vec_dot_type = GGML_TYPE_Q8_K,
445445
.nrows = 1,
446446
},
447447
[GGML_TYPE_IQ2_KT] = {
448-
.from_float = quantize_row_iq2_kt,
449-
.vec_dot = ggml_vec_dot_iq2_kt_q8_K,
448+
// .from_float = quantize_row_iq2_kt,
449+
// .vec_dot = vec_dot_iq2_kt_q8_k,
450450
.vec_dot_type = GGML_TYPE_Q8_K,
451451
.nrows = 1,
452452
},
453453
[GGML_TYPE_IQ3_KT] = {
454-
.from_float = quantize_row_iq3_kt,
455-
.vec_dot = ggml_vec_dot_iq3_kt_q8_K,
454+
// .from_float = quantize_row_iq3_kt,
455+
// .vec_dot = vec_dot_iq3_kt_q8_k,
456456
.vec_dot_type = GGML_TYPE_Q8_K,
457457
.nrows = 1,
458458
},
459459
[GGML_TYPE_IQ4_KT] = {
460-
.from_float = quantize_row_iq4_kt,
461-
.vec_dot = ggml_vec_dot_iq4_kt_q8_K,
460+
// .from_float = quantize_row_iq4_kt,
461+
// .vec_dot = vec_dot_iq4_kt_q8_k,
462462
.vec_dot_type = GGML_TYPE_Q8_K,
463463
.nrows = 1,
464464
},
465465
[GGML_TYPE_IQ3_K] = {
466466
.from_float = quantize_row_iq3_k,
467-
.vec_dot = ggml_vec_dot_iq3_k_q8_K,
467+
.vec_dot = vec_dot_iq3_k_q8_k,
468468
.vec_dot_type = GGML_TYPE_Q8_K,
469469
.nrows = 1,
470470
},
471471
[GGML_TYPE_IQ4_K] = {
472472
.from_float = quantize_row_iq4_k,
473-
.vec_dot = ggml_vec_dot_iq4_k_q8_K,
473+
.vec_dot = vec_dot_iq4_k_q8_k,
474474
.vec_dot_type = GGML_TYPE_Q8_K,
475475
.nrows = 1,
476476
},
477477
[GGML_TYPE_IQ5_K] = {
478478
.from_float = quantize_row_iq5_k,
479-
.vec_dot = ggml_vec_dot_iq5_k_q8_K,
479+
.vec_dot = vec_dot_iq5_k_q8_k,
480480
.vec_dot_type = GGML_TYPE_Q8_K,
481481
.nrows = 1,
482482
},
483483
[GGML_TYPE_IQ6_K] = {
484484
.from_float = quantize_row_iq6_k,
485-
.vec_dot = ggml_vec_dot_iq6_k_q8_K,
485+
.vec_dot = vec_dot_iq6_k_q8_k,
486486
.vec_dot_type = GGML_TYPE_Q8_K,
487487
.nrows = 1,
488488
},

‎ggml/src/ggml-cuda/iqk_mmvq.cu

+1-16
Original file line numberDiff line numberDiff line change
@@ -6,8 +6,6 @@
66

77
#include "iqk_mmvq.cuh"
88

9-
// #include "mmvq.cu"
10-
119
typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
1210

1311
// Reminder:
@@ -109,8 +107,7 @@ void iqk_mul_mat_vec_q_cuda(
109107
int64_t nwarps = 1;
110108
int64_t rows_per_cuda_block = 1;
111109

112-
if (ggml_cuda_info().devices[id].cc < GGML_CUDA_CC_CDNA ||
113-
ggml_cuda_info().devices[id].cc == GGML_CUDA_CC_RDNA1) { // NVIDIA and AMD older than RDNA2
110+
if (ggml_cuda_info().devices[id].cc < GGML_CUDA_CC_CDNA || ggml_cuda_info().devices[id].cc == GGML_CUDA_CC_RDNA1) { // NVIDIA and AMD older than RDNA2 but not CDNA
114111
switch(ncols_y) {
115112
case 1:
116113
nwarps = 4;
@@ -545,11 +542,6 @@ __device__ __forceinline__ float vec_dot_iq2_ks_q8_1(
545542
+ __low2float(bq8_1[4*(i4/4)+3].ds) * sumi4);
546543
}
547544

548-
__device__ __forceinline__ float vec_dot_iq2_kt_q8_1(
549-
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
550-
return 0.f;
551-
}
552-
553545
#define VDR_IQ3_K_Q8_1_MMVQ 4
554546
#define VDR_IQ3_K_Q8_1_MMQ 4
555547

@@ -778,13 +770,6 @@ void mul_mat_vec_iq2_ks_q8_1_cuda(
778770
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ2_KS, VDR_IQ2_KS_Q8_1_MMVQ, vec_dot_iq2_ks_q8_1>(vx, vy, dst, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, stream);
779771
}
780772

781-
void mul_mat_vec_iq2_kt_q8_1_cuda(
782-
const void * vx, const void * vy, float * dst,
783-
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst, cudaStream_t stream) {
784-
785-
iqk_mul_mat_vec_q_cuda<GGML_TYPE_IQ2_KT, VDR_IQ2_KS_Q8_1_MMVQ, vec_dot_iq2_kt_q8_1>(vx, vy, dst, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, stream);
786-
}
787-
788773
void mul_mat_vec_iq5_k_q8_1_cuda(
789774
const void * vx, const void * vy, float * dst,
790775
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst, cudaStream_t stream) {

0 commit comments

Comments
 (0)