From 770ce1ad2150da992228e76a8e4aa6d69ae0c5c5 Mon Sep 17 00:00:00 2001 From: Clemens Wasser Date: Wed, 7 Oct 2026 07:28:23 +0200 Subject: [PATCH] sycl: add IQ3_S multi-column MMVQ (llama/29500) --- ggml/src/ggml-sycl/mmvq.cpp | 133 ++++++++++++++++++++++++++++-------- 1 file changed, 104 insertions(+), 29 deletions(-) diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index d0f090703..ceaacb59e 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -6,6 +6,39 @@ #include "quants.hpp" #include "vecdotq.hpp" +// vec_dot_q_sycl_t adapters for the IQ vec_dots that take their codebook tables as extra +// arguments: bind the constant tables here (as vec_dot_iq2_s_q8_1 / vec_dot_iq1_m_q8_1 already do +// internally) so they can be used as template arguments of mul_mat_vec_q_moe. +static __dpct_inline__ float vec_dot_iq2_xxs_q8_1_moe(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + return vec_dot_iq2_xxs_q8_1(vbq, bq8_1, iqs, iq2xxs_grid, ksigns_iq2xs, kmask_iq2xs); +} + +static __dpct_inline__ float vec_dot_iq2_xs_q8_1_moe(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + return vec_dot_iq2_xs_q8_1(vbq, bq8_1, iqs, iq2xs_grid, ksigns64); +} + +static __dpct_inline__ float vec_dot_iq3_xxs_q8_1_moe(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + return vec_dot_iq3_xxs_q8_1(vbq, bq8_1, iqs, iq3xxs_grid, ksigns64); +} + +static __dpct_inline__ float vec_dot_iq3_s_q8_1_moe(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + return vec_dot_iq3_s_q8_1(vbq, bq8_1, iqs, iq3s_grid); +} + +static __dpct_inline__ float vec_dot_iq1_s_q8_1_moe(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + return vec_dot_iq1_s_q8_1(vbq, bq8_1, iqs, iq1s_grid_gpu); +} + // Minimum weight-row count at which the Q4_K multi-column MMVQ kernel handles two output rows per // subgroup (rows_per_sg == 2) instead of one, when ncols_dst == 2. // @@ -2190,6 +2223,67 @@ static void mul_mat_vec_iq3_s_q8_1_sycl(const void *vx, const void *vy, } } +template +static void mul_mat_vec_iq3_s_q8_1_sycl_ncols(const void * vx, + const void * vy, + float * dst, + const int ncols, + const int nrows, + const int stride_col_y, + const int stride_col_dst, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE); + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_ncols( + vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, item_ct1); + }); + }); +} + +static void mul_mat_vec_iq3_s_q8_1_sycl_switch_ncols(const void * vx, + const void * vy, + float * dst, + const int ncols, + const int nrows, + const int ncols_dst, + const int stride_col_y, + const int stride_col_dst, + dpct::queue_ptr stream) { + switch (ncols_dst) { + case 1: + mul_mat_vec_iq3_s_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); + break; + case 2: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 3: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 4: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 5: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 6: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 7: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + case 8: + mul_mat_vec_iq3_s_q8_1_sycl_ncols<8>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); + break; + default: + GGML_ABORT("unsupported ncols_dst=%d for IQ3_S multi-col MMVQ", ncols_dst); + } +} + static void mul_mat_vec_iq1_s_q8_1_sycl(const void *vx, const void *vy, float *dst, const int ncols, const int nrows, @@ -2627,7 +2721,16 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens mul_mat_vec_iq3_xxs_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); break; case GGML_TYPE_IQ3_S: - mul_mat_vec_iq3_s_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + const int stride_col_y = src1_padded_col_size / QK8_1; + const int stride_col_dst = dst->ne[0]; + GGML_SYCL_DEBUG("Calling mul_mat_vec_iq3_s_q8_1_sycl_switch_ncols ncols=%d\n", (int) src1_ncols); + mul_mat_vec_iq3_s_q8_1_sycl_switch_ncols(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, + src1_ncols, stride_col_y, stride_col_dst, stream); + return; + } else if (i == 0 || src1_ncols == 1) { + mul_mat_vec_iq3_s_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + } break; case GGML_TYPE_IQ4_NL: mul_mat_vec_iq4_nl_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); @@ -2681,34 +2784,6 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens GGML_UNUSED(ctx); } -// vec_dot_q_sycl_t adapters for the IQ vec_dots that take their codebook tables as extra -// arguments: bind the constant tables here (as vec_dot_iq2_s_q8_1 / vec_dot_iq1_m_q8_1 already do -// internally) so they can be used as template arguments of mul_mat_vec_q_moe. -static __dpct_inline__ float vec_dot_iq2_xxs_q8_1_moe(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, const int & iqs) { - return vec_dot_iq2_xxs_q8_1(vbq, bq8_1, iqs, iq2xxs_grid, ksigns_iq2xs, kmask_iq2xs); -} - -static __dpct_inline__ float vec_dot_iq2_xs_q8_1_moe(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, const int & iqs) { - return vec_dot_iq2_xs_q8_1(vbq, bq8_1, iqs, iq2xs_grid, ksigns64); -} - -static __dpct_inline__ float vec_dot_iq3_xxs_q8_1_moe(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, const int & iqs) { - return vec_dot_iq3_xxs_q8_1(vbq, bq8_1, iqs, iq3xxs_grid, ksigns64); -} - -static __dpct_inline__ float vec_dot_iq3_s_q8_1_moe(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, const int & iqs) { - return vec_dot_iq3_s_q8_1(vbq, bq8_1, iqs, iq3s_grid); -} - -static __dpct_inline__ float vec_dot_iq1_s_q8_1_moe(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, const int & iqs) { - return vec_dot_iq1_s_q8_1(vbq, bq8_1, iqs, iq1s_grid_gpu); -} - // src1_row_stride: 0 for shared src1 (gate/up proj), else per-expert stride (down proj). template static void mul_mat_vec_q_moe(