Skip to content

Commit 1f06a2d

Browse files
AndgihatAndgihat
committed
CUDA: add Q2_0 (ternary, group-64) weight support
1 parent 86ec665 commit 1f06a2d

17 files changed

Lines changed: 297 additions & 0 deletions

‎ggml/src/ggml-cuda/common.cuh‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -977,6 +977,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q1_0> {
977977
static constexpr int qi = QI1_0;
978978
};
979979

980+
template<>
981+
struct ggml_cuda_type_traits<GGML_TYPE_Q2_0> {
982+
static constexpr int qk = QK2_0;
983+
static constexpr int qr = QR2_0;
984+
static constexpr int qi = QI2_0;
985+
};
986+
980987
template<>
981988
struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {
982989
static constexpr int qk = QK4_0;

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -508,6 +508,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
508508
switch (type) {
509509
case GGML_TYPE_Q1_0:
510510
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
511+
case GGML_TYPE_Q2_0:
512+
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
511513
case GGML_TYPE_Q4_0:
512514
return dequantize_row_q4_0_cuda;
513515
case GGML_TYPE_Q4_1:
@@ -563,6 +565,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
563565
switch (type) {
564566
case GGML_TYPE_Q1_0:
565567
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
568+
case GGML_TYPE_Q2_0:
569+
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
566570
case GGML_TYPE_Q4_0:
567571
return dequantize_row_q4_0_cuda;
568572
case GGML_TYPE_Q4_1:
@@ -633,6 +637,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
633637
switch (type) {
634638
case GGML_TYPE_Q1_0:
635639
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
640+
case GGML_TYPE_Q2_0:
641+
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
636642
case GGML_TYPE_Q4_0:
637643
return dequantize_row_q4_0_cuda;
638644
case GGML_TYPE_Q4_1:
@@ -702,6 +708,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
702708
return convert_unary_cuda<float>;
703709
case GGML_TYPE_Q1_0:
704710
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
711+
case GGML_TYPE_Q2_0:
712+
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
705713
case GGML_TYPE_Q4_0:
706714
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
707715
case GGML_TYPE_Q4_1:
@@ -739,6 +747,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
739747
return convert_unary_cuda<float, nv_bfloat16>;
740748
case GGML_TYPE_Q1_0:
741749
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
750+
case GGML_TYPE_Q2_0:
751+
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
742752
case GGML_TYPE_Q4_0:
743753
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
744754
case GGML_TYPE_Q4_1:
@@ -774,6 +784,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
774784
return convert_unary_cuda<half, float>;
775785
case GGML_TYPE_Q1_0:
776786
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
787+
case GGML_TYPE_Q2_0:
788+
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
777789
case GGML_TYPE_Q4_0:
778790
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
779791
case GGML_TYPE_Q4_1:

‎ggml/src/ggml-cuda/dequantize.cuh‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,26 @@ static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const in
3232
v.y = (2*bit_1 - 1) * d;
3333
}
3434

35+
static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
36+
const block_q2_0 * x = (const block_q2_0 *) vx;
37+
38+
const float d = x[ib].d;
39+
40+
// Q2_0: 2 bits per element, 4 elements per byte.
41+
// Stored code c in {0,1,2,3} maps to symbol s = c - 1 in {-1, 0, +1, +2}.
42+
const int byte_index_0 = iqs / 4;
43+
const int bit_offset_0 = (iqs % 4) * 2;
44+
45+
const int byte_index_1 = (iqs + 1) / 4;
46+
const int bit_offset_1 = ((iqs + 1) % 4) * 2;
47+
48+
const int c0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 0x3;
49+
const int c1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 0x3;
50+
51+
v.x = (c0 - 1) * d;
52+
v.y = (c1 - 1) * d;
53+
}
54+
3555
static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
3656
const block_q4_0 * x = (const block_q4_0 *) vx;
3757

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -320,6 +320,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
320320
get_rows_cuda_q<QK1_0, QR1_0, dequantize_q1_0>(src0_d, src1_d, dst_d,
321321
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
322322
break;
323+
case GGML_TYPE_Q2_0:
324+
get_rows_cuda_q<QK2_0, QR2_0, dequantize_q2_0>(src0_d, src1_d, dst_d,
325+
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
326+
break;
323327
case GGML_TYPE_Q4_0:
324328
get_rows_cuda_q<QK4_0, QR4_0, dequantize_q4_0>(src0_d, src1_d, dst_d,
325329
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4944,6 +4944,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
49444944
case GGML_TYPE_F32:
49454945
case GGML_TYPE_F16:
49464946
case GGML_TYPE_Q1_0:
4947+
case GGML_TYPE_Q2_0:
49474948
case GGML_TYPE_Q4_0:
49484949
case GGML_TYPE_Q4_1:
49494950
case GGML_TYPE_Q5_0:
@@ -5002,6 +5003,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
50025003
case GGML_TYPE_BF16:
50035004
case GGML_TYPE_I32:
50045005
case GGML_TYPE_Q1_0:
5006+
case GGML_TYPE_Q2_0:
50055007
case GGML_TYPE_Q4_0:
50065008
case GGML_TYPE_Q4_1:
50075009
case GGML_TYPE_Q5_0:

‎ggml/src/ggml-cuda/mmq-config-ampere.cuh‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,23 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
1616
CASE(GGML_TYPE_Q1_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
1717
CASE(GGML_TYPE_Q1_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
1818

19+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
20+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
21+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
22+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
23+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
24+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
25+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
26+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
27+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
28+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
29+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
30+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
31+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
32+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
33+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
34+
CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
35+
1936
CASE(GGML_TYPE_Q4_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
2037
CASE(GGML_TYPE_Q4_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
2138
CASE(GGML_TYPE_Q4_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);

‎ggml/src/ggml-cuda/mmq-config-cdna.cuh‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,14 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
77
CASE(GGML_TYPE_Q1_0, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
88
CASE(GGML_TYPE_Q1_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
99

10+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
11+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
12+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
13+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
14+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
15+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
16+
CASE(GGML_TYPE_Q2_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
17+
1018
CASE(GGML_TYPE_Q4_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
1119
CASE(GGML_TYPE_Q4_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
1220
CASE(GGML_TYPE_Q4_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);

‎ggml/src/ggml-cuda/mmq-config-pascal.cuh‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
1111
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1212
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1313

14+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
15+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
16+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
17+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
18+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
19+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
20+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
21+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
22+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
23+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
24+
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
25+
1426
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1527
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1628
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

‎ggml/src/ggml-cuda/mmq-config-rdna2.cuh‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
1111
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1212
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1313

14+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
15+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
16+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
17+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
18+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
19+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
20+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
21+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
22+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
23+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
24+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
25+
1426
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1527
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1628
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

‎ggml/src/ggml-cuda/mmq-config-rdna4.cuh‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,19 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
1212
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1313
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
1414

15+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
16+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
17+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
18+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
19+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
20+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
21+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
22+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
23+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
24+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
25+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
26+
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
27+
1528
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1629
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
1730
CASE(GGML_TYPE_Q4_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

0 commit comments

Comments
 (0)