Commit 78651c410 for llama.cpp
commit 78651c410dd8d97e3e22e533e0e7117889e3863a
Author: Clemens Wasser <clemens.wasser@gmail.com>
Date: Wed Oct 7 07:28:23 2026 +0200
sycl: add IQ3_S multi-column MMVQ (#29500)
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 <int ncols_dst>
+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<QK_K, QI3_S / 2, block_iq3_s, 1, vec_dot_iq3_s_q8_1_moe, ncols_dst>(
+ 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 <int qk, int qi, typename block_q_t, int vdr, vec_dot_q_sycl_t vec_dot_q_sycl>
static void mul_mat_vec_q_moe(