diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 65593402e7de..0e9dde7e662a 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -639,9 +639,59 @@ static void convert_unary_sycl(const void * vx, dst_t * y, const int64_t k, dpct convert_unary_nc_sycl(vx, y, k, 1, 1, 1, k, k, k, queue); } +// Q1_0 uses the official PrismML 128-element block layout. +// Bit 1 maps to +d and bit 0 maps to -d. +template +static void dequantize_row_q1_0_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) { + GGML_ASSERT(k % QK1_0 == 0); + const int64_t num_threads = SYCL_DEQUANTIZE_BLOCK_SIZE; + const int64_t num_blocks = (k + num_threads - 1) / num_threads; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, num_threads), + sycl::range<3>(1, 1, num_threads)), + [=](sycl::nd_item<3> item_ct1) { + const int64_t i = item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2); + if (i >= k) { + return; + } + const block_q1_0 * x = static_cast(vx); + const int64_t ib = i / QK1_0; + const int bit = i % QK1_0; + const float d = static_cast(x[ib].d); + y[i] = static_cast(((x[ib].qs[bit / 8] >> (bit % 8)) & 1) ? d : -d); + }); +} + +// Q2_0 uses one FP16 scale for 128 values and four 2-bit codes per byte. +// Stored code 0..3 maps to -1, 0, +1, +2 multiplied by the block scale. +template +static void dequantize_row_q2_0_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) { + GGML_ASSERT(k % QK2_0 == 0); + const int64_t num_threads = SYCL_DEQUANTIZE_BLOCK_SIZE; + const int64_t num_blocks = (k + num_threads - 1) / num_threads; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, num_threads), + sycl::range<3>(1, 1, num_threads)), + [=](sycl::nd_item<3> item_ct1) { + const int64_t i = item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2); + if (i >= k) { + return; + } + const block_q2_0 * x = static_cast(vx); + const int64_t ib = i / QK2_0; + const int local = static_cast(i % QK2_0); + const int code = (x[ib].qs[local / 4] >> ((local % 4) * 2)) & 0x3; + const float d = static_cast(x[ib].d); + y[i] = static_cast((code - 1) * d); + }); +} to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { switch (type) { + case GGML_TYPE_Q1_0: + return dequantize_row_q1_0_sycl; + case GGML_TYPE_Q2_0: + return dequantize_row_q2_0_sycl; case GGML_TYPE_Q4_0: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) { @@ -724,6 +774,10 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { switch (type) { + case GGML_TYPE_Q1_0: + return dequantize_row_q1_0_sycl; + case GGML_TYPE_Q2_0: + return dequantize_row_q2_0_sycl; case GGML_TYPE_Q4_0: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) { diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index ca8cd96c08c3..0966141ea9b4 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -296,6 +296,21 @@ static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib, v.y() = (2 * bit_1 - 1) * d; } +static __dpct_inline__ void dequantize_q2_0(const void *vx, const int64_t ib, + const int iqs, dfloat2 &v) { + const block_q2_0 * x = (const block_q2_0 *) vx; + const dfloat d = x[ib].d; + + auto dequantize_one = [&](const int idx) -> dfloat { + const uint8_t packed = x[ib].qs[idx / 4]; + const int code = (packed >> (2 * (idx % 4))) & 0x3; + return (code - 1) * d; + }; + + v.x() = dequantize_one(iqs + 0); + v.y() = dequantize_one(iqs + 1); +} + static __dpct_inline__ void dequantize_nvfp4(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { const block_nvfp4 & xb = ((const block_nvfp4 *) vx)[ib]; diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 298f247f84e0..86caa253d9be 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -199,6 +199,10 @@ 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_Q2_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 9e9242ba3b1c..1ff74f2f8949 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5272,11 +5272,6 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g struct ggml_tensor * a = op->src[0]; struct ggml_tensor * b = op->src[1]; - // disable Q1_0 until implementation - if (a->type == GGML_TYPE_Q1_0 || b->type == GGML_TYPE_Q1_0) { - return false; - } - if (a->ne[3] != b->ne[3]) { return false; } @@ -5308,6 +5303,7 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_BF16: case GGML_TYPE_F32: case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: case GGML_TYPE_IQ2_XXS: diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index cf2b59576aa1..fe850dd415ff 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -715,6 +715,37 @@ static void reorder_mul_mat_vec_q4_0_q8_1_sycl_switch_ncols( } } +static void mul_mat_vec_q1_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 % QK1_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); + }); + }); +} +static void mul_mat_vec_q2_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 % QK2_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); + }); + }); +} + static void mul_mat_vec_q4_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 % QK4_0 == 0); @@ -2029,6 +2060,14 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens const char * src1_ddq_i_bs = src1_ddq_i + src1_ddq_i_offset; float * dst_dd_i_bs = dst_dd_i + i * dst->ne[0]; switch (src0->type) { + case GGML_TYPE_Q1_0: + GGML_SYCL_DEBUG("Calling mul_mat_vec_q1_0_q8_1_sycl\n"); + 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_Q2_0: + GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_0_q8_1_sycl\n"); + mul_mat_vec_q2_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + break; case GGML_TYPE_Q4_0: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 4b58b09ab2c5..7f32a17c6436 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -1533,4 +1533,55 @@ vec_dot_iq4_xs_q8_1(const void *__restrict__ vbq, #endif } +#define VDR_Q1_0_Q8_1_MMVQ 1 + +static __dpct_inline__ float vec_dot_q1_0_q8_1( + const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + // Official PrismML Q1_0: one 128-element weight block and four + // aligned 32-element Q8_1 activation blocks. iqs selects one chunk. + const block_q1_0 * bq = static_cast(vbq); + const float d_q1 = static_cast(bq->d); + const sycl::float2 ds8 = bq8_1[iqs].ds.convert(); + const float d_q8 = ds8.x(); + const int8_t * q8 = bq8_1[iqs].qs; + const int base_bit = iqs * QK8_1; + + float sum = 0.0f; +#pragma unroll + for (int bit = 0; bit < QK8_1; ++bit) { + const int absolute_bit = base_bit + bit; + const float sign = ((bq->qs[absolute_bit / 8] >> (absolute_bit % 8)) & 1) ? 1.0f : -1.0f; + sum += sign * static_cast(q8[bit]); + } + return d_q1 * d_q8 * sum; +} + +#define VDR_Q2_0_Q8_1_MMVQ 1 + +static __dpct_inline__ float vec_dot_q2_0_q8_1( + const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & iqs) { + // Official PrismML Q2_0: one 128-element weight block and four + // aligned 32-element Q8_1 activation blocks. iqs selects one chunk. + const block_q2_0 * bq = static_cast(vbq); + const float d_q2 = static_cast(bq->d); + const sycl::float2 ds8 = bq8_1[iqs].ds.convert(); + const float d_q8 = ds8.x(); + const float sum_q8_real = ds8.y(); + const int8_t * q8 = bq8_1[iqs].qs; + + int sumi = 0; +#pragma unroll + for (int j = 0; j < QK8_1; ++j) { + const int absolute = iqs * QK8_1 + j; + const int code = (bq->qs[absolute / 4] >> ((absolute % 4) * 2)) & 0x3; + sumi += code * static_cast(q8[j]); + } + + return d_q2 * (d_q8 * static_cast(sumi) - sum_q8_real); +} + #endif // GGML_SYCL_VECDOTQ_HPP