Superminaren/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-04 02:37:27 +02:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source 
* sycl: Q4_K Weight unpack optimization and reuse between destination Columns

* sycl: Q4_K small N (N=2..4) + two output rows by subgroup reuse of activation between two rows.

* sycl: gate Q4_K two-row reuse for small N=2

* sycl: Fix on magic number now uses Q4_K_MMVQ_ROW_PAIR_MIN_NROWS=6272 for it, added tests for coverage around Q4_K_MMVQ_ROW_PAIR_MIN_NROWS with perf support to test Q4_K MUL_MAT, applied the same  reuse pattern to the activation as the weights.

Assisted-by: GPT-5.6 Sol

---------

Co-authored-by: RaulAbejonDelgado <[email protected]>
This commit is contained in:
Eurekatic
2026-09-03 14:59:06 +08:00
committed by GitHub
co-authored by RaulAbejonDelgado
parent c61b98b875
commit 4aa6ffba25
3 changed files with 293 additions and 80 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/mmvq.cpp
+158 -39
View File
@@ -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) {
@@ -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,
@@ -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;
@@ -87,34 +108,96 @@ 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));
partial_sum[j] += reorder_vec_dot_q_sycl()(vx, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
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) {
partial_gate[j] +=
reorder_vec_dot_q_sycl()(vgate, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
const auto wg = 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));
// 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);
}
}
}
}
@@ -122,17 +205,20 @@ static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void
#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;
}
}
}
}
@@ -1671,8 +1757,8 @@ 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,
@@ -1680,20 +1766,31 @@ static void reorder_mul_mat_vec_q4_k_q8_1_sycl_ncols(
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,
@@ -1701,7 +1798,13 @@ static void reorder_mul_mat_vec_q4_k_q8_1_sycl_switch_ncols(
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;
@@ -2839,8 +2942,8 @@ 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) {
@@ -2848,20 +2951,33 @@ static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate
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,
@@ -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,
ggml/src/ggml-sycl/vecdotq.hpp
+94 -41
View File
@@ -351,6 +351,25 @@ template <ggml_type T> struct reorder_vec_dot_q_sycl {
static_assert(T != T, "ggml_type for reorder vecdot not implemented");
};
// For some types the weight side of the dot product does not depend on the destination column, so a
// multi-column mul_mat_vec can unpack it once per block instead of once per column. Such a type adds
// load() and dot() next to operator() and opts in here. See reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>.
template <ggml_type T> struct reorder_vec_dot_shared_weights {
static constexpr bool value = false;
};
template <> struct reorder_vec_dot_shared_weights<GGML_TYPE_Q4_K> {
static constexpr bool value = true;
};
template <ggml_type T> struct reorder_vec_dot_shared_activations {
static constexpr bool value = false;
};
template <> struct reorder_vec_dot_shared_activations<GGML_TYPE_Q4_K> {
static constexpr bool value = true;
};
template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0> {
static constexpr ggml_type gtype = GGML_TYPE_Q4_0;
@@ -540,50 +559,84 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K> {
using q4_k_block = ggml_sycl_reordered::block_q_t<GGML_TYPE_Q4_K>;
using q4_k_traits = typename q4_k_block::traits;
struct weights {
int v[2];
uint16_t aux[2];
ggml_half2 dm;
int bq8_offset;
};
struct activations {
int u[2 * QR4_K];
float d8[QR4_K];
};
__dpct_inline__ static weights load(const void * __restrict__ vbq, const std::pair<int, int> ibx_offset,
const std::pair<int, int> d_offset, const int & iqs) {
const uint8_t * base = static_cast<const uint8_t *>(vbq);
const uint8_t * qs = base + ibx_offset.first;
const uint8_t * scs = base + d_offset.first;
const ggml_half2 * dms = reinterpret_cast<const ggml_half2 *>(base + d_offset.second);
weights w;
w.bq8_offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
const int * q4 = (const int *) (qs + 16 * w.bq8_offset + 4 * ((iqs / 2) % 4));
const uint16_t * scales = (const uint16_t *) scs;
w.v[0] = q4[0];
w.v[1] = q4[4];
const int j = (QR4_K * ((iqs / 2) / (QI8_1 / 2))) / 2;
if (j < 2) {
w.aux[0] = scales[j + 0] & 0x3f3f;
w.aux[1] = scales[j + 2] & 0x3f3f;
} else {
w.aux[0] = ((scales[j + 2] >> 0) & 0x0f0f) | ((scales[j - 2] & 0xc0c0) >> 2);
w.aux[1] = ((scales[j + 2] >> 4) & 0x0f0f) | ((scales[j - 0] & 0xc0c0) >> 2);
}
w.dm = *dms;
return w;
}
__dpct_inline__ static activations load_activations(const int8_t * q8_1_quant_ptr,
const sycl::half2 * q8_1_ds, const int & iqs) {
activations a;
const int bq8_offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
for (int i = 0; i < QR4_K; ++i) {
const int8_t * quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1;
sycl::half2 ds_values = *(q8_1_ds + bq8_offset + i);
a.d8[i] = ds_values[0];
const int * q8 = (const int *) quant_base_ptr + ((iqs / 2) % 4);
a.u[2 * i + 0] = q8[0];
a.u[2 * i + 1] = q8[4];
}
return a;
}
__dpct_inline__ static float apply(const weights & w, const activations & a) {
const uint8_t * sc = (const uint8_t *) w.aux;
const uint8_t * m = sc + 2;
return vec_dot_q4_K_q8_1_impl_vmmq(w.v, a.u, sc, m, w.dm, a.d8);
}
__dpct_inline__ static float dot(const weights & w, const int8_t * q8_1_quant_ptr,
const sycl::half2 * q8_1_ds, const int & iqs) {
const auto a = load_activations(q8_1_quant_ptr, q8_1_ds, iqs);
return apply(w, a);
}
__dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair<int, int> ibx_offset,
const std::pair<int, int> d_offset, const int8_t * q8_1_quant_ptr,
const sycl::half2 * q8_1_ds, const int & iqs) {
const uint8_t * base = static_cast<const uint8_t *>(vbq);
const uint8_t * qs = base + ibx_offset.first;
const uint8_t * scs = base + d_offset.first;
const ggml_half2 * dms = reinterpret_cast<const ggml_half2 *>(base + d_offset.second);
const int bq8_offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
const int * q4 = (const int *) (qs + 16 * bq8_offset + 4 * ((iqs / 2) % 4));
const uint16_t * scales = (const uint16_t *) scs;
int v[2];
int u[2 * QR4_K];
float d8[QR4_K];
v[0] = q4[0];
v[1] = q4[4];
uint16_t aux[2];
const int j = (QR4_K * ((iqs / 2) / (QI8_1 / 2))) / 2;
if (j < 2) {
aux[0] = scales[j + 0] & 0x3f3f;
aux[1] = scales[j + 2] & 0x3f3f;
} else {
aux[0] = ((scales[j + 2] >> 0) & 0x0f0f) | ((scales[j - 2] & 0xc0c0) >> 2);
aux[1] = ((scales[j + 2] >> 4) & 0x0f0f) | ((scales[j - 0] & 0xc0c0) >> 2);
}
const uint8_t * sc = (const uint8_t *) aux;
const uint8_t * m = sc + 2;
for (int i = 0; i < QR4_K; ++i) {
const int8_t* quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1;
sycl::half2 ds_values = *(q8_1_ds + bq8_offset + i);
d8[i] = ds_values[0];
const int * q8 = (const int *) quant_base_ptr + ((iqs / 2) % 4);
u[2 * i + 0] = q8[0];
u[2 * i + 1] = q8[4];
}
return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, *dms, d8);
return dot(load(vbq, ibx_offset, d_offset, iqs), q8_1_quant_ptr, q8_1_ds, iqs);
}
};
tests/test-backend-ops.cpp
+41
View File
@@ -9435,6 +9435,21 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 16, 8, 16*256, { 1, 1}, {1, 1}));
}
// Multi-column MMVQ coverage for the Q4_K weight-reuse path and a Q5_K control.
for (ggml_type type_a : { GGML_TYPE_Q4_K, GGML_TYPE_Q5_K }) {
for (int n = 1; n <= 8; ++n) {
test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 4096, n, 1024, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 1023, n, 4096, { 1, 1 }, { 1, 1 }));
}
}
// The SYCL backend picks between one and two output rows per subgroup by row count when there
// are two destination columns (Q4_K_MMVQ_ROW_PAIR_MIN_NROWS in ggml-sycl/mmvq.cpp). Cover both
// sides of that boundary, including an odd row count above it for the row-pair tail.
for (int64_t m : {6271, 6272, 6273}) {
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_K, GGML_TYPE_F32, m, 2, 1024, { 1, 1 }, { 1, 1 }));
}
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_MXFP4, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
@@ -10364,6 +10379,22 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
true, 16, 8, b, false, true, false));
}
// Fused row-pair coverage: minimum rows, an even pair, and an odd tail.
for (ggml_glu_op glu_op : { GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU }) {
for (int64_t m_batch : { 2, 3, 4 }) {
for (int64_t rows : { 1, 2, 3 }) {
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, m_batch, rows, 256,
false, 16, 8, false, false, true, false, { 1, 1 }));
}
}
}
// Both sides of the same row-count boundary as above, on the fused path.
for (int64_t rows : {6271, 6272, 6273}) {
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 2, rows, 256,
false, 16, 8, false, false, true, false, { 1, 1 }));
}
for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT, GATING_FUNC_SQRT_SOFTPLUS}) {
for (bool with_norm : {false, true}) {
for (bool bias_probs : {false, true}) {
@@ -10651,6 +10682,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
}
}
// Q4_K multi-column mat-vec, at ffn_up/ffn_gate geometry (k = n_embd, m = n_ff): n sweeps the
// per-column specializations used for short prompts and speculative/MTP verify, and m brackets
// the row count at which the SYCL backend switches to two output rows per subgroup
// (Q4_K_MMVQ_ROW_PAIR_MIN_NROWS in ggml-sycl/mmvq.cpp), so both sides of it can be measured.
for (int64_t m : {4096, 6144, 6272, 14336}) {
for (int bs : {1, 2, 3, 4, 8}) {
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_K, GGML_TYPE_F32, m, bs, 4096, {1, 1}, {1, 1}));
}
}
// qwen3-30b-a3b
for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) {
for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) {