From 782cc392128de9064db3867188b0225314a31f03 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Tue, 1 Sep 2026 11:41:15 -0700 Subject: [PATCH 1/5] metal: multi-column Q1_0 mul_mv for small verify batches Speculative-decode verify batches put N = k+1 rows (2..8) through the target. For Q1_0 that hit the generic small-batch mul_mv_ext path, which dequantizes every 1-bit weight through the float path and runs at a flat ~2.5 TFLOPS for every type: a 4-row forward on the 27B Q1_0 target cost 3.5 decode steps on M5 Pro (pp4 44.7 tok/s vs tg128 39.1). Add kernel_mul_mv_q1_0_f32_nr1_{2,3,4}: each q1_0 block is read once and dotted against nr1 src1 columns (tpb threads per block, SW = QK1_0/tpb). Route Q1_0 off the ext path, pick nr1 = 3 for exactly 3 columns and 2 otherwise, and keep Q1_0 on mul_mv up to 16 rows before mul_mm. The single-column kernel is untouched, so decode numerics are unchanged. M5 Pro, 27B Q1_0, -fa 1 (tok/s, 3-run medians, interleaved A/B): pp2 29.6 -> 54.2 pp4 44.7 -> 69.2 pp8 48.6 -> 77.4 pp12 49.1 -> 79.6 pp16 65.3 -> 80.4 tg128 unchanged (38.8 vs 37.7, thermal spread) test-backend-ops MUL_MAT q1_0: 45/45 vs CPU. Env knobs for A/B: GGML_METAL_Q1_0_NR1 (1 disables, 2..4 forces), GGML_METAL_Q1_0_EXT_ENABLE, GGML_METAL_Q1_0_MV_MAX (default 16). --- ggml/src/ggml-metal/ggml-metal-device.cpp | 13 ++ ggml/src/ggml-metal/ggml-metal-ops.cpp | 11 +- ggml/src/ggml-metal/kernels/mul_mv.metal | 147 ++++++++++++++++++++++ 3 files changed, 169 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 4870bdcb968d..a7b2a223f6f3 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1,3 +1,4 @@ +#include #include "ggml-metal-device.h" #include "ggml-metal-impl.h" @@ -860,6 +861,18 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta { nsg = N_SG_Q1_0; nr0 = N_R0_Q1_0; + // multi-column variants stream each weight block once per nr1 src1 + // columns (spec-decode verify batches). GGML_METAL_Q1_0_NR1=1 disables, + // 2..4 forces a variant. + static const int nr1_env = getenv("GGML_METAL_Q1_0_NR1") ? atoi(getenv("GGML_METAL_Q1_0_NR1")) : 0; + if (nr1_env >= 2 && nr1_env <= 4) { + nr1 = nr1_env; + } else if (nr1_env != 1 && ne11 == 3) { + nr1 = 3; + } else if (nr1_env != 1 && ne11 >= 2) { + nr1 = 2; + } + suffix = nr1 == 2 ? "_nr1_2" : nr1 == 3 ? "_nr1_3" : nr1 == 4 ? "_nr1_4" : ""; } break; case GGML_TYPE_Q2_0: { diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 6be0f52015b2..8d6fbe28f3bf 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -1,3 +1,4 @@ +#include #include "ggml-metal-ops.h" #include "ggml.h" @@ -2551,7 +2552,13 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { // find the break-even point where the matrix-matrix kernel becomes more efficient compared // to the matrix-vector kernel - const int ne11_mm_min = 8; + // Q1_0: the generic small-batch mul_mv_ext path is ALU-bound (flat ~2.5 TFLOPS for + // 2..8 columns, 2.5-4x slower per weight pass than the Q1_0 mul_mv kernel), so Q1_0 + // stays on the (multi-column) mul_mv kernels up to GGML_METAL_Q1_0_MV_MAX rows. + static const bool q1_0_ext_enable = getenv("GGML_METAL_Q1_0_EXT_ENABLE") != nullptr; + static const int q1_0_mv_max = getenv("GGML_METAL_Q1_0_MV_MAX") ? atoi(getenv("GGML_METAL_Q1_0_MV_MAX")) : 16; + + const int ne11_mm_min = op->src[0]->type == GGML_TYPE_Q1_0 ? std::max(8, q1_0_mv_max) : 8; // first try to use small-batch mat-mv kernels // these should be efficient for BS [2, ~8] @@ -2562,7 +2569,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { op->src[0]->type == GGML_TYPE_F32 || // TODO: helper function 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_Q1_0 && q1_0_ext_enable) || op->src[0]->type == GGML_TYPE_Q2_0 || op->src[0]->type == GGML_TYPE_PQ2_0 || op->src[0]->type == GGML_TYPE_Q4_0 || diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index bd4482ecf205..39d0b21b94ea 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -430,6 +430,153 @@ kernel void kernel_mul_mv_q1_0_f32( kernel_mul_mv_q1_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +// Q1_0 multi-column mat-vec: each q1_0 block is read once and dotted against nr1 +// src1 columns (the spec-decode verify batch, ne11 = k+1 rows). tpb threads +// cooperate on one block, slice width SW = QK1_0/tpb. Register budget (M5 Pro): +// keep nr1*SW <= 32 staged y values per thread. +template +static inline float q1_0_dot_y_sw(thread const uint8_t * b, const float d, const float sumy, thread const float * yl) { + float acc = 0.0f; + + FOR_UNROLL (short i = 0; i < SW; i++) { + acc += select(0.0f, yl[i], bool(b[i/8] & (1u << (i%8)))); + } + + return d * (2.0f * acc - sumy); +} + +template +void kernel_mul_mv_q1_0_f32_nc_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 short SW = QK1_0/tpb; + + const int nb = args.ne00/QK1_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 int c0 = r1 * nr1; + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + device const block_q1_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_q1_0 *) ((device char *) src0 + offset0); + } + + float yl[nr1][SW]; + float sumy[nr1]; + float sumf[nr0][nr1]; + FOR_UNROLL (short row = 0; row < nr0; row++) { + FOR_UNROLL (short c = 0; c < nr1; c++) { + sumf[row][c] = 0.f; + } + } + + const short ix = (tiisg/tpb); // block in flight + const short il = (tiisg%tpb)*SW; // element offset within the block + + device const float * yb[nr1]; + FOR_UNROLL (short c = 0; c < nr1; c++) { + // tail columns are clamped (computed but not stored) + const int ic = min(c0 + c, args.ne11 - 1); + const uint64_t offset1 = (uint64_t)ic*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + yb[c] = (device const float *) (src1 + offset1) + ix*QK1_0 + il; + } + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/tpb) { + FOR_UNROLL (short c = 0; c < nr1; c++) { + sumy[c] = 0.f; + FOR_UNROLL (short i = 0; i < SW; i++) { + yl[c][i] = yb[c][i]; + sumy[c] += yb[c][i]; + } + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + device const block_q1_0 * qb = ax[row] + ib; + device const uint8_t * qs = qb->qs + il/8; + + uint8_t b[SW/8]; + FOR_UNROLL (short i = 0; i < SW/8; i++) { + b[i] = qs[i]; + } + const float d = qb->d; + + FOR_UNROLL (short c = 0; c < nr1; c++) { + sumf[row][c] += q1_0_dot_y_sw(b, d, sumy[c], yl[c]); + } + } + + FOR_UNROLL (short c = 0; c < nr1; c++) { + yb[c] += QK1_0 * (N_SIMDWIDTH/tpb); + } + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1; + + for (short c = 0; c < nr1; c++) { + if (c0 + c >= args.ne11) { + break; + } + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row][c]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[(uint64_t)(c0 + c)*args.ne0 + first_row + row] = tot; + } + } + } +} + +[[host_name("kernel_mul_mv_q1_0_f32_nr1_2")]] +kernel void kernel_mul_mv_q1_0_f32_nr1_2( + 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_q1_0_f32_nc_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + +[[host_name("kernel_mul_mv_q1_0_f32_nr1_3")]] +kernel void kernel_mul_mv_q1_0_f32_nr1_3( + 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_q1_0_f32_nc_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + +[[host_name("kernel_mul_mv_q1_0_f32_nr1_4")]] +kernel void kernel_mul_mv_q1_0_f32_nr1_4( + 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_q1_0_f32_nc_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + template void kernel_mul_mv_q2_0_f32_impl( args_t args, From c560a5e4640814ae8a23fb80d9e37389501d6cd3 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Tue, 1 Sep 2026 12:06:15 -0700 Subject: [PATCH 2/5] metal: experimental Q1_0 small-batch verify kernels (half2, word-parallel popcount) Both opt-in and default-off, kept for the measurements they produced. The multi-column kernel in the parent commit is still ALU-bound: a timing probe that removes only the per-weight bit test runs pp4 69.2 -> 118.9 tok/s, so the bit extraction, not the weight read, is what is left to attack. 1. GGML_METAL_Q1_0_H2: two src1 columns per half2 lane, fp32 block sums, fp32 promotion every q1_0 block. Correct (45/45 MUL_MAT q1_0 vs CPU) and worth pp2 55.7 -> 57.7, pp4 69.2 -> 72.1, pp8 78.0 -> 81.1 = +3.5 to +4.5%. Too small to justify the fp16 accumulator, so it stays off. 2. GGML_METAL_Q1_0_POPCNT: quantize each activation column to int8 once per matmul and store it as 8 bit-planes, then consume 32 weights per AND+popcount (~0.8 integer ops per weight per column against ~3 for select-per-weight). Adds kernel_q1_0_build_planes, a scratch record in the dst buffer tail sized by ggml_metal_op_mul_mat_extra_q1_0_planes, and a barrier between the two dispatches. Correct (45/45) and end-to-end accept counters are unchanged (80.645%, identical to the fp32 path). Timing is pending an uncontended GPU. Note for whoever touches Metal type gating next: this backend's supports_op accepts every type for MUL_MAT and only rejects NVFP4, so a type with no good kernel silently takes a generic path. That is the same shape as the TQ1_0 and iq4_nl gaps found the same day; those abort, Q1_0 small batch merely ran slow. --- ggml/src/ggml-metal/ggml-metal-device.cpp | 62 +++- ggml/src/ggml-metal/ggml-metal-device.h | 2 + ggml/src/ggml-metal/ggml-metal-impl.h | 11 + ggml/src/ggml-metal/ggml-metal-ops.cpp | 104 +++++++ ggml/src/ggml-metal/ggml-metal-ops.h | 1 + ggml/src/ggml-metal/ggml-metal.cpp | 4 + ggml/src/ggml-metal/kernels/mul_mv.metal | 334 ++++++++++++++++++++++ 7 files changed, 517 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index a7b2a223f6f3..77bb09213450 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -411,6 +411,59 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_sum_rows(ggml_me return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_q1_0_planes(ggml_metal_library_t lib) { + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_q1_0_build_planes"); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_q1_0_pc(ggml_metal_library_t lib, const ggml_tensor * op, int nr1) { + GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne); + GGML_TENSOR_LOCALS( int32_t, ne1, op->src[1], ne); + + char base[256]; + char name[256]; + + const int nsg = N_SG_Q1_0; + const int nr0 = N_R0_Q1_0_PC; + + const int16_t r2 = (int16_t) (ne12 / ne02); + const int16_t r3 = (int16_t) (ne13 / ne03); + + snprintf(base, 256, "kernel_mul_mv_q1_0_f32_pc_nr1_%d", nr1); + snprintf(name, 256, "%s_nsg=%d_ne12=%d_r2=%d_r3=%d", base, nsg, ne12, r2, r3); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + ggml_metal_cv_t cv = ggml_metal_cv_init(); + + ggml_metal_cv_set_int16(cv, nsg, FC_MUL_MV + 0); + ggml_metal_cv_set_int16(cv, (int16_t) ne12, FC_MUL_MV + 2); + ggml_metal_cv_set_int16(cv, r2, FC_MUL_MV + 3); + ggml_metal_cv_set_int16(cv, r3, FC_MUL_MV + 4); + + res = ggml_metal_library_compile_pipeline(lib, base, name, cv); + + ggml_metal_cv_free(cv); + } + + res.nr0 = nr0; + res.nr1 = nr1; + res.nsg = nsg; + res.smem = 0; + + return res; +} + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_cumsum_blk(ggml_metal_library_t lib, const ggml_tensor * op) { GGML_ASSERT(op->op == GGML_OP_CUMSUM); @@ -872,7 +925,14 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta } else if (nr1_env != 1 && ne11 >= 2) { nr1 = 2; } - suffix = nr1 == 2 ? "_nr1_2" : nr1 == 3 ? "_nr1_3" : nr1 == 4 ? "_nr1_4" : ""; + // half2 variants pack two src1 columns per lane (opt-in, see the kernel comment) + static const bool h2 = getenv("GGML_METAL_Q1_0_H2") != nullptr; + if (h2 && nr1 >= 2) { + nr1 = nr1 == 2 ? 2 : 4; + suffix = nr1 == 2 ? "_h2_nr1_2" : "_h2_nr1_4"; + } else { + suffix = nr1 == 2 ? "_nr1_2" : nr1 == 3 ? "_nr1_3" : nr1 == 4 ? "_nr1_4" : ""; + } } break; case GGML_TYPE_Q2_0: { diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 46708c020288..d1c615633a45 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -123,6 +123,8 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_glu struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_sum (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_sum_rows (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_cumsum_blk (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_q1_0_planes (ggml_metal_library_t lib); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_q1_0_pc (ggml_metal_library_t lib, const struct ggml_tensor * op, int nr1); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_cumsum_add (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_tri (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_soft_max (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 836cbbcb5c9e..17f9a25b17c5 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -24,6 +24,11 @@ #define N_R0_Q1_0 8 #define N_SG_Q1_0 2 +// Q1_0 word-parallel (popcount) verify path: rows per simdgroup, and the uint32 +// stride of one activation bit-plane record (8 planes x 4 words + scale + sum + pad). +#define N_R0_Q1_0_PC 4 +#define Q1_0_PLANE_STRIDE 36 + #define N_R0_Q2_0 8 #define N_SG_Q2_0 2 @@ -503,6 +508,12 @@ typedef struct { int16_t r3; } ggml_metal_kargs_mul_mv; +typedef struct { + int32_t nblk; + int32_t ne11; + uint64_t nb11; +} ggml_metal_kargs_q1_0_planes; + typedef struct { int32_t ne00; int32_t ne01; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 8d6fbe28f3bf..4538e15a2653 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2560,6 +2560,90 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { const int ne11_mm_min = op->src[0]->type == GGML_TYPE_Q1_0 ? std::max(8, q1_0_mv_max) : 8; + // Q1_0 word-parallel verify path (opt-in): quantize the activation columns into + // int8 bit-planes once, then consume 32 weights per AND+popcount instead of one + // select per weight. See the kernel comment in mul_mv.metal. + static const bool q1_0_pc = getenv("GGML_METAL_Q1_0_POPCNT") != nullptr; + + if (q1_0_pc && + op->src[0]->type == GGML_TYPE_Q1_0 && + op->src[1]->type == GGML_TYPE_F32 && + ne11 >= 2 && ne11 <= 16 && + ne00 % 128 == 0 && + ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1 && + ggml_metal_op_mul_mat_extra_q1_0_planes(op) > 0) { + const int32_t nblk = ne00/128; + + ggml_metal_buffer_id bid_src0 = ggml_metal_get_buffer_id(op->src[0]); + ggml_metal_buffer_id bid_src1 = ggml_metal_get_buffer_id(op->src[1]); + ggml_metal_buffer_id bid_dst = ggml_metal_get_buffer_id(op); + ggml_metal_buffer_id bid_planes = bid_dst; + + bid_planes.offs += ggml_nbytes(op); + + { + auto pipeline = ggml_metal_library_get_pipeline_q1_0_planes(lib); + + ggml_metal_kargs_q1_0_planes args = { + /*.nblk =*/ nblk, + /*.ne11 =*/ (int32_t) ne11, + /*.nb11 =*/ nb11, + }; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, bid_src1, 1); + ggml_metal_encoder_set_buffer (enc, bid_planes, 2); + + ggml_metal_encoder_dispatch_threadgroups(enc, nblk, ne11, 1, 32, 1, 1); + } + + // the matmul reads what the plane pass just wrote + ggml_metal_op_concurrency_reset(ctx); + + { + const int nr1 = ne11 == 2 ? 2 : 4; + + auto pipeline = ggml_metal_library_get_pipeline_mul_mv_q1_0_pc(lib, op, nr1); + + ggml_metal_kargs_mul_mv args = { + /*.ne00 =*/ ne00, + /*.ne01 =*/ ne01, + /*.ne02 =*/ ne02, + /*.nb00 =*/ nb00, + /*.nb01 =*/ nb01, + /*.nb02 =*/ nb02, + /*.nb03 =*/ nb03, + /*.ne10 =*/ ne10, + /*.ne11 =*/ ne11, + /*.ne12 =*/ ne12, + /*.nb10 =*/ nb10, + /*.nb11 =*/ nb11, + /*.nb12 =*/ nb12, + /*.nb13 =*/ nb13, + /*.ne0 =*/ ne0, + /*.ne1 =*/ ne1, + /*.nr0 =*/ pipeline.nr0, + /*.r2 =*/ r2, + /*.r3 =*/ r3, + }; + + const int nr0 = pipeline.nr0; + const int nsg = pipeline.nsg; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, bid_src0, 1); + ggml_metal_encoder_set_buffer (enc, bid_src1, 2); + ggml_metal_encoder_set_buffer (enc, bid_dst, 3); + ggml_metal_encoder_set_buffer (enc, bid_planes, 4); + + ggml_metal_encoder_dispatch_threadgroups(enc, (ne01 + nr0*nsg - 1)/(nr0*nsg), (ne11 + nr1 - 1)/nr1, 1, 32, nsg, 1); + } + + return 1; + } + // first try to use small-batch mat-mv kernels // these should be efficient for BS [2, ~8] if (op->src[1]->type == GGML_TYPE_F32 && (ne00%128 == 0) && @@ -2766,6 +2850,26 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { return 1; } +// scratch for the Q1_0 word-parallel verify path: int8 activation bit-planes, one +// record per (column, 128-element block). Only the small-batch shapes that path +// serves are padded, so prefill dst buffers are unaffected. +size_t ggml_metal_op_mul_mat_extra_q1_0_planes(const ggml_tensor * op) { + assert(op->op == GGML_OP_MUL_MAT); + + if (!op->src[0] || !op->src[1] || op->src[0]->type != GGML_TYPE_Q1_0) { + return 0; + } + + const int64_t ne00 = op->src[0]->ne[0]; + const int64_t ne11 = op->src[1]->ne[1]; + + if (ne11 < 2 || ne11 > 16 || ne00 % 128 != 0) { + return 0; + } + + return (size_t) (ne00/128) * ne11 * Q1_0_PLANE_STRIDE * sizeof(uint32_t); +} + size_t ggml_metal_op_mul_mat_id_extra_tpe(const ggml_tensor * op) { assert(op->op == GGML_OP_MUL_MAT_ID); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 159a628d04a7..21a064091d15 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -31,6 +31,7 @@ int ggml_metal_op_encode(ggml_metal_op_t ctx, int idx); // // tokens per expert +size_t ggml_metal_op_mul_mat_extra_q1_0_planes(const struct ggml_tensor * op); size_t ggml_metal_op_mul_mat_id_extra_tpe(const struct ggml_tensor * op); // id map [n_tokens, n_expert] diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp index e5b2ee8a55b1..a2fdde0b3ae8 100644 --- a/ggml/src/ggml-metal/ggml-metal.cpp +++ b/ggml/src/ggml-metal/ggml-metal.cpp @@ -216,6 +216,10 @@ static size_t ggml_backend_metal_buffer_type_get_alloc_size(ggml_backend_buffer_ // some operations require additional memory for fleeting data: switch (tensor->op) { + case GGML_OP_MUL_MAT: + { + res += ggml_metal_op_mul_mat_extra_q1_0_planes(tensor); + } break; case GGML_OP_MUL_MAT_ID: { res += ggml_metal_op_mul_mat_id_extra_tpe(tensor); diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index 39d0b21b94ea..284b2ecc3f2b 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -577,6 +577,340 @@ kernel void kernel_mul_mv_q1_0_f32_nr1_4( kernel_mul_mv_q1_0_f32_nc_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +// Q1_0 multi-column mat-vec, half2 form. Two src1 columns share one half2 lane, so +// a single select+add covers both and the ALU slots per weight per column halve -- +// the multi-column kernel above is ALU-bound (2 fp32 lane-ops per weight per column), +// not weight-read-bound, at the 2..8 rows a speculative verify batch uses. +// Only the per-row accumulator is half: the block sum of y stays fp32 (it is computed +// once per block for all nr0 rows, so its cost is amortized) and the accumulator is +// promoted to fp32 at every q1_0 block, bounding the rounding to at most SW halves. +// Opt-in via GGML_METAL_Q1_0_H2=1; gate on measured logit-KLD before enabling by default. +template +void kernel_mul_mv_q1_0_f32_h2_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 short SW = QK1_0/tpb; + + const int nb = args.ne00/QK1_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 int c0 = r1 * (2*npair); + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + device const block_q1_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_q1_0 *) ((device char *) src0 + offset0); + } + + half2 yh[npair][SW]; + float sumy[2*npair]; + float sumf[nr0][2*npair]; + FOR_UNROLL (short row = 0; row < nr0; row++) { + FOR_UNROLL (short c = 0; c < 2*npair; c++) { + sumf[row][c] = 0.f; + } + } + + const short ix = (tiisg/tpb); // block in flight + const short il = (tiisg%tpb)*SW; // element offset within the block + + device const float * yb[2*npair]; + FOR_UNROLL (short c = 0; c < 2*npair; c++) { + // tail columns are clamped (computed but not stored) + const int ic = min(c0 + c, args.ne11 - 1); + const uint64_t offset1 = (uint64_t)ic*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + yb[c] = (device const float *) (src1 + offset1) + ix*QK1_0 + il; + } + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/tpb) { + FOR_UNROLL (short p = 0; p < npair; p++) { + float s0 = 0.f; + float s1 = 0.f; + FOR_UNROLL (short i = 0; i < SW; i++) { + const float y0 = yb[2*p ][i]; + const float y1 = yb[2*p + 1][i]; + yh[p][i] = half2((half) y0, (half) y1); + s0 += y0; + s1 += y1; + } + sumy[2*p ] = s0; + sumy[2*p + 1] = s1; + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + device const block_q1_0 * qb = ax[row] + ib; + device const uint8_t * qs = qb->qs + il/8; + + uint8_t b[SW/8]; + FOR_UNROLL (short i = 0; i < SW/8; i++) { + b[i] = qs[i]; + } + const float d = qb->d; + + FOR_UNROLL (short p = 0; p < npair; p++) { + half2 acc = half2(0.0h); + + FOR_UNROLL (short i = 0; i < SW; i++) { + acc += bool(b[i/8] & (1u << (i%8))) ? yh[p][i] : half2(0.0h); + } + + sumf[row][2*p ] += d*(2.0f*(float) acc.x - sumy[2*p ]); + sumf[row][2*p + 1] += d*(2.0f*(float) acc.y - sumy[2*p + 1]); + } + } + + FOR_UNROLL (short c = 0; c < 2*npair; c++) { + yb[c] += QK1_0 * (N_SIMDWIDTH/tpb); + } + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1; + + for (short c = 0; c < 2*npair; c++) { + if (c0 + c >= args.ne11) { + break; + } + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row][c]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[(uint64_t)(c0 + c)*args.ne0 + first_row + row] = tot; + } + } + } +} + +[[host_name("kernel_mul_mv_q1_0_f32_h2_nr1_2")]] +kernel void kernel_mul_mv_q1_0_f32_h2_nr1_2( + 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_q1_0_f32_h2_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + +[[host_name("kernel_mul_mv_q1_0_f32_h2_nr1_4")]] +kernel void kernel_mul_mv_q1_0_f32_h2_nr1_4( + 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_q1_0_f32_h2_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + +// --------------------------------------------------------------------------- +// Q1_0 word-parallel (popcount) mat-vec for small verify batches. +// +// The select-per-weight kernels above are ALU-bound: a timing probe with the bit +// test removed runs 1.7x faster at 4 rows, so the per-weight bit extraction, not +// the weight read, is what costs. This path removes it. Activations are quantized +// to int8 once per matmul (kernel_q1_0_build_planes) and stored as 8 bit-planes, +// so 32 weights are consumed per AND+popcount pair: +// +// sum_{bit=1} q = sum_b 2^b * popcount(w & plane_b) - 128*popcount(w) +// dot = d * sy * (2*sum_{bit=1} q - sum q) +// +// Costs ~0.8 integer ops per weight per column against ~3 for select-per-weight. +// The int8 activation step is the same numeric class as the CUDA MMVQ path but is +// NOT free: gate on measured logit-KLD before enabling by default. +// Opt-in via GGML_METAL_Q1_0_POPCNT=1. +// --------------------------------------------------------------------------- + +// 32 threads: thread t builds word (t%4) of bit-plane (t/4) for one 128-element block. +kernel void kernel_q1_0_build_planes( + constant ggml_metal_kargs_q1_0_planes & args, + device const char * src1, + device uint * planes, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]]) { + const int ib = tgpig.x; + const int c = tgpig.y; + + device const float * y = (device const float *) (src1 + (uint64_t)c*args.nb11) + (uint64_t)ib*QK1_0; + + float amax = 0.0f; + for (short i = tiisg; i < QK1_0; i += 32) { + amax = max(amax, fabs(y[i])); + } + amax = simd_max(amax); + + const float sy = amax > 0.0f ? amax/127.0f : 1.0f; + const float isy = amax > 0.0f ? 127.0f/amax : 0.0f; + + int sq = 0; + for (short i = tiisg; i < QK1_0; i += 32) { + sq += clamp((int) rint(y[i]*isy), -127, 127); + } + sq = simd_sum(sq); + + const short b = tiisg / 4; // bit plane + const short j = tiisg % 4; // 32-element word + + uint word = 0; + FOR_UNROLL (short i = 0; i < 32; i++) { + const int q = clamp((int) rint(y[32*j + i]*isy), -127, 127); + const uint u = (uint) (q + 128); + word |= ((u >> b) & 1u) << i; + } + + device uint * p = planes + (uint64_t)(c*args.nblk + ib)*Q1_0_PLANE_STRIDE; + + p[b*4 + j] = word; + + if (tiisg == 0) { + p[32] = as_type(sy); + p[33] = as_type((float) sq); + } +} + +template +void kernel_mul_mv_q1_0_f32_pc_impl( + args_t args, + device const char * src0, + device const char * src1, + device char * dst, + device const uint * planes, + uint3 tgpig, + ushort tiisg, + ushort sgitg) { + const short NSG = FC_mul_mv_nsg; + + const int nb = args.ne00/QK1_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 int c0 = r1 * nr1; + + device const block_q1_0 * ax[nr0]; + for (int row = 0; row < nr0; ++row) { + ax[row] = (device const block_q1_0 *) ((device char *) src0 + (first_row + row)*args.nb01); + } + + float sumf[nr0][nr1]; + FOR_UNROLL (short row = 0; row < nr0; row++) { + FOR_UNROLL (short c = 0; c < nr1; c++) { + sumf[row][c] = 0.f; + } + } + + for (int ib = tiisg; ib < nb; ib += N_SIMDWIDTH) { + uint w [nr0][4]; + float d [nr0]; + int pcw[nr0]; + + FOR_UNROLL (short row = 0; row < nr0; row++) { + device const block_q1_0 * qb = ax[row] + ib; + device const uint8_t * qs = qb->qs; // 2-byte aligned: assemble from bytes + + int pc = 0; + FOR_UNROLL (short j = 0; j < 4; j++) { + const uint wj = (uint) qs[4*j + 0] | ((uint) qs[4*j + 1] << 8) | + (((uint) qs[4*j + 2]) << 16) | ((uint) qs[4*j + 3] << 24); + w[row][j] = wj; + pc += popcount(wj); + } + + d [row] = qb->d; + pcw[row] = pc; + } + + FOR_UNROLL (short c = 0; c < nr1; c++) { + // tail columns are clamped (computed but not stored) + const int ic = min(c0 + c, args.ne11 - 1); + + device const uint * p = planes + (uint64_t)(ic*nb + ib)*Q1_0_PLANE_STRIDE; + + uint pl[32]; + FOR_UNROLL (short t = 0; t < 32; t++) { + pl[t] = p[t]; + } + + const float sy = as_type(p[32]); + const float sq = as_type(p[33]); + + FOR_UNROLL (short row = 0; row < nr0; row++) { + uint tot = 0; + + FOR_UNROLL (short b = 0; b < 8; b++) { + const uint cnt = popcount(w[row][0] & pl[b*4 + 0]) + + popcount(w[row][1] & pl[b*4 + 1]) + + popcount(w[row][2] & pl[b*4 + 2]) + + popcount(w[row][3] & pl[b*4 + 3]); + tot += cnt << b; + } + + const int sbit = (int) tot - 128*pcw[row]; + + sumf[row][c] += d[row]*sy*(2.0f*(float) sbit - sq); + } + } + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1; + + for (short c = 0; c < nr1; c++) { + if (c0 + c >= args.ne11) { + break; + } + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row][c]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[(uint64_t)(c0 + c)*args.ne0 + first_row + row] = tot; + } + } + } +} + +[[host_name("kernel_mul_mv_q1_0_f32_pc_nr1_2")]] +kernel void kernel_mul_mv_q1_0_f32_pc_nr1_2( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + device const uint * planes, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + kernel_mul_mv_q1_0_f32_pc_impl(args, src0, src1, dst, planes, tgpig, tiisg, sgitg); +} + +[[host_name("kernel_mul_mv_q1_0_f32_pc_nr1_4")]] +kernel void kernel_mul_mv_q1_0_f32_pc_nr1_4( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + device const uint * planes, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + kernel_mul_mv_q1_0_f32_pc_impl(args, src0, src1, dst, planes, tgpig, tiisg, sgitg); +} + template void kernel_mul_mv_q2_0_f32_impl( args_t args, From a460ccc750e6b4fd89789242e27aa414a370b617 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Tue, 1 Sep 2026 12:20:59 -0700 Subject: [PATCH 3/5] metal: default the Q1_0 popcount path to nr1=2 (measured best at every batch height) pp4 84.4 vs 62.8 tok/s and pp8 102.7 vs 82.3 against nr1=4, which spills registers (pl[32] + w[nr0][4] + sumf). GGML_METAL_Q1_0_PC_NR1 forces either. --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 4538e15a2653..0bc9d4003ca4 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2602,7 +2602,11 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { ggml_metal_op_concurrency_reset(ctx); { - const int nr1 = ne11 == 2 ? 2 : 4; + static const int nr1_env = getenv("GGML_METAL_Q1_0_PC_NR1") ? atoi(getenv("GGML_METAL_Q1_0_PC_NR1")) : 0; + + // nr1=2 measured best at every batch height; nr1=4 spills registers + // (pp4 84.4 vs 62.8 tok/s, pp8 102.7 vs 82.3) + const int nr1 = (nr1_env == 2 || nr1_env == 4) ? nr1_env : 2; auto pipeline = ggml_metal_library_get_pipeline_mul_mv_q1_0_pc(lib, op, nr1); From 0afb1d1f1aa8802d9ea9e6b1924df6a4a715c58f Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Tue, 1 Sep 2026 12:24:04 -0700 Subject: [PATCH 4/5] metal: allocate the Q1_0 plane scratch only when the popcount path is on The scratch record lives in the dst-buffer tail, so sizing it unconditionally would grow every small-batch Q1_0 mat-mul destination even with the path off. Gate the size helper on the same knob the dispatch reads. --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 0bc9d4003ca4..070b74c0d8a9 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2860,7 +2860,11 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { size_t ggml_metal_op_mul_mat_extra_q1_0_planes(const ggml_tensor * op) { assert(op->op == GGML_OP_MUL_MAT); - if (!op->src[0] || !op->src[1] || op->src[0]->type != GGML_TYPE_Q1_0) { + // the planes are only built when the word-parallel path is enabled, so this + // costs nothing (not even address space) when it is off + static const bool q1_0_pc = getenv("GGML_METAL_Q1_0_POPCNT") != nullptr; + + if (!q1_0_pc || !op->src[0] || !op->src[1] || op->src[0]->type != GGML_TYPE_Q1_0) { return 0; } From 443f0168582966c8a2fc7528970bd81e98c45a09 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Tue, 1 Sep 2026 20:48:48 -0700 Subject: [PATCH 5/5] metal: tighten the Q1_0 plane-scratch gate and drop the half2 kernels The plane-scratch size helper and the encoder disagreed. The encoder only takes the word-parallel path for F32 activations with no src0/src1 broadcast, while the size helper reserved a record for any Q1_0 mat-mul in the 2..16 column range, so unsupported shapes were padding their destination buffer for a path that would never run. Both now call one predicate, ggml_metal_op_mul_mat_q1_0_pc_supported, which keeps the two conditions from drifting apart again. Remove the half2 kernel family: kernel_mul_mv_q1_0_f32_h2_impl, both _h2_nr1_2 and _h2_nr1_4 entry points, and the GGML_METAL_Q1_0_H2 routing branch. It was worth 3.5 to 4.5% at the batch heights it targeted, which does not pay for an fp16 accumulator sitting in the tree with nothing enabling it by default. The multi-column and word-parallel paths are unchanged. Also give the plane-scratch declaration its own comment in ggml-metal-ops.h. It had landed under the "tokens per expert" label that belongs to the declaration below it. test-backend-ops MUL_MAT q1_0 on Metal: 45 OK / 0 FAIL, both with GGML_METAL_Q1_0_POPCNT unset and with it set. Interleaved before/after llama-bench at 2, 4 and 8 columns plus tg128 is unchanged on both arms, all deltas inside run-to-run spread. --- ggml/src/ggml-metal/ggml-metal-device.cpp | 9 +- ggml/src/ggml-metal/ggml-metal-ops.cpp | 59 +++++---- ggml/src/ggml-metal/ggml-metal-ops.h | 4 +- ggml/src/ggml-metal/kernels/mul_mv.metal | 141 ---------------------- 4 files changed, 42 insertions(+), 171 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 77bb09213450..6c0fbbcd2680 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -925,14 +925,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta } else if (nr1_env != 1 && ne11 >= 2) { nr1 = 2; } - // half2 variants pack two src1 columns per lane (opt-in, see the kernel comment) - static const bool h2 = getenv("GGML_METAL_Q1_0_H2") != nullptr; - if (h2 && nr1 >= 2) { - nr1 = nr1 == 2 ? 2 : 4; - suffix = nr1 == 2 ? "_h2_nr1_2" : "_h2_nr1_4"; - } else { - suffix = nr1 == 2 ? "_nr1_2" : nr1 == 3 ? "_nr1_3" : nr1 == 4 ? "_nr1_4" : ""; - } + suffix = nr1 == 2 ? "_nr1_2" : nr1 == 3 ? "_nr1_3" : nr1 == 4 ? "_nr1_4" : ""; } break; case GGML_TYPE_Q2_0: { diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 070b74c0d8a9..37b764cd586d 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2515,6 +2515,40 @@ int ggml_metal_op_pool_2d(ggml_metal_op_t ctx, int idx) { return 1; } +// Q1_0 word-parallel verify path (opt-in): quantize the activation columns into +// int8 bit-planes once, then consume 32 weights per AND+popcount instead of one +// select per weight. See the kernel comment in mul_mv.metal. +// +// The plane scratch is carved out of the padding the allocator adds behind dst, so +// the encoder must take this path exactly when the allocator reserved for it. Both +// call this predicate rather than repeating the shape test. +static bool ggml_metal_op_mul_mat_q1_0_pc_supported(const ggml_tensor * op) { + static const bool q1_0_pc = getenv("GGML_METAL_Q1_0_POPCNT") != nullptr; + + if (!q1_0_pc || !op->src[0] || !op->src[1]) { + return false; + } + + if (op->src[0]->type != GGML_TYPE_Q1_0 || op->src[1]->type != GGML_TYPE_F32) { + return false; + } + + const int64_t ne00 = op->src[0]->ne[0]; + const int64_t ne02 = op->src[0]->ne[2]; + const int64_t ne03 = op->src[0]->ne[3]; + + const int64_t ne11 = op->src[1]->ne[1]; + const int64_t ne12 = op->src[1]->ne[2]; + const int64_t ne13 = op->src[1]->ne[3]; + + if (ne11 < 2 || ne11 > 16 || ne00 < 128 || ne00 % 128 != 0) { + return false; + } + + // the kernel indexes src0/src1 without the broadcast strides + return ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1; +} + int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); @@ -2560,18 +2594,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { const int ne11_mm_min = op->src[0]->type == GGML_TYPE_Q1_0 ? std::max(8, q1_0_mv_max) : 8; - // Q1_0 word-parallel verify path (opt-in): quantize the activation columns into - // int8 bit-planes once, then consume 32 weights per AND+popcount instead of one - // select per weight. See the kernel comment in mul_mv.metal. - static const bool q1_0_pc = getenv("GGML_METAL_Q1_0_POPCNT") != nullptr; - - if (q1_0_pc && - op->src[0]->type == GGML_TYPE_Q1_0 && - op->src[1]->type == GGML_TYPE_F32 && - ne11 >= 2 && ne11 <= 16 && - ne00 % 128 == 0 && - ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1 && - ggml_metal_op_mul_mat_extra_q1_0_planes(op) > 0) { + if (ggml_metal_op_mul_mat_q1_0_pc_supported(op)) { const int32_t nblk = ne00/128; ggml_metal_buffer_id bid_src0 = ggml_metal_get_buffer_id(op->src[0]); @@ -2860,21 +2883,15 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { size_t ggml_metal_op_mul_mat_extra_q1_0_planes(const ggml_tensor * op) { assert(op->op == GGML_OP_MUL_MAT); - // the planes are only built when the word-parallel path is enabled, so this - // costs nothing (not even address space) when it is off - static const bool q1_0_pc = getenv("GGML_METAL_Q1_0_POPCNT") != nullptr; - - if (!q1_0_pc || !op->src[0] || !op->src[1] || op->src[0]->type != GGML_TYPE_Q1_0) { + // the encoder gates on this exact predicate, so nothing is reserved for a shape + // that path would not take -- and nothing at all when the path is off + if (!ggml_metal_op_mul_mat_q1_0_pc_supported(op)) { return 0; } const int64_t ne00 = op->src[0]->ne[0]; const int64_t ne11 = op->src[1]->ne[1]; - if (ne11 < 2 || ne11 > 16 || ne00 % 128 != 0) { - return 0; - } - return (size_t) (ne00/128) * ne11 * Q1_0_PLANE_STRIDE * sizeof(uint32_t); } diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 21a064091d15..e40f2bdb4d64 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -30,8 +30,10 @@ int ggml_metal_op_encode(ggml_metal_op_t ctx, int idx); // available ops: // -// tokens per expert +// activation bit-planes for the Q1_0 word-parallel path size_t ggml_metal_op_mul_mat_extra_q1_0_planes(const struct ggml_tensor * op); + +// tokens per expert size_t ggml_metal_op_mul_mat_id_extra_tpe(const struct ggml_tensor * op); // id map [n_tokens, n_expert] diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index 284b2ecc3f2b..7be9f511357d 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -577,147 +577,6 @@ kernel void kernel_mul_mv_q1_0_f32_nr1_4( kernel_mul_mv_q1_0_f32_nc_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } -// Q1_0 multi-column mat-vec, half2 form. Two src1 columns share one half2 lane, so -// a single select+add covers both and the ALU slots per weight per column halve -- -// the multi-column kernel above is ALU-bound (2 fp32 lane-ops per weight per column), -// not weight-read-bound, at the 2..8 rows a speculative verify batch uses. -// Only the per-row accumulator is half: the block sum of y stays fp32 (it is computed -// once per block for all nr0 rows, so its cost is amortized) and the accumulator is -// promoted to fp32 at every q1_0 block, bounding the rounding to at most SW halves. -// Opt-in via GGML_METAL_Q1_0_H2=1; gate on measured logit-KLD before enabling by default. -template -void kernel_mul_mv_q1_0_f32_h2_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 short SW = QK1_0/tpb; - - const int nb = args.ne00/QK1_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 int c0 = r1 * (2*npair); - - const uint i12 = im%FC_mul_mv_ne12; - const uint i13 = im/FC_mul_mv_ne12; - - device const block_q1_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_q1_0 *) ((device char *) src0 + offset0); - } - - half2 yh[npair][SW]; - float sumy[2*npair]; - float sumf[nr0][2*npair]; - FOR_UNROLL (short row = 0; row < nr0; row++) { - FOR_UNROLL (short c = 0; c < 2*npair; c++) { - sumf[row][c] = 0.f; - } - } - - const short ix = (tiisg/tpb); // block in flight - const short il = (tiisg%tpb)*SW; // element offset within the block - - device const float * yb[2*npair]; - FOR_UNROLL (short c = 0; c < 2*npair; c++) { - // tail columns are clamped (computed but not stored) - const int ic = min(c0 + c, args.ne11 - 1); - const uint64_t offset1 = (uint64_t)ic*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; - yb[c] = (device const float *) (src1 + offset1) + ix*QK1_0 + il; - } - - for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/tpb) { - FOR_UNROLL (short p = 0; p < npair; p++) { - float s0 = 0.f; - float s1 = 0.f; - FOR_UNROLL (short i = 0; i < SW; i++) { - const float y0 = yb[2*p ][i]; - const float y1 = yb[2*p + 1][i]; - yh[p][i] = half2((half) y0, (half) y1); - s0 += y0; - s1 += y1; - } - sumy[2*p ] = s0; - sumy[2*p + 1] = s1; - } - - FOR_UNROLL (short row = 0; row < nr0; row++) { - device const block_q1_0 * qb = ax[row] + ib; - device const uint8_t * qs = qb->qs + il/8; - - uint8_t b[SW/8]; - FOR_UNROLL (short i = 0; i < SW/8; i++) { - b[i] = qs[i]; - } - const float d = qb->d; - - FOR_UNROLL (short p = 0; p < npair; p++) { - half2 acc = half2(0.0h); - - FOR_UNROLL (short i = 0; i < SW; i++) { - acc += bool(b[i/8] & (1u << (i%8))) ? yh[p][i] : half2(0.0h); - } - - sumf[row][2*p ] += d*(2.0f*(float) acc.x - sumy[2*p ]); - sumf[row][2*p + 1] += d*(2.0f*(float) acc.y - sumy[2*p + 1]); - } - } - - FOR_UNROLL (short c = 0; c < 2*npair; c++) { - yb[c] += QK1_0 * (N_SIMDWIDTH/tpb); - } - } - - device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1; - - for (short c = 0; c < 2*npair; c++) { - if (c0 + c >= args.ne11) { - break; - } - for (int row = 0; row < nr0; ++row) { - const float tot = simd_sum(sumf[row][c]); - - if (tiisg == 0 && first_row + row < args.ne01) { - dst_f32[(uint64_t)(c0 + c)*args.ne0 + first_row + row] = tot; - } - } - } -} - -[[host_name("kernel_mul_mv_q1_0_f32_h2_nr1_2")]] -kernel void kernel_mul_mv_q1_0_f32_h2_nr1_2( - 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_q1_0_f32_h2_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); -} - -[[host_name("kernel_mul_mv_q1_0_f32_h2_nr1_4")]] -kernel void kernel_mul_mv_q1_0_f32_h2_nr1_4( - 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_q1_0_f32_h2_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); -} - // --------------------------------------------------------------------------- // Q1_0 word-parallel (popcount) mat-vec for small verify batches. //