diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 15290c3d1..870f93df0 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -805,6 +805,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta nsg = N_SG_Q1_0; nr0 = N_R0_Q1_0; } break; + case GGML_TYPE_Q2_0: + { + nsg = N_SG_Q2_0; + nr0 = N_R0_Q2_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; @@ -1029,6 +1034,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m nsg = N_SG_Q1_0; nr0 = N_R0_Q1_0; } break; + case GGML_TYPE_Q2_0: + { + nsg = N_SG_Q2_0; + nr0 = N_R0_Q2_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 5d29250f6..1dfe0bdd5 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1289,6 +1289,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_BF16: case GGML_TYPE_Q8_0: case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1316,6 +1317,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te return false; } case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index d6761023b..a45f3ac67 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -24,6 +24,9 @@ #define N_R0_Q1_0 8 #define N_SG_Q1_0 2 +#define N_R0_Q2_0 8 +#define N_SG_Q2_0 2 + #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 45909c477..805bc4093 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2077,6 +2077,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_BF16 || op->src[0]->type == GGML_TYPE_Q1_0 || + op->src[0]->type == GGML_TYPE_Q2_0 || op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || op->src[0]->type == GGML_TYPE_Q5_0 || diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 6b6f9fd87..38919ebad 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -170,6 +170,39 @@ void dequantize_q1_0_t4(device const block_q1_0 * xb, short il, thread type4 & r reg = (type4) reg_f; } +template +void dequantize_q2_0(device const block_q2_0 * xb, short il, thread type4x4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + + const int byte_offset = il * 4; // il*16 elements = il*4 bytes (4 elements per byte) + float4x4 reg_f; + + for (int i = 0; i < 4; i++) { + const uint8_t b = qs[byte_offset + i]; + reg_f[i][0] = ((float)((b >> 0) & 3) - 1.0f) * d; + reg_f[i][1] = ((float)((b >> 2) & 3) - 1.0f) * d; + reg_f[i][2] = ((float)((b >> 4) & 3) - 1.0f) * d; + reg_f[i][3] = ((float)((b >> 6) & 3) - 1.0f) * d; + } + + reg = (type4x4) reg_f; +} + +template +void dequantize_q2_0_t4(device const block_q2_0 * xb, short il, thread type4 & reg) { + const float d = xb->d; + const uint8_t b = xb->qs[il]; + + float4 reg_f; + reg_f[0] = ((float)((b >> 0) & 3) - 1.0f) * d; + reg_f[1] = ((float)((b >> 2) & 3) - 1.0f) * d; + reg_f[2] = ((float)((b >> 4) & 3) - 1.0f) * d; + reg_f[3] = ((float)((b >> 6) & 3) - 1.0f) * d; + + reg = (type4) reg_f; +} + template void dequantize_q4_0(device const block_q4_0 * xb, short il, thread type4x4 & reg) { device const uint16_t * qs = ((device const uint16_t *)xb + 1); @@ -221,6 +254,27 @@ void quantize_q1_0(device const float * src, device block_q1_0 & dst) { } } +void quantize_q2_0(device const float * src, device block_q2_0 & dst) { + float amax = 0.0f; + for (int j = 0; j < QK2_0; j++) { + float a = fabs(src[j]); + if (a > amax) amax = a; + } + const float d = amax; + dst.d = d; + + const float id = d > 0.0f ? 1.0f / d : 0.0f; + + for (int j = 0; j < QK2_0 / 4; j++) { + dst.qs[j] = 0; + } + for (int j = 0; j < QK2_0; j++) { + int q = (int)round(src[j] * id) + 1; + q = max(0, min(3, q)); + dst.qs[j / 4] |= (q << (2 * (j % 4))); + } +} + void quantize_q4_0(device const float * src, device block_q4_0 & dst) { #pragma METAL fp math_mode(safe) float amax = 0.0f; // absolute max @@ -3289,6 +3343,55 @@ inline float block_q_n_dot_y(device const block_q1_0 * qb_curr, float sumy, thre return qb_curr->d * (2.0f * acc - sumy); } +// Q2_0 dot: d * (sum_lo(y) + 2*sum_hi(y) - sumy) via per-bit conditional adds +inline float block_q_n_dot_y(device const block_q2_0 * qb_curr, float sumy, thread float * yl, int il) { + device const uint8_t * qs = qb_curr->qs + (il / 4); + const uint8_t b0 = qs[0]; + const uint8_t b1 = qs[1]; + const uint8_t b2 = qs[2]; + const uint8_t b3 = qs[3]; + + // Accumulate where low bit is set (bits 0,2,4,6 of each byte) + float acc_lo = 0.0f; + acc_lo += select(0.0f, yl[ 0], bool(b0 & 0x01)); + acc_lo += select(0.0f, yl[ 1], bool(b0 & 0x04)); + acc_lo += select(0.0f, yl[ 2], bool(b0 & 0x10)); + acc_lo += select(0.0f, yl[ 3], bool(b0 & 0x40)); + acc_lo += select(0.0f, yl[ 4], bool(b1 & 0x01)); + acc_lo += select(0.0f, yl[ 5], bool(b1 & 0x04)); + acc_lo += select(0.0f, yl[ 6], bool(b1 & 0x10)); + acc_lo += select(0.0f, yl[ 7], bool(b1 & 0x40)); + acc_lo += select(0.0f, yl[ 8], bool(b2 & 0x01)); + acc_lo += select(0.0f, yl[ 9], bool(b2 & 0x04)); + acc_lo += select(0.0f, yl[10], bool(b2 & 0x10)); + acc_lo += select(0.0f, yl[11], bool(b2 & 0x40)); + acc_lo += select(0.0f, yl[12], bool(b3 & 0x01)); + acc_lo += select(0.0f, yl[13], bool(b3 & 0x04)); + acc_lo += select(0.0f, yl[14], bool(b3 & 0x10)); + acc_lo += select(0.0f, yl[15], bool(b3 & 0x40)); + + // Accumulate where high bit is set (bits 1,3,5,7 of each byte) + float acc_hi = 0.0f; + acc_hi += select(0.0f, yl[ 0], bool(b0 & 0x02)); + acc_hi += select(0.0f, yl[ 1], bool(b0 & 0x08)); + acc_hi += select(0.0f, yl[ 2], bool(b0 & 0x20)); + acc_hi += select(0.0f, yl[ 3], bool(b0 & 0x80)); + acc_hi += select(0.0f, yl[ 4], bool(b1 & 0x02)); + acc_hi += select(0.0f, yl[ 5], bool(b1 & 0x08)); + acc_hi += select(0.0f, yl[ 6], bool(b1 & 0x20)); + acc_hi += select(0.0f, yl[ 7], bool(b1 & 0x80)); + acc_hi += select(0.0f, yl[ 8], bool(b2 & 0x02)); + acc_hi += select(0.0f, yl[ 9], bool(b2 & 0x08)); + acc_hi += select(0.0f, yl[10], bool(b2 & 0x20)); + acc_hi += select(0.0f, yl[11], bool(b2 & 0x80)); + acc_hi += select(0.0f, yl[12], bool(b3 & 0x02)); + acc_hi += select(0.0f, yl[13], bool(b3 & 0x08)); + acc_hi += select(0.0f, yl[14], bool(b3 & 0x20)); + acc_hi += select(0.0f, yl[15], bool(b3 & 0x80)); + + return qb_curr->d * (acc_lo + 2.0f * acc_hi - sumy); +} + // function for calculate inner product between half a q4_0 block and 16 floats (yl), sumy is SUM(yl[i]) // il indicates where the q4 quants begin (0 or QK4_0/4) // we assume that the yl's have been multiplied with the appropriate scale factor @@ -3592,6 +3695,86 @@ kernel void kernel_mul_mv_q1_0_f32( kernel_mul_mv_q1_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +template +void kernel_mul_mv_q2_0_f32_impl( + args_t args, + device const char * src0, + device const char * src1, + device char * dst, + threadgroup char * shmem, + uint3 tgpig, + ushort tiisg, + ushort sgitg) { + const short NSG = FC_mul_mv_nsg; + + const int nb = args.ne00/QK2_0; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * NSG + sgitg) * nr0; + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + device const block_q2_0 * ax[nr0]; + for (int row = 0; row < nr0; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03; + ax[row] = (device const block_q2_0 *) ((device char *) src0 + offset0); + } + + float yl[16]; + float sumf[nr0] = {0.f}; + + // group 64: 4 sub-blocks of 16 weights per Q2_0 block + const short ix = (tiisg/4); + const short il = (tiisg%4)*16; + + device const float * yb = y + ix*QK2_0 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/4) { + float sumy = 0.f; + + FOR_UNROLL (short i = 0; i < 16; i++) { + yl[i] = yb[i]; + sumy += yb[i]; + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + sumf[row] += block_q_n_dot_y(ax[row] + ib, sumy, yl, il); + } + + yb += QK2_0 * (N_SIMDWIDTH/4); + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + +[[host_name("kernel_mul_mv_q2_0_f32")]] +kernel void kernel_mul_mv_q2_0_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + kernel_mul_mv_q2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + kernel void kernel_mul_mv_q4_0_f32( constant ggml_metal_kargs_mul_mv & args, device const char * src0, @@ -3989,6 +4172,11 @@ template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_3")]] kernel mul_mv_ext_q4 template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q1_0, 128, dequantize_q1_0_t4>; template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_q1_0, 128, dequantize_q1_0_t4>; +template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q2_0, 64, dequantize_q2_0_t4>; +template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_q2_0, 64, dequantize_q2_0_t4>; +template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q2_0, 64, dequantize_q2_0_t4>; +template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_q2_0, 64, dequantize_q2_0_t4>; + template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q4_0, 32, dequantize_q4_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_q4_0, 32, dequantize_q4_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q4_0, 32, dequantize_q4_0_t4>; @@ -7700,6 +7888,7 @@ typedef decltype(kernel_cpy_f32_q) cpy_f_q_ template [[host_name("kernel_cpy_f32_q8_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_q1_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; +template [[host_name("kernel_cpy_f32_q2_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_q4_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_q4_1")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_q5_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; @@ -7745,6 +7934,7 @@ kernel void kernel_cpy_q_f32( typedef decltype(kernel_cpy_q_f32) cpy_q_f_t; template [[host_name("kernel_cpy_q1_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; +template [[host_name("kernel_cpy_q2_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q4_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q4_1_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q5_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; @@ -7752,6 +7942,7 @@ template [[host_name("kernel_cpy_q5_1_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32< template [[host_name("kernel_cpy_q8_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q1_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; +template [[host_name("kernel_cpy_q2_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q4_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q4_1_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q5_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; @@ -9596,6 +9787,7 @@ template [[host_name("kernel_get_rows_bf16")]] kernel get_rows_f_t kernel_get_ro typedef decltype(kernel_get_rows_q) get_rows_q_t; template [[host_name("kernel_get_rows_q1_0")]] kernel get_rows_q_t kernel_get_rows_q; +template [[host_name("kernel_get_rows_q2_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_1")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q5_0")]] kernel get_rows_q_t kernel_get_rows_q; @@ -10466,6 +10658,7 @@ template [[host_name("kernel_mul_mm_f16_f32")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_mm; #endif template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_q2_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_mm; @@ -10490,6 +10683,7 @@ template [[host_name("kernel_mul_mm_iq4_xs_f32")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_f32_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_f16_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q1_0_f16")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_q2_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f16")]] kernel mul_mm_t kernel_mul_mm; @@ -10523,6 +10717,7 @@ template [[host_name("kernel_mul_mm_id_f16_f32")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_bf16_f32")]] kernel mul_mm_id kernel_mul_mm_id; #endif template [[host_name("kernel_mul_mm_id_q1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_q2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; @@ -10547,6 +10742,7 @@ template [[host_name("kernel_mul_mm_id_iq4_xs_f32")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_f32_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_f16_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_q2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; @@ -10702,6 +10898,7 @@ template [[host_name("kernel_mul_mv_id_bf16_f32_4")]] kernel kernel_mul_mv_id_4 template [[host_name("kernel_mul_mv_id_q8_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_q1_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; +template [[host_name("kernel_mul_mv_id_q2_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_q4_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_q4_1_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_q5_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>;