From 0dc2eed1d307446eb0b3c0e0ee1dd34cd36442b5 Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Mon, 21 Sep 2026 05:32:00 -0700 Subject: [PATCH 1/7] sycl: add PTQ1_0 and PQ2_0 support with MMVQ vector dot kernel - Implement base-3 trit unpacking and dequantization (dequantize_ptq1_0, dequantize_pq2_0) in dequantize.hpp and convert.cpp - Implement vec_dot_ptq1_0_q8_1 and mul_mat_vec_ptq1_0_q8_1_sycl in vecdotq.hpp and mmvq.cpp - Add get_rows and cpy dispatch for PTQ1_0 and PQ2_0 - Enable PTQ1_0 in can_use_mul_mat_vec_q() for fast in-register Level-Zero vector execution - Verified on Intel Arc Pro B50 with oneAPI 2026.1 / Level-Zero (2.20 tok/s on 27B) --- ggml/src/ggml-sycl/convert.cpp | 12 +++++ ggml/src/ggml-sycl/cpy.cpp | 32 +++++++++++++ ggml/src/ggml-sycl/dequantize.hpp | 53 ++++++++++++++++++++++ ggml/src/ggml-sycl/getrows.cpp | 8 ++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 10 ++++- ggml/src/ggml-sycl/mmvq.cpp | 74 +++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/vecdotq.hpp | 56 +++++++++++++++++++++++ 7 files changed, 244 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index b660b56ab39b..b046091512ea 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -657,6 +657,10 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { switch (type) { case GGML_TYPE_Q1_0: return dequantize_block_sycl; + case GGML_TYPE_PTQ1_0: + return dequantize_block_sycl; + case GGML_TYPE_PQ2_0: + return dequantize_block_sycl; case GGML_TYPE_Q2_0: return dequantize_block_sycl; case GGML_TYPE_Q4_0: @@ -747,6 +751,10 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { switch (type) { case GGML_TYPE_Q1_0: return dequantize_block_sycl; + case GGML_TYPE_PTQ1_0: + return dequantize_block_sycl; + case GGML_TYPE_PQ2_0: + return dequantize_block_sycl; case GGML_TYPE_Q2_0: return dequantize_block_sycl; case GGML_TYPE_Q4_0: @@ -861,6 +869,10 @@ to_fp16_nc_sycl_t ggml_get_to_fp16_nc_sycl(ggml_type type) { #endif case GGML_TYPE_Q1_0: return dequantize_block_nc_sycl; + case GGML_TYPE_PTQ1_0: + return dequantize_block_nc_sycl; + case GGML_TYPE_PQ2_0: + return dequantize_block_nc_sycl; case GGML_TYPE_Q4_0: return dequantize_block_nc_sycl; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-sycl/cpy.cpp b/ggml/src/ggml-sycl/cpy.cpp index ef7413abd880..1cfb7890ae7c 100644 --- a/ggml/src/ggml-sycl/cpy.cpp +++ b/ggml/src/ggml-sycl/cpy.cpp @@ -619,6 +619,8 @@ static void ggml_cpy_f16_q5_0_sycl(const char * cx, char * cdst, const int ne, c static bool ggml_sycl_is_quantized_type(enum ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: + case GGML_TYPE_PQ2_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -955,6 +957,32 @@ static void ggml_cpy_q1_0_q1_0(const char * cx, char * cdst, const int ne, const }); } +static void ggml_cpy_ptq1_0_ptq1_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, + const int ne02, const int nb00, const int nb01, const int nb02, const int nb03, + const int ne10, const int ne11, const int ne12, const int nb10, const int nb11, + const int nb12, const int nb13, queue_ptr stream) { + GGML_ASSERT(ne % QK_PTQ1_0 == 0); + const int num_blocks = ceil_div(ne / QK_PTQ1_0, SYCL_CPY_BLOCK_SIZE); + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), + [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + cpy_q_q(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1); + }); +} + +static void ggml_cpy_pq2_0_pq2_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, + const int ne02, const int nb00, const int nb01, const int nb02, const int nb03, + const int ne10, const int ne11, const int ne12, const int nb10, const int nb11, + const int nb12, const int nb13, queue_ptr stream) { + GGML_ASSERT(ne % QK_PQ2_0 == 0); + const int num_blocks = ceil_div(ne / QK_PQ2_0, SYCL_CPY_BLOCK_SIZE); + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), + [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + cpy_q_q(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1); + }); +} + static void ggml_cpy_q2_0_q2_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11, const int ne12, const int nb10, const int nb11, @@ -1366,6 +1394,10 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co ggml_cpy_q4_1_q4_1(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); } else if (src0->type == GGML_TYPE_Q1_0 && src1->type == GGML_TYPE_Q1_0) { ggml_cpy_q1_0_q1_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } else if (src0->type == GGML_TYPE_PTQ1_0 && src1->type == GGML_TYPE_PTQ1_0) { + ggml_cpy_ptq1_0_ptq1_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } else if (src0->type == GGML_TYPE_PQ2_0 && src1->type == GGML_TYPE_PQ2_0) { + ggml_cpy_pq2_0_pq2_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); } else if (src0->type == GGML_TYPE_Q2_0 && src1->type == GGML_TYPE_Q2_0) { ggml_cpy_q2_0_q2_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); } else if (src0->type == GGML_TYPE_MXFP4 && src1->type == GGML_TYPE_MXFP4) { diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 1b13e0f1a31a..6bb4dd4515a9 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -123,6 +123,59 @@ static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib, v.y() = (2 * bit_1 - 1) * d; } +static __dpct_inline__ int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; + int n; + if (e < 80) { + b = x->qs[e & 15]; + n = e >> 4; + } else if (e < 120) { + const int t = e - 80; + b = x->qs[16 + (t & 7)]; + n = t >> 3; + } else { + const int t = e - 120; + b = x->qh[t & 1]; + n = t >> 1; + } + + uint32_t v = b; +#pragma unroll + for (int i = 0; i < 4; ++i) { + if (i < n) { + v = (v * 3) & 0xFF; + } + } + return (int) ((v * 3) >> 8) - 1; +} + +static __dpct_inline__ void dequantize_ptq1_0(const void * vx, const int64_t ib, + const int iqs, dfloat2 & v) { + const block_ptq1_0 * x = (const block_ptq1_0 *) vx; + const dfloat d = x[ib].d; + + v.x() = ptq1_0_trit(&x[ib], iqs + 0) * d; + v.y() = ptq1_0_trit(&x[ib], iqs + 1) * d; +} + +static __dpct_inline__ void dequantize_pq2_0(const void * vx, const int64_t ib, + const int iqs, dfloat2 & v) { + const block_pq2_0 * x = (const block_pq2_0 *) vx; + const dfloat d = x[ib].d; + + const int byte_index_0 = iqs / 4; + const int bit_offset_0 = (iqs % 4) * 2; + + const int byte_index_1 = (iqs + 1) / 4; + const int bit_offset_1 = ((iqs + 1) % 4) * 2; + + const int c0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 0x3; + const int c1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 0x3; + + v.x() = (c0 - 1) * d; + v.y() = (c1 - 1) * d; +} + static __dpct_inline__ void dequantize_q4_1(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { const block_q4_1 * x = (const block_q4_1 *) vx; diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 2113f3563398..36f840e6f5d1 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -276,6 +276,14 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; + case GGML_TYPE_PTQ1_0: + get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + break; + case GGML_TYPE_PQ2_0: + get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + break; case GGML_TYPE_MXFP4: get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 23bce9a01794..53d771bc1ebc 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -1053,6 +1053,8 @@ static int64_t get_row_rounding(ggml_type type, const std::array= VER_GEN9 ? 128 : 64; @@ -4485,7 +4487,9 @@ static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_ } static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { - return ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && + return ggml_is_quantized(src0->type) && + src0->type != GGML_TYPE_PQ2_0 && + src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; } @@ -6034,6 +6038,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons case GGML_TYPE_BF16: case GGML_TYPE_F32: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: + case GGML_TYPE_PQ2_0: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: case GGML_TYPE_IQ2_XXS: @@ -6159,6 +6165,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons if (src1_type == GGML_TYPE_F32) { if (src0_type == GGML_TYPE_Q1_0 || + src0_type == GGML_TYPE_PTQ1_0 || + src0_type == GGML_TYPE_PQ2_0 || src0_type == GGML_TYPE_NVFP4 || src0_type == GGML_TYPE_Q2_K || src0_type == GGML_TYPE_Q3_K || diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index f6a8591e3b40..e8ed8eba4590 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1278,6 +1278,66 @@ static void mul_mat_vec_q1_0_q8_1_sycl_switch_ncols( } } +static void mul_mat_vec_ptq1_0_q8_1_sycl(const void * vx, const void * vy, + float * dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_PTQ1_0 == 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( + vx, vy, dst, ncols, nrows, item_ct1); + }); + }); +} + +template +static void mul_mat_vec_ptq1_0_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_PTQ1_0 == 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_ptq1_0_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_ptq1_0_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break; + case 2: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 3: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 4: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 5: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 6: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 7: mul_mat_vec_ptq1_0_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 8: mul_mat_vec_ptq1_0_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 PTQ1_0 multi-col MMVQ", ncols_dst); + } +} + static void mul_mat_vec_q2_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, @@ -2341,6 +2401,20 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens mul_mat_vec_q1_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); } break; + case GGML_TYPE_PTQ1_0: + 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_ptq1_0_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); + mul_mat_vec_ptq1_0_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) { + GGML_SYCL_DEBUG("Calling mul_mat_vec_ptq1_0_q8_1_sycl\n"); + mul_mat_vec_ptq1_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + } + break; case GGML_TYPE_Q2_0: if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { const int stride_col_y = src1_padded_col_size / QK8_1; diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 3ad4cee93a14..8e8b17718338 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -344,6 +344,62 @@ vec_dot_q1_0_q8_1(const void *__restrict__ vbq, return d1 * bq8_1_chunk->ds[0] * sumi; } +#define VDR_PTQ1_0_Q8_1_MMVQ 4 + +static __dpct_inline__ float +vec_dot_ptq1_0_q8_1(const void *__restrict__ vbq, + const block_q8_1 *__restrict__ bq8_1, const int &iqs) { + GGML_UNUSED(iqs); + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq; + int sumi[4] = { 0, 0, 0, 0 }; + +#pragma unroll + for (int m = 0; m < 16; ++m) { + uint32_t v = bq->qs[m]; +#pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = t * 16 + m; + sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; + } + } + +#pragma unroll + for (int m = 0; m < 8; ++m) { + uint32_t v = bq->qs[16 + m]; +#pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 80 + t * 8 + m; + sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; + } + } + +#pragma unroll + for (int h = 0; h < 2; ++h) { + uint32_t v = bq->qh[h]; +#pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 120 + t * 2 + h; + sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; + } + } + + float acc = 0.0f; +#pragma unroll + for (int k = 0; k < 4; ++k) { + acc += ((const float) bq8_1[k].ds[0]) * (float) sumi[k]; + } + return (float) bq->d * acc; +} + // VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called // MMVQ = mul_mat_vec_q, MMQ = mul_mat_q From 016be855cdb22b23125705af851c81d2e8ad9177 Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Mon, 21 Sep 2026 05:58:36 -0700 Subject: [PATCH 2/7] sycl: vectorize PTQ1_0 dot product with SIMD trit unpacking and DP4A --- ggml/src/ggml-sycl/vecdotq.hpp | 75 +++++++++++++++++++++++----------- 1 file changed, 51 insertions(+), 24 deletions(-) diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 8e8b17718338..9cf49e9b0a93 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -353,43 +353,70 @@ vec_dot_ptq1_0_q8_1(const void *__restrict__ vbq, const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq; int sumi[4] = { 0, 0, 0, 0 }; + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes #pragma unroll - for (int m = 0; m < 16; ++m) { - uint32_t v = bq->qs[m]; + for (int g = 0; g < 4; ++g) { + const uint32_t packed = get_int_from_uint8_aligned(bq->qs, g); + uint32_t v_lo = (packed & 0x000000FF) | ((packed & 0x0000FF00) << 8); + uint32_t v_hi = ((packed >> 16) & 0x000000FF) | ((packed & 0xFF000000) >> 8); + #pragma unroll for (int t = 0; t < 5; ++t) { - const uint32_t w = v * 3; - const int q = (int) (w >> 8) - 1; - v = w & 0xFF; - const int e = t * 16 + m; - sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const uint32_t perm = ((w_lo >> 8) & 0x000000FF) | + ((w_lo >> 16) & 0x0000FF00) | + ((w_hi << 8) & 0x00FF0000) | + (w_hi & 0xFF000000); + const int q = byte_sub_4(perm, 0x01010101); + const int e = t * 16 + 4 * g; + const int u = get_int_from_int8_aligned(bq8_1[e >> 5].qs, (e & 31) >> 2); + sumi[e >> 5] = dpct::dp4a(q, u, sumi[e >> 5]); } } #pragma unroll - for (int m = 0; m < 8; ++m) { - uint32_t v = bq->qs[16 + m]; + for (int g = 0; g < 2; ++g) { + const uint32_t packed = get_int_from_uint8_aligned(bq->qs + 16, g); + uint32_t v_lo = (packed & 0x000000FF) | ((packed & 0x0000FF00) << 8); + uint32_t v_hi = ((packed >> 16) & 0x000000FF) | ((packed & 0xFF000000) >> 8); + #pragma unroll for (int t = 0; t < 5; ++t) { - const uint32_t w = v * 3; - const int q = (int) (w >> 8) - 1; - v = w & 0xFF; - const int e = 80 + t * 8 + m; - sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const uint32_t perm = ((w_lo >> 8) & 0x000000FF) | + ((w_lo >> 16) & 0x0000FF00) | + ((w_hi << 8) & 0x00FF0000) | + (w_hi & 0xFF000000); + const int q = byte_sub_4(perm, 0x01010101); + const int e = 80 + t * 8 + 4 * g; + const int u = get_int_from_int8_aligned(bq8_1[e >> 5].qs, (e & 31) >> 2); + sumi[e >> 5] = dpct::dp4a(q, u, sumi[e >> 5]); } } + uint32_t v = (uint32_t) bq->qh[0] | ((uint32_t) bq->qh[1] << 16); #pragma unroll - for (int h = 0; h < 2; ++h) { - uint32_t v = bq->qh[h]; -#pragma unroll - for (int t = 0; t < 4; ++t) { - const uint32_t w = v * 3; - const int q = (int) (w >> 8) - 1; - v = w & 0xFF; - const int e = 120 + t * 2 + h; - sumi[e >> 5] += q * (int) bq8_1[e >> 5].qs[e & 31]; - } + for (int t = 0; t < 4; t += 2) { + const uint32_t w0 = v * 3; + v = w0 & 0x00FF00FF; + const uint32_t w1 = v * 3; + v = w1 & 0x00FF00FF; + + const uint32_t perm = ((w0 >> 8) & 0x000000FF) | + ((w0 >> 16) & 0x0000FF00) | + ((w1 << 8) & 0x00FF0000) | + (w1 & 0xFF000000); + const int q = byte_sub_4(perm, 0x01010101); + const int u = get_int_from_int8_aligned(bq8_1[3].qs, 6 + t / 2); + sumi[3] = dpct::dp4a(q, u, sumi[3]); } float acc = 0.0f; From c5f6c9718820d526111b9f895a6ac826ec9954e9 Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Mon, 21 Sep 2026 06:20:05 -0700 Subject: [PATCH 3/7] sycl: add native PQ2_0 MMVQ vector dot kernel with DP4A --- ggml/src/ggml-sycl/ggml-sycl.cpp | 1 - ggml/src/ggml-sycl/mmvq.cpp | 74 ++++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/vecdotq.hpp | 26 +++++++++++ 3 files changed, 100 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 53d771bc1ebc..5bbb39175e62 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4488,7 +4488,6 @@ static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_ static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { return ggml_is_quantized(src0->type) && - src0->type != GGML_TYPE_PQ2_0 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; } diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index e8ed8eba4590..abc88b7bc5c3 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1338,6 +1338,66 @@ static void mul_mat_vec_ptq1_0_q8_1_sycl_switch_ncols( } } +static void mul_mat_vec_pq2_0_q8_1_sycl(const void * vx, const void * vy, + float * dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_PQ2_0 == 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( + vx, vy, dst, ncols, nrows, item_ct1); + }); + }); +} + +template +static void mul_mat_vec_pq2_0_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_PQ2_0 == 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_pq2_0_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_pq2_0_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break; + case 2: mul_mat_vec_pq2_0_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 3: mul_mat_vec_pq2_0_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 4: mul_mat_vec_pq2_0_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 5: mul_mat_vec_pq2_0_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 6: mul_mat_vec_pq2_0_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 7: mul_mat_vec_pq2_0_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break; + case 8: mul_mat_vec_pq2_0_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 PQ2_0 multi-col MMVQ", ncols_dst); + } +} + static void mul_mat_vec_q2_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, @@ -2415,6 +2475,20 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens mul_mat_vec_ptq1_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); } break; + case GGML_TYPE_PQ2_0: + 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_pq2_0_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); + mul_mat_vec_pq2_0_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) { + GGML_SYCL_DEBUG("Calling mul_mat_vec_pq2_0_q8_1_sycl\n"); + mul_mat_vec_pq2_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + } + break; case GGML_TYPE_Q2_0: if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { const int stride_col_y = src1_padded_col_size / QK8_1; diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 9cf49e9b0a93..de894268a578 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -427,6 +427,32 @@ vec_dot_ptq1_0_q8_1(const void *__restrict__ vbq, return (float) bq->d * acc; } +#define VDR_PQ2_0_Q8_1_MMVQ 1 + +static __dpct_inline__ float +vec_dot_pq2_0_q8_1(const void *__restrict__ vbq, + const block_q8_1 *__restrict__ bq8_1, const int &iqs) { + const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq; + const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; + const float d2 = bq2_0->d; + const uint8_t * qs_chunk = bq2_0->qs + iqs * 8; + + int sumi = 0; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const uint8_t q = qs_chunk[j]; + const uint32_t vi = ((q >> 0) & 0x3) | + (((q >> 2) & 0x3) << 8) | + (((q >> 4) & 0x3) << 16) | + (((q >> 6) & 0x3) << 24); + const int w = byte_sub_4(vi, 0x01010101); + const int u = get_int_from_int8_aligned(bq8_1_chunk->qs, j); + sumi = dpct::dp4a(w, u, sumi); + } + + return d2 * ((const float) bq8_1_chunk->ds[0]) * (float) sumi; +} + // VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called // MMVQ = mul_mat_vec_q, MMQ = mul_mat_q From de7a01bc2a7be72e1fe94257f57ad90b5b5f387f Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Mon, 21 Sep 2026 19:47:46 -0700 Subject: [PATCH 4/7] sycl: spread PQ2_0 MMVQ across 16 lanes per block (~2x tg) The MMVQ kernel derives its lane split from qi/vdr. PQ2_0 was passing QI_PQ2_0 (= QK_PQ2_0/32 = 4), which put only 4 lanes on each 128-element block and left every lane an 8-deep serial dp4a chain over 8 bytes of qs. Pass a dedicated MMVQ_PQ2_0_QI of 16 instead, so a block is covered by 16 lanes reading 2 bytes / 8 elements each with 2 dp4a per call, and a sub-group keeps more blocks in flight. Measured on an Arc iGPU with test-backend-ops perf, MUL_MAT m=4096 n=1 k=14336, alternating the two builds to cancel out background load: 827/846/837 us before, 397/395/384 us after (~2.1x). On an idle machine the same case goes 247.6 -> 123.0 us. 8 and 32 lanes per block were both tried and were slower; at 32 the per-call scale setup dominates the single remaining dp4a. Also fold the 2-bit-to-byte-lane unpack shared with the q2_0 path into unpack_2bit_to_byte_lanes(), and switch the epilogue to the q4_0-style sumi*ds.x - zp*ds.y form. The previous PQ2_0 epilogue applied the zero point with byte_sub_4 and then used only ds[0], dropping the ds.y term. test-backend-ops test -o MUL_MAT -p type_a=pq2_0: 38/38 on both SYCL devices, 3/3 backends passed. The unrelated q1_0 failures on SYCL1 are present with and without this change. --- ggml/src/ggml-sycl/mmvq.cpp | 4 +-- ggml/src/ggml-sycl/vecdotq.hpp | 55 ++++++++++++++++++++-------------- 2 files changed, 34 insertions(+), 25 deletions(-) diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index abc88b7bc5c3..36c4c1778769 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1351,7 +1351,7 @@ static void mul_mat_vec_pq2_0_q8_1_sycl(const void * vx, const void * vy, 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( vx, vy, dst, ncols, nrows, item_ct1); }); @@ -1373,7 +1373,7 @@ static void mul_mat_vec_pq2_0_q8_1_sycl_ncols( 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); }); diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index de894268a578..dcfded35b2a0 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -427,30 +427,46 @@ vec_dot_ptq1_0_q8_1(const void *__restrict__ vbq, return (float) bq->d * acc; } +static __dpct_inline__ uint32_t unpack_2bit_to_byte_lanes(const uint32_t b) { + const uint32_t bits = (b | (b << 12)) & 0x000F000Fu; + return (bits | (bits << 6)) & 0x03030303u; +} + +// PQ2_0 packs 128 elements into 32 bytes. The kernel derives its lane split from qi/vdr, and +// QI_PQ2_0 (= 4) would put only 4 lanes on a block, leaving each one an 8-deep serial dp4a chain. +// MMVQ_PQ2_0_QI spreads a block over 16 lanes instead (2 bytes / 8 elements, 2 dp4a each), which +// measures ~2x faster on Arc iGPUs; 8 and 32 lanes were both slower, the latter because the +// per-call scale setup then dominates a single dp4a. #define VDR_PQ2_0_Q8_1_MMVQ 1 +#define MMVQ_PQ2_0_QI 16 static __dpct_inline__ float vec_dot_pq2_0_q8_1(const void *__restrict__ vbq, const block_q8_1 *__restrict__ bq8_1, const int &iqs) { - const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq; - const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; - const float d2 = bq2_0->d; - const uint8_t * qs_chunk = bq2_0->qs + iqs * 8; + const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq; + const float d2 = bq2_0->d; + + // iqs indexes 2-byte groups of qs, i.e. 8 elements. Four groups share one 32-element q8_1 + // sub-block. + const block_q8_1 * bq8_1_chunk = bq8_1 + (iqs >> 2); + const int u_base = 2 * (iqs & 3); + + uint32_t val = (uint32_t) *(const uint16_t *) (bq2_0->qs + 2 * iqs); int sumi = 0; #pragma unroll - for (int j = 0; j < 8; ++j) { - const uint8_t q = qs_chunk[j]; - const uint32_t vi = ((q >> 0) & 0x3) | - (((q >> 2) & 0x3) << 8) | - (((q >> 4) & 0x3) << 16) | - (((q >> 6) & 0x3) << 24); - const int w = byte_sub_4(vi, 0x01010101); - const int u = get_int_from_int8_aligned(bq8_1_chunk->qs, j); - sumi = dpct::dp4a(w, u, sumi); + for (int l = 0; l < 2; ++l) { + const uint32_t vi = unpack_2bit_to_byte_lanes(val & 0xFFu); + val >>= 8; + const int u = get_int_from_int8_aligned(bq8_1_chunk->qs, u_base + l); + sumi = dpct::dp4a(vi, u, sumi); } - return d2 * ((const float) bq8_1_chunk->ds[0]) * (float) sumi; + const sycl::float2 ds8f = bq8_1_chunk->ds.convert(); + + // The unpack leaves quants as 0..2; the second term subtracts the zero point of 1 from each, + // scaled by the fraction of the q8_1 sub-block this call covers (8 of 32 elements). + return d2 * (sumi * ds8f.x() - 0.25f * ds8f.y()); } // VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called @@ -814,15 +830,8 @@ static __dpct_inline__ float vec_dot_q2_0_q8_1_impl( for (int i = 0; i < vdr; ++i) { #pragma unroll for (int j = 0; j < 4; ++j) { - const uint8_t q = (uint8_t) ((uint32_t) v[i] >> (8 * j)); - - // unpack 2-bit values to byte lanes (0..3), then apply zero-point - // correction with ds8f.y() below, mirroring the q4_0 style. - int vi = 0; - vi |= (((q >> 0) & 0x3) & 0xFF) << 0; - vi |= (((q >> 2) & 0x3) & 0xFF) << 8; - vi |= (((q >> 4) & 0x3) & 0xFF) << 16; - vi |= (((q >> 6) & 0x3) & 0xFF) << 24; + const uint8_t q = (uint8_t) ((uint32_t) v[i] >> (8 * j)); + const uint32_t vi = unpack_2bit_to_byte_lanes(q); sumi = dpct::dp4a(vi, u[4 * i + j], sumi); } From 8a08dd0192209d81e1666ba1d6fb2df1e7cbb8ea Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Tue, 22 Sep 2026 21:14:36 -0700 Subject: [PATCH 5/7] sycl: gate mmvq on implemented types instead of aborting can_use_mul_mat_vec_q() claimed every quantized type via ggml_is_quantized(), so any type without a case in mul_mat_vec_q_switch_type() reached the GGML_ABORT in its default arm and killed the process. test-backend-ops -o MUL_MAT could not complete on SYCL for this reason: it died at mmvq.cpp:2710 on tq2_0, which also meant every case generated after that point never ran. Gate on an explicit list of the types the switch actually implements. The list matches the switch exactly, and is also identical to the set ggml_get_to_fp16_sycl() handles, so a type this declines cannot be served by the dequantize fallback either - hence the matching check in supports_op(), which lets the scheduler place those nodes on CPU rather than abort mid-graph. Arc Pro B50, oneAPI 2026.1, Level Zero: before: test -o MUL_MAT aborts, exit 134 after: 1098/1098 tests passed, exit 0 tq2_0 goes from "fatal error: unsupport data type" to "not supported". --- ggml/src/ggml-sycl/ggml-sycl.cpp | 43 +++++++++++++++++++++++++++++++- 1 file changed, 42 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 5bbb39175e62..1d091a970f78 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3848,6 +3848,39 @@ static bool ggml_sycl_supports_dmmv(enum ggml_type type) { } } +static bool ggml_sycl_supports_mmvq(enum ggml_type type) { + switch (type) { + case GGML_TYPE_Q4_0: + case GGML_TYPE_Q4_1: + case GGML_TYPE_Q5_0: + case GGML_TYPE_Q5_1: + case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_Q2_0: + case GGML_TYPE_Q2_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: + case GGML_TYPE_Q6_K: + case GGML_TYPE_IQ1_S: + case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_XXS: + case GGML_TYPE_IQ2_XS: + case GGML_TYPE_IQ2_S: + case GGML_TYPE_IQ3_XXS: + case GGML_TYPE_IQ3_S: + case GGML_TYPE_IQ4_NL: + case GGML_TYPE_IQ4_XS: + case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: + return true; + default: + return false; + } +} + // Helper functions to unify device memory allocation for both async and sync paths static inline void * sycl_ext_malloc_device(dpct::queue_ptr stream, size_t size) { bool use_async = g_ggml_sycl_use_async_mem_op; @@ -4487,7 +4520,7 @@ static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_ } static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { - return ggml_is_quantized(src0->type) && + return ggml_sycl_supports_mmvq(src0->type) && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; } @@ -6021,6 +6054,14 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons a->ne[0] > 128 && a->ne[2] == 1 && src0_type == GGML_TYPE_F16) { return false; } + + if (ggml_is_quantized(src0_type) && + !ggml_sycl_supports_mmvq(src0_type) && + !ggml_sycl_supports_dmmv(src0_type) && + !ggml_sycl_supports_mmq(src0_type)) { + return false; + } + return true; } case GGML_OP_OUT_PROD: From 450c5c3ebc9bef8e53647aa5e3d26c04c35e77c1 Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Tue, 22 Sep 2026 21:14:42 -0700 Subject: [PATCH 6/7] test-backend-ops: add Bonsai-shape MUL_MAT cases for PTQ1_0/PQ2_0 The stock generator only emits 16 x 256 cases for these types, all of which take the single-column mmvq path, so the multi-column dispatchers (ncols_dst 2..8) had no coverage at all. Add the real Bonsai row lengths (k = 1024/5120/6144/17408) at odd row counts (m = 67/70, landing a partial row group) with n sweeping 1..8. Arc Pro B50, oneAPI 2026.1, Level Zero, per type: MUL_MAT 9 -> 102/102 passed MUL_MAT_ID 75/75 passed --- tests/test-backend-ops.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index b28bc29ef94a..1363837294df 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9301,6 +9301,19 @@ static std::vector> make_test_cases_eval() { 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})); + // Ternary formats at real Bonsai weight shapes. The k values are the model's actual row + // lengths and the odd m values land a partial row group; n sweeps 1..8 so the multi-column + // mmvq dispatchers run, which the 16 x 256 cases above never reach. + for (ggml_type type_a : {GGML_TYPE_PTQ1_0, GGML_TYPE_PQ2_0}) { + for (int64_t k : {1024, 5120, 6144, 17408}) { + for (int64_t m : {67, 70}) { + for (int64_t n = 1; n <= 8; ++n) { + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, m, n, k, {1, 1}, {1, 1})); + } + } + } + } + // m == 1, with n on both sides of MMVF_MAX_BATCH_SIZE (8): mmvf below, operand swap above for (int64_t n : {1, 7, 8, 9, 16, 128, 512}) { test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 1, n, 2048, {1, 1}, {1, 1})); From 42ce76d0196baf2e2a995123dd303e3727b6634a Mon Sep 17 00:00:00 2001 From: Scott J Guyton Date: Wed, 23 Sep 2026 12:35:37 -0700 Subject: [PATCH 7/7] sycl: decline float -> PTQ1_0/PQ2_0 copies instead of asserting Adding the two types to ggml_sycl_is_quantized_type() made supports_op accept f32/f16/bf16 -> PTQ1_0/PQ2_0 copies, but there is no quantize-rows kernel for either, so ggml_sycl_cpy() took the float -> quantized branch and died on GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(...)) at cpy.cpp:1298/1302/1308. Neither format can be produced by a row-at-a-time GPU kernel: the packing assumes the Hadamard rotation the offline converter applies, so the right answer is to decline the pair and let the scheduler place it elsewhere, not to add a kernel. Quantized -> same-quantized copies are unaffected and still take the memcpy path. Reported by @dwymark on an Arc B390 (Panther Lake Xe3); reproduced on an Arc Pro B50. test -b SYCL0 -o CPY -p "ptq1_0|pq2_0" before: abort at cpy.cpp:1298 after: 18/18 tests passed, float -> ternary reports "not supported" No regression: MUL_MAT 204/204, MUL_MAT_ID 150/150, GET_ROWS 8/8 for the two types. The unrelated tq2_0 -> tq2_0 abort at cpy.cpp:1455 and the SET_ROWS abort at set_rows.cpp:549 are both present on prism and untouched here. --- ggml/src/ggml-sycl/ggml-sycl.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 1d091a970f78..761b7f9c664d 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -6146,6 +6146,16 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons ggml_type src0_type = op->src[0]->type; ggml_type src1_type = op->src[1]->type; + // Quantizing a float row into PTQ1_0 or PQ2_0 has no kernel: both are + // produced offline by the converter, which also applies the Hadamard + // rotation the packing assumes. ggml_sycl_cpy() would take the + // float -> quantized branch and assert, so decline the pair here and let + // the scheduler fall back. The quant -> same-quant copies are handled. + if ((src1_type == GGML_TYPE_PTQ1_0 || src1_type == GGML_TYPE_PQ2_0) && + src0_type != src1_type) { + return false; + } + if (src0_type == GGML_TYPE_F16) { if (src1_type == GGML_TYPE_Q2_K || src1_type == GGML_TYPE_Q3_K ||