Skip to content

Commit bc8122e

Browse files
khosravipashakhosravipasha
authored
Merge pull request #1 from PrismML-Eng/mmq
[cuda] Fix mmq/mma path
2 parents 59f2b84 + 43e67bd commit bc8122e

7 files changed

Lines changed: 128 additions & 71 deletions

File tree

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

Lines changed: 12 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -6,12 +6,12 @@
66
static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {
77
switch (args.type_x) {
88
// TODO: Q1_0/Q1_0_g128 MMQ disabled due to accuracy issues; for now commenting these to use cuBLAS fallback
9-
// case GGML_TYPE_Q1_0:
10-
// mul_mat_q_case<GGML_TYPE_Q1_0>(ctx, args, stream);
11-
// break;
12-
// case GGML_TYPE_Q1_0_g128:
13-
// mul_mat_q_case<GGML_TYPE_Q1_0_g128>(ctx, args, stream);
14-
// break;
9+
case GGML_TYPE_Q1_0:
10+
mul_mat_q_case<GGML_TYPE_Q1_0>(ctx, args, stream);
11+
break;
12+
case GGML_TYPE_Q1_0_g128:
13+
mul_mat_q_case<GGML_TYPE_Q1_0_g128>(ctx, args, stream);
14+
break;
1515
case GGML_TYPE_Q4_0:
1616
mul_mat_q_case<GGML_TYPE_Q4_0>(ctx, args, stream);
1717
break;
@@ -275,8 +275,8 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
275275

276276
switch (type) {
277277
// TODO: Q1_0 and Q1_0_g128 MMQ implementation exists but is currently disabled due to accuracy issues
278-
// case GGML_TYPE_Q1_0:
279-
// case GGML_TYPE_Q1_0_g128:
278+
case GGML_TYPE_Q1_0:
279+
case GGML_TYPE_Q1_0_g128:
280280
case GGML_TYPE_Q4_0:
281281
case GGML_TYPE_Q4_1:
282282
case GGML_TYPE_Q5_0:
@@ -307,6 +307,10 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
307307
return false;
308308
}
309309

310+
if ((type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q1_0_g128) && !turing_mma_available(cc)) {
311+
return false;
312+
}
313+
310314
if (turing_mma_available(cc)) {
311315
return true;
312316
}

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

Lines changed: 102 additions & 61 deletions
Original file line numberDiff line numberDiff line change
@@ -174,8 +174,6 @@ static constexpr __device__ int get_mmq_y_device() {
174174
// in terms of 32 bit elements that means K % 2 == 1 for dp4a or K % 8 == 4 for mma.
175175
#define MMQ_TILE_NE_K 32
176176

177-
#define MMQ_DP4A_TXS_Q1_0 tile_x_sizes{mmq_y*MMQ_TILE_NE_K + mmq_y, mmq_y*MMQ_TILE_NE_K/QI1_0 + mmq_y/QI1_0, 0}
178-
#define MMQ_DP4A_TXS_Q1_0_g128 tile_x_sizes{mmq_y*MMQ_TILE_NE_K*4 + mmq_y, mmq_y*MMQ_TILE_NE_K*4/QI1_0_g128 + mmq_y/(QI1_0_g128/4), 0}
179177
#define MMQ_DP4A_TXS_Q4_0 tile_x_sizes{mmq_y*MMQ_TILE_NE_K + mmq_y, mmq_y*MMQ_TILE_NE_K/QI4_0 + mmq_y/QI4_0, 0}
180178
#define MMQ_DP4A_TXS_Q4_1 tile_x_sizes{mmq_y*MMQ_TILE_NE_K + mmq_y, mmq_y*MMQ_TILE_NE_K/QI4_1 + mmq_y/QI4_1, 0}
181179
#define MMQ_DP4A_TXS_Q8_0 tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K*2/QI8_0 + mmq_y/(QI8_0/2), 0}
@@ -189,8 +187,6 @@ static constexpr __device__ int get_mmq_y_device() {
189187

190188
static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml_type type, int mmq_y) {
191189
switch (type) {
192-
case GGML_TYPE_Q1_0: return MMQ_DP4A_TXS_Q1_0;
193-
case GGML_TYPE_Q1_0_g128: return MMQ_DP4A_TXS_Q1_0_g128;
194190
case GGML_TYPE_Q4_0: return MMQ_DP4A_TXS_Q4_0;
195191
case GGML_TYPE_Q4_1: return MMQ_DP4A_TXS_Q4_1;
196192
case GGML_TYPE_Q5_0: return MMQ_DP4A_TXS_Q8_0;
@@ -234,7 +230,7 @@ static_assert(MMQ_MMA_TILE_X_K_FP4 == MMQ_MMA_TILE_X_K_Q8_1, "Wrong tile size fo
234230
static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) {
235231
switch (type) {
236232
case GGML_TYPE_Q1_0: return MMQ_MMA_TILE_X_K_Q8_0;
237-
case GGML_TYPE_Q1_0_g128: return MMQ_MMA_TILE_X_K_Q8_0_g128;
233+
case GGML_TYPE_Q1_0_g128: return MMQ_MMA_TILE_X_K_Q8_0;
238234
case GGML_TYPE_Q4_0: return MMQ_MMA_TILE_X_K_Q8_0;
239235
case GGML_TYPE_Q4_1: return MMQ_MMA_TILE_X_K_Q8_1;
240236
case GGML_TYPE_Q5_0: return MMQ_MMA_TILE_X_K_Q8_0;
@@ -309,20 +305,19 @@ static constexpr __device__ int mmq_get_nwarps_device() {
309305

310306
template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_q1_0(
311307
const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
308+
#if !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE))
309+
GGML_UNUSED_VARS(x, x_tile, kbx0, i_max, stride, mmq_y, need_check);
310+
NO_DEVICE_CODE;
311+
#else
312312
constexpr int nwarps = mmq_get_nwarps_device();
313313
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
314314

315-
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
316-
int * x_qs = (int *) x_tile;
317-
float * x_df = (float *) (x_qs + MMQ_TILE_NE_K);
318-
#else
319-
constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q1_0, mmq_y);
320315
int * x_qs = (int *) x_tile;
321-
float * x_df = (float *) (x_qs + txs.qs);
322-
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
323-
324-
constexpr int threads_per_row = MMQ_ITER_K_Q1_0 / (4 * QR1_0);
316+
float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K);
317+
constexpr int blocks_per_iter = MMQ_ITER_K / QK1_0;
318+
constexpr int threads_per_row = blocks_per_iter * QI1_0;
325319
constexpr int nrows = warp_size / threads_per_row;
320+
constexpr int scale_entries_per_row = blocks_per_iter * (QK1_0 / QK8_1);
326321
const int txi = threadIdx.x % threads_per_row;
327322
const int kbx = txi / QI1_0;
328323

@@ -341,7 +336,6 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
341336
// Read all 4 bytes safely to avoid alignment issues
342337
const int qs0 = bxi->qs[0] | (bxi->qs[1] << 8) | (bxi->qs[2] << 16) | (bxi->qs[3] << 24);
343338

344-
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
345339
// For MMA: unpack 1-bit values to signed bytes (-1 or +1)
346340
// Process all 32 bits, 4 at a time
347341
int unpacked_bytes[8];
@@ -358,34 +352,98 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
358352
// Store unpacked values
359353
#pragma unroll
360354
for (int j = 0; j < 8; ++j) {
361-
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + kbx*8 + j] = unpacked_bytes[j];
355+
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + kbx*QI8_0 + j] = unpacked_bytes[j];
362356
}
363-
#else
364-
// For DP4A: store raw bits, will unpack in vec_dot
365-
x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0;
366-
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
367357
}
368358

369-
constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI1_0;
370-
constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row;
371-
const int kbxd = threadIdx.x % blocks_per_tile_x_row;
359+
constexpr int rows_per_warp = warp_size / scale_entries_per_row;
360+
const int kbxd = threadIdx.x % scale_entries_per_row;
372361

373362
#pragma unroll
374363
for (int i0 = 0; i0 < mmq_y; i0 += nwarps * rows_per_warp) {
375-
int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row;
364+
int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / scale_entries_per_row;
376365

377366
if (need_check) {
378367
i = min(i, i_max);
379368
}
380369

381370
const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbxd;
382371

383-
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
384372
x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + kbxd] = bxi->d;
385-
#else
386-
x_df[i*(MMQ_TILE_NE_K/QI1_0) + i/QI1_0 + kbxd] = bxi->d;
373+
}
387374
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
375+
}
376+
377+
template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_q1_0_g128(
378+
const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
379+
#if !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE))
380+
GGML_UNUSED_VARS(x, x_tile, kbx0, i_max, stride, mmq_y, need_check);
381+
NO_DEVICE_CODE;
382+
#else
383+
constexpr int nwarps = mmq_get_nwarps_device();
384+
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
385+
386+
int * x_qs = (int *) x_tile;
387+
float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K);
388+
389+
constexpr int blocks_per_iter = MMQ_ITER_K / QK1_0_g128;
390+
constexpr int threads_per_row = blocks_per_iter * QI1_0_g128;
391+
constexpr int nrows = warp_size / threads_per_row;
392+
constexpr int scale_entries_per_block = QK1_0_g128 / QK8_1;
393+
constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block;
394+
395+
const int txi = threadIdx.x % threads_per_row;
396+
const int kbx = txi / QI1_0_g128;
397+
const int kqsx = txi % QI1_0_g128;
398+
399+
#pragma unroll
400+
for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) {
401+
int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row;
402+
403+
if (need_check) {
404+
i = min(i, i_max);
405+
}
406+
407+
const block_q1_0_g128 * bxi = (const block_q1_0_g128 *) x + kbx0 + i*stride + kbx;
408+
const int qs_offset = 4*kqsx;
409+
const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) |
410+
(bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24);
411+
412+
int unpacked_bytes[8];
413+
#pragma unroll
414+
for (int j = 0; j < 8; ++j) {
415+
const int shift = j * 4;
416+
const int bits4 = (qs0 >> shift) & 0x0F;
417+
const int b0 = (bits4 & 0x01) ? 1 : -1;
418+
const int b1 = (bits4 & 0x02) ? 1 : -1;
419+
const int b2 = (bits4 & 0x04) ? 1 : -1;
420+
const int b3 = (bits4 & 0x08) ? 1 : -1;
421+
unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24);
422+
}
423+
424+
const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0;
425+
#pragma unroll
426+
for (int j = 0; j < 8; ++j) {
427+
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j] = unpacked_bytes[j];
428+
}
388429
}
430+
431+
const int ksx = threadIdx.x % scale_entries_per_row;
432+
const int scale_block = ksx / scale_entries_per_block;
433+
434+
#pragma unroll
435+
for (int i0 = 0; i0 < mmq_y; i0 += nwarps) {
436+
int i = i0 + threadIdx.y;
437+
438+
if (need_check) {
439+
i = min(i, i_max);
440+
}
441+
442+
const block_q1_0_g128 * bxi = (const block_q1_0_g128 *) x + kbx0 + i*stride + scale_block;
443+
444+
x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + ksx] = bxi->d;
445+
}
446+
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
389447
}
390448

391449
template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_q4_0(
@@ -450,38 +508,12 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
450508
}
451509

452510
template <int mmq_x, int mmq_y>
453-
static __device__ __forceinline__ void vec_dot_q1_0_q8_1_dp4a(
511+
static __device__ __forceinline__ void vec_dot_q1_mmq_dp4a_disabled(
454512
const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) {
455-
constexpr int nwarps = mmq_get_nwarps_device();
456-
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
457-
458-
constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q1_0, mmq_y);
459-
const int * x_qs = (const int *) x;
460-
const float * x_df = (const float *) x_qs + txs.qs;
461-
const int * y_qs = (const int *) y + 4;
462-
const float * y_df = (const float *) y;
463-
464-
// #pragma unroll
465-
for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += VDR_Q1_0_Q8_1_MMQ) {
466-
const int k0 = k00 + k01;
467-
468-
#pragma unroll
469-
for (int j0 = 0; j0 < mmq_x; j0 += nwarps) {
470-
const int j = j0 + threadIdx.y;
471-
472-
#pragma unroll
473-
for (int i0 = 0; i0 < mmq_y; i0 += warp_size) {
474-
const int i = i0 + threadIdx.x;
475-
476-
// For MMQ, we have separate float scales, need to combine them into half2 format
477-
const half2 ds8 = __floats2half2_rn(y_df[j*MMQ_TILE_Y_K + (k0/QI8_1) % (MMQ_TILE_NE_K/QI8_1)], 0.0f);
478-
479-
sum[j0/nwarps*mmq_y/warp_size + i0/warp_size] += vec_dot_q1_0_q8_1_impl<VDR_Q1_0_Q8_1_MMQ>
480-
(&x_qs[i*(MMQ_TILE_NE_K + 1) + k0], &y_qs[j*MMQ_TILE_Y_K + k0 % MMQ_TILE_NE_K],
481-
x_df[i*(MMQ_TILE_NE_K/QI1_0) + i/QI1_0 + k0/QI1_0], ds8);
482-
}
483-
}
484-
}
513+
// Q1_0 and Q1_0_g128 intentionally target the MMA path only on this branch.
514+
// If DP4A support is needed later for older GPUs, it should be reintroduced and validated separately.
515+
GGML_UNUSED_VARS(x, y, sum, k00, mmq_x, mmq_y);
516+
NO_DEVICE_CODE;
485517
}
486518

487519
template <int mmq_x, int mmq_y>
@@ -3341,7 +3373,16 @@ struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_Q1_0> {
33413373
static constexpr int vdr = VDR_Q1_0_Q8_1_MMQ;
33423374
static constexpr load_tiles_mmq_t load_tiles = load_tiles_q1_0<mmq_y, need_check>;
33433375
static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, MMQ_Q8_1_DS_LAYOUT_D4>;
3344-
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q1_0_q8_1_dp4a<mmq_x, mmq_y>;
3376+
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q1_mmq_dp4a_disabled<mmq_x, mmq_y>;
3377+
};
3378+
3379+
template <int mmq_x, int mmq_y, bool need_check>
3380+
struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_Q1_0_g128> {
3381+
static constexpr int vdr = VDR_Q1_0_g128_Q8_1_MMQ;
3382+
static constexpr load_tiles_mmq_t load_tiles = load_tiles_q1_0_g128<mmq_y, need_check>;
3383+
static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, MMQ_Q8_1_DS_LAYOUT_D4>;
3384+
// The DP4A path is intentionally disabled; keep the MMA path as the validated route.
3385+
static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q1_mmq_dp4a_disabled<mmq_x, mmq_y>;
33453386
};
33463387

33473388
template <int mmq_x, int mmq_y, bool need_check>
@@ -4199,8 +4240,8 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda
41994240
#define DECL_MMQ_CASE(type) \
42004241
template void mul_mat_q_case<type>(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \
42014242

4202-
// extern DECL_MMQ_CASE(GGML_TYPE_Q1_0); // disabled, cuBLAS fallback used
4203-
// extern DECL_MMQ_CASE(GGML_TYPE_Q1_0_g128); // disabled, cuBLAS fallback used
4243+
extern DECL_MMQ_CASE(GGML_TYPE_Q1_0);
4244+
extern DECL_MMQ_CASE(GGML_TYPE_Q1_0_g128);
42044245
extern DECL_MMQ_CASE(GGML_TYPE_Q4_0);
42054246
extern DECL_MMQ_CASE(GGML_TYPE_Q4_1);
42064247
extern DECL_MMQ_CASE(GGML_TYPE_Q5_0);

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -297,6 +297,7 @@ void quantize_mmq_q8_1_cuda(
297297
const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ);
298298
const dim3 num_blocks(ne1, block_num_y, ne2*ne3);
299299
const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1);
300+
300301
switch (mmq_get_q8_1_ds_layout(type_src0)) {
301302
case MMQ_Q8_1_DS_LAYOUT_D4:
302303
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4>

‎ggml/src/ggml-cuda/template-instances/generate_cu_files.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,7 @@
3232
SOURCE_FATTN_MMA_CASE = "DECL_FATTN_MMA_F16_CASE({head_size_kq}, {head_size_v}, {ncols1}, {ncols2});\n"
3333

3434
TYPES_MMQ = [
35+
"GGML_TYPE_Q1_0", "GGML_TYPE_Q1_0_g128",
3536
"GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0",
3637
"GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K",
3738
"GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S",

‎ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0.cu‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
2+
3+
#include "../mmq.cuh"
4+
5+
DECL_MMQ_CASE(GGML_TYPE_Q1_0);

‎ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0_g128.cu‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
2+
3+
#include "../mmq.cuh"
4+
5+
DECL_MMQ_CASE(GGML_TYPE_Q1_0_g128);

‎src/llama-model-loader.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,8 @@ static std::string llama_model_ftype_name(llama_ftype ftype) {
3131
case LLAMA_FTYPE_ALL_F32: return "all F32";
3232
case LLAMA_FTYPE_MOSTLY_F16: return "F16";
3333
case LLAMA_FTYPE_MOSTLY_BF16: return "BF16";
34-
case LLAMA_FTYPE_MOSTLY_Q1_0: return "Q1_0 - 1.5 bpw";
35-
case LLAMA_FTYPE_MOSTLY_Q1_0_g128: return "Q1_0_g128 - 1.125 bpw";
34+
case LLAMA_FTYPE_MOSTLY_Q1_0: return "Q1_0";
35+
case LLAMA_FTYPE_MOSTLY_Q1_0_g128: return "Q1_0_g128";
3636
case LLAMA_FTYPE_MOSTLY_Q4_0: return "Q4_0";
3737
case LLAMA_FTYPE_MOSTLY_Q4_1: return "Q4_1";
3838
case LLAMA_FTYPE_MOSTLY_Q5_0: return "Q5_0";

0 commit comments

Comments
 (0)