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..761b7f9c664d 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; @@ -3846,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; @@ -4485,7 +4520,8 @@ 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_sycl_supports_mmvq(src0->type) && + src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; } @@ -6018,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: @@ -6034,6 +6078,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: @@ -6100,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 || @@ -6159,6 +6215,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..36c4c1778769 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1278,6 +1278,126 @@ 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_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, @@ -2341,6 +2461,34 @@ 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_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 3ad4cee93a14..dcfded35b2a0 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -344,6 +344,131 @@ 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 }; + + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes +#pragma unroll + 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_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 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_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 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; +#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; +} + +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 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 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); + } + + 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 // MMVQ = mul_mat_vec_q, MMQ = mul_mat_q @@ -705,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); } 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}));