diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index e85c76e0cca..7cf163b677f 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" @@ -410,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); @@ -860,6 +914,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-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 5d68090250b..07aead10c38 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 01562df5739..394bbb1b2a7 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 @@ -505,6 +510,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 33bcd76220a..35bf8caf743 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" @@ -2521,6 +2522,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); @@ -2558,7 +2593,90 @@ 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; + + 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]); + 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); + + { + 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); + + 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] @@ -2569,7 +2687,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_PTQ1_0 || @@ -2767,6 +2885,24 @@ 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); + + // 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]; + + 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 159a628d04a..e40f2bdb4d6 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -30,6 +30,9 @@ int ggml_metal_op_encode(ggml_metal_op_t ctx, int idx); // available ops: // +// 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); diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp index e5b2ee8a55b..a2fdde0b3ae 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 222902cc7e0..4997a611b3d 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -430,6 +430,346 @@ 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); +} + +// --------------------------------------------------------------------------- +// 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,