sycl: add Prism Q1_0 and Q2_0 g128 support

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<src_t>(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 <typename dst_t>
+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<const block_q1_0 *>(vx);
+            const int64_t      ib  = i / QK1_0;
+            const int          bit = i % QK1_0;
+            const float        d   = static_cast<float>(x[ib].d);
+            y[i] = static_cast<dst_t>(((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 <typename dst_t>
+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<const block_q2_0 *>(vx);
+            const int64_t ib     = i / QK2_0;
+            const int local      = static_cast<int>(i % QK2_0);
+            const int code       = (x[ib].qs[local / 4] >> ((local % 4) * 2)) & 0x3;
+            const float d        = static_cast<float>(x[ib].d);
+            y[i] = static_cast<dst_t>((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) {

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];

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<QK1_0, 1, dequantize_q1_0>(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<QK2_0, 1, dequantize_q2_0>(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<QK_MXFP4, 2, dequantize_mxfp4>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());

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:

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<QK1_0, QI1_0, block_q1_0, VDR_Q1_0_Q8_1_MMVQ, vec_dot_q1_0_q8_1>(
+                    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<QK2_0, QI2_0, block_q2_0, VDR_Q2_0_Q8_1_MMVQ, vec_dot_q2_0_q8_1>(
+                    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) {

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<const block_q1_0 *>(vbq);
+    const float d_q1 = static_cast<float>(bq->d);
+    const sycl::float2 ds8 = bq8_1[iqs].ds.convert<float, sycl::rounding_mode::automatic>();
+    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<float>(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<const block_q2_0 *>(vbq);
+    const float d_q2 = static_cast<float>(bq->d);
+    const sycl::float2 ds8 = bq8_1[iqs].ds.convert<float, sycl::rounding_mode::automatic>();
+    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<int>(q8[j]);
+    }
+
+    return d_q2 * (d_q8 * static_cast<float>(sumi) - sum_q8_real);
+}
+
 #endif // GGML_SYCL_VECDOTQ_HPP