Skip to content

sycl: reduce redundant work in Q4_K multi-column MMVQ #27062

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
arthw merged 4 commits into from
Sep 3, 2026
Merged

sycl: reduce redundant work in Q4_K multi-column MMVQ #27062

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b9af702
sycl: Q4_K Weight unpack optimization and reuse between destination C…
RaulAbejonDelgado Aug 14, 2026
9cd4b8c
sycl: Q4_K small N (N=2..4) + two output rows by subgroup reuse of ac…
RaulAbejonDelgado Aug 14, 2026
4d82396
sycl: gate Q4_K two-row reuse for small N=2
RaulAbejonDelgado Aug 15, 2026
87e2596
sycl: Fix on magic number now uses Q4_K_MMVQ_ROW_PAIR_MIN_NROWS=6272 …
RaulAbejonDelgado Aug 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
195 changes: 157 additions & 38 deletions ggml/src/ggml-sycl/mmvq.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,24 @@
#include "quants.hpp"
#include "vecdotq.hpp"

// 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.
//
// Pairing rows lets a subgroup load each activation block once and apply it to two rows, at the cost
// of halving the number of subgroups in the launch. With only two destination columns there is too
// little work per row to hide that loss of parallelism, so pairing only pays off once there are
// enough rows to keep the device occupied. This is a measured performance crossover, not a
// correctness or hardware limit - both variants compute the same result for any nrows.
//
// Derived on Intel Arc Pro B70 with `test-backend-ops perf -o MUL_MAT` (Q4_K, ncols_dst == 2),
// sweeping nrows over 5120..6912 at ncols 17408 and 19968: one row per subgroup was up to 9% faster
// below the crossover, two rows per subgroup 8-15% faster above it, and the crossover fell inside
// (6144, 6272] for both ncols with no measurable ncols dependence. A later 32-row granularity sweep
// narrowed it to (6144, 6176], so 6272 is a conservative gate rather than the exact crossover.
// ncols_dst >= 3 amortizes the activation loads over more columns and is faster with two rows at
// every row count, so it does not consult this threshold.
static constexpr int Q4_K_MMVQ_ROW_PAIR_MIN_NROWS = 6272;

template <typename reorder_vec_dot_q_sycl>
static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __restrict__ vy, float * __restrict__ dst,
const int ncols, const int nrows, const sycl::nd_item<3> & nd_item) {
Expand Down Expand Up @@ -59,7 +77,7 @@ static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __r

// With has_fusion, `vgate` is a second weight matrix sharing vx's shape, stride and reorder
// layout: one pass computes both row dot products and the epilogue writes glu(gate, up).
template <typename reorder_vec_dot_q_sycl, int ncols_dst, bool has_fusion = false>
template <typename reorder_vec_dot_q_sycl, int ncols_dst, bool has_fusion = false, int rows_per_sg = 1>
static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void * __restrict__ vgate,
const void * __restrict__ vy, float * __restrict__ dst, const int ncols,
const int nrows, const int stride_col_y_bytes, const int stride_col_dst,
Expand All @@ -71,14 +89,17 @@ static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void
const int sg_range = sg.get_group_linear_range();
const int workgroup_id = nd_item.get_group_linear_id();
const int sg_id = sg.get_group_linear_id();
const int row = workgroup_id * sg_range + sg_id;
const int row0 = (workgroup_id * sg_range + sg_id) * rows_per_sg;

// row is sub-group uniform, so this retires whole sub-groups and the collectives below
// stay convergent
if (row >= nrows) {
if (row0 >= nrows) {
return;
}

static_assert(rows_per_sg == 1 ||
reorder_vec_dot_shared_activations<reorder_vec_dot_q_sycl::gtype>::value);

const int blocks_per_row = ncols / block_traits::qk;
constexpr int blocks_per_subgroup = ceil_div(block_traits::vdr_mmvq * WARP_SIZE, block_traits::qi);
constexpr int block_elements_per_subgroup = block_traits::qi / block_traits::vdr_mmvq;
Expand All @@ -87,52 +108,117 @@ static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void
static_assert(blocks_per_subgroup > 0);
static_assert(block_elements_per_subgroup > 0);

float partial_sum[ncols_dst] = { 0.0f };
float partial_sum[ncols_dst][rows_per_sg] = {};
// sized 1 rather than 0 when unused: zero-length arrays are not standard C++, and the
// array is dead and eliminated in that case
[[maybe_unused]] float partial_gate[has_fusion ? ncols_dst : 1] = { 0.0f };
[[maybe_unused]] float partial_gate[has_fusion ? ncols_dst : 1][has_fusion ? rows_per_sg : 1] = {};
for (int i = sg.get_local_linear_id() / block_elements_per_subgroup; i < blocks_per_row; i += blocks_per_subgroup) {
const int ibx = row * blocks_per_row + i;

// the offsets depend only on the block index and the matrix shape, never on the base
// pointer, which is what lets vgate reuse them
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
const int iby = i * block_type::block_to_q8_1_ratio();

#pragma unroll
for (int elem = 0; elem < block_elements_per_subgroup; elem += WARP_SIZE) {
const int iqs = elem + block_traits::vdr_mmvq * (sg.get_local_linear_id() % block_elements_per_subgroup);

if constexpr (rows_per_sg > 1) {
typename reorder_vec_dot_q_sycl::weights wx[rows_per_sg];
[[maybe_unused]] typename reorder_vec_dot_q_sycl::weights wg[rows_per_sg];
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
const sycl::half2 * q8_1_ds_ptr = (const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));
for (int r = 0; r < rows_per_sg; ++r) {
const int row = sycl::min(row0 + r, nrows - 1);
const int ibx = row * blocks_per_row + i;
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
wx[r] = reorder_vec_dot_q_sycl::load(vx, bx_offset, d_offset, iqs);
if constexpr (has_fusion) {
wg[r] = reorder_vec_dot_q_sycl::load(vgate, bx_offset, d_offset, iqs);
}
}
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
const sycl::half2 * q8_1_ds_ptr =
(const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));
const auto a = reorder_vec_dot_q_sycl::load_activations(q8_1_quant_ptr, q8_1_ds_ptr, iqs);
#pragma unroll
for (int r = 0; r < rows_per_sg; ++r) {
partial_sum[j][r] += reorder_vec_dot_q_sycl::apply(wx[r], a);
if constexpr (has_fusion) {
partial_gate[j][r] += reorder_vec_dot_q_sycl::apply(wg[r], a);
}
}
}
} else if constexpr (reorder_vec_dot_shared_weights<reorder_vec_dot_q_sycl::gtype>::value) {
const int ibx = row0 * blocks_per_row + i;
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
const auto wx = reorder_vec_dot_q_sycl::load(vx, bx_offset, d_offset, iqs);
if constexpr (has_fusion) {
const auto wg = reorder_vec_dot_q_sycl::load(vgate, bx_offset, d_offset, iqs);

partial_sum[j] += reorder_vec_dot_q_sycl()(vx, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
const sycl::half2 * q8_1_ds_ptr =
(const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));

if constexpr (has_fusion) {
partial_gate[j] +=
reorder_vec_dot_q_sycl()(vgate, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
// up and gate share the activation, so load it once and apply it twice
const auto a = reorder_vec_dot_q_sycl::load_activations(q8_1_quant_ptr, q8_1_ds_ptr, iqs);

partial_sum[j][0] += reorder_vec_dot_q_sycl::apply(wx, a);
partial_gate[j][0] += reorder_vec_dot_q_sycl::apply(wg, a);
}
} else {
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
const sycl::half2 * q8_1_ds_ptr =
(const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));

partial_sum[j][0] += reorder_vec_dot_q_sycl::dot(wx, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
}
}
} else {
const int ibx = row0 * blocks_per_row + i;
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
const sycl::half2 * q8_1_ds_ptr =
(const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));

partial_sum[j][0] +=
reorder_vec_dot_q_sycl()(vx, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);

if constexpr (has_fusion) {
partial_gate[j][0] +=
reorder_vec_dot_q_sycl()(vgate, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
}
}
}
}
}

#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
float sum = sycl::reduce_over_group(nd_item.get_sub_group(), partial_sum[j], std::plus<>());
#pragma unroll
for (int r = 0; r < rows_per_sg; ++r) {
float sum = sycl::reduce_over_group(nd_item.get_sub_group(), partial_sum[j][r], std::plus<>());

if constexpr (has_fusion) {
const float gate = sycl::reduce_over_group(nd_item.get_sub_group(), partial_gate[j], std::plus<>());
if constexpr (has_fusion) {
const float gate = sycl::reduce_over_group(nd_item.get_sub_group(), partial_gate[j][r], std::plus<>());

// uniform across the launch; the launcher only instantiates SWIGLU and GEGLU
sum *= glu_op == GGML_GLU_OP_SWIGLU ? op_silu(gate) : op_gelu(gate);
}
// uniform across the launch; the launcher only instantiates SWIGLU and GEGLU
sum *= glu_op == GGML_GLU_OP_SWIGLU ? op_silu(gate) : op_gelu(gate);
}

if (sg.leader()) {
dst[j * stride_col_dst + row] = sum;
if (sg.leader() && row0 + r < nrows) {
dst[j * stride_col_dst + row0 + r] = sum;
}
}
}
}
Expand Down Expand Up @@ -1671,37 +1757,54 @@ static void reorder_mul_mat_vec_q4_k_q8_1_sycl(const void * vx, const void * vy,
});
}

template <int ncols_dst>
static void reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols(
template <int ncols_dst, int rows_per_sg>
static void reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols_impl(
const void * vx, const void * vy, float * dst,
const int ncols, const int nrows,
const int stride_col_y_bytes, const int stride_col_dst,
dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK_K == 0);

constexpr size_t num_subgroups = WARP_SIZE;
const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups);
const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups * rows_per_sg);
const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);

stream->submit([&](sycl::handler & cgh) {
cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
mul_mat_vec_q_reorder_ncols<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>, ncols_dst>(
mul_mat_vec_q_reorder_ncols<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>, ncols_dst,
/*has_fusion=*/ false, rows_per_sg>(
vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
/*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}

template <int ncols_dst>
static void reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols(
const void * vx, const void * vy, float * dst,
const int ncols, const int nrows,
const int stride_col_y_bytes, const int stride_col_dst,
dpct::queue_ptr stream) {
constexpr int rows_per_sg = ncols_dst >= 3 && ncols_dst <= 4 ? 2 : 1;
reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols_impl<ncols_dst, rows_per_sg>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream);
}

static void reorder_mul_mat_vec_q4_k_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_bytes, const int stride_col_dst,
dpct::queue_ptr stream) {
switch (ncols_dst) {
case 1: reorder_mul_mat_vec_q4_k_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break;
case 2: reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break;
case 2:
if (nrows >= Q4_K_MMVQ_ROW_PAIR_MIN_NROWS) {
reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols_impl<2, 2>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream);
} else {
reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols_impl<2, 1>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream);
}
break;
case 3: reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break;
case 4: reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break;
case 5: reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break;
Expand Down Expand Up @@ -2839,29 +2942,42 @@ bool ggml_sycl_mul_mat_vec_q_id_reorder(
}
}

template <typename reorder_vec_dot_q_sycl, int ncols_dst>
static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate, const void * vy, float * dst,
template <typename reorder_vec_dot_q_sycl, int ncols_dst, int rows_per_sg>
static void launch_mul_mat_vec_q_reorder_glu_impl(const void * vx, const void * vgate, const void * vy, float * dst,
const int ncols, const int nrows, const int stride_col_y_bytes,
const int stride_col_dst, const ggml_glu_op glu_op,
dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK_K == 0);

constexpr size_t num_subgroups = WARP_SIZE;

const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups);
const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups * rows_per_sg);
const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);

stream->submit([&](sycl::handler & cgh) {
cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
mul_mat_vec_q_reorder_ncols<reorder_vec_dot_q_sycl, ncols_dst, /*has_fusion=*/ true>(
mul_mat_vec_q_reorder_ncols<reorder_vec_dot_q_sycl, ncols_dst, /*has_fusion=*/ true,
rows_per_sg>(
vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op,
nd_item);
});
});
}

template <typename reorder_vec_dot_q_sycl, int ncols_dst>
static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate, const void * vy, float * dst,
const int ncols, const int nrows, const int stride_col_y_bytes,
const int stride_col_dst, const ggml_glu_op glu_op,
dpct::queue_ptr stream) {
constexpr int rows_per_sg =
reorder_vec_dot_shared_activations<reorder_vec_dot_q_sycl::gtype>::value && ncols_dst >= 3 && ncols_dst <= 4
? 2
: 1;
launch_mul_mat_vec_q_reorder_glu_impl<reorder_vec_dot_q_sycl, ncols_dst, rows_per_sg>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream);
}

bool ggml_sycl_mul_mat_vec_q_glu_reorder(enum ggml_type src0_type, enum ggml_glu_op glu_op, const void * vx,
const void * vgate, const void * vy, float * dst, int ncols, int nrows,
int ncols_dst, int stride_col_y_bytes, int stride_col_dst,
Expand All @@ -2881,8 +2997,11 @@ bool ggml_sycl_mul_mat_vec_q_glu_reorder(enum ggml_type src0_type, enum ggml_glu
stride_col_dst, glu_op, stream);
return true;
case 2:
launch_mul_mat_vec_q_reorder_glu<vec_dot, 2>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
stride_col_dst, glu_op, stream);
if (nrows >= Q4_K_MMVQ_ROW_PAIR_MIN_NROWS) {
launch_mul_mat_vec_q_reorder_glu_impl<vec_dot, 2, 2>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream);
} else {
launch_mul_mat_vec_q_reorder_glu_impl<vec_dot, 2, 1>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream);
}
return true;
case 3:
launch_mul_mat_vec_q_reorder_glu<vec_dot, 3>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
Expand Down
Loading
Loading