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.

Author: kiljoy001

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<QK_PQ2_0, QI_PQ2_0, block_pq2_0,
+                mul_mat_vec_q<QK_PQ2_0, MMVQ_PQ2_0_QI, block_pq2_0,
                               VDR_PQ2_0_Q8_1_MMVQ, vec_dot_pq2_0_q8_1>(
                     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<QK_PQ2_0, QI_PQ2_0, block_pq2_0,
+                mul_mat_vec_q_ncols<QK_PQ2_0, MMVQ_PQ2_0_QI, block_pq2_0,
                                     VDR_PQ2_0_Q8_1_MMVQ, vec_dot_pq2_0_q8_1, ncols_dst>(
                     vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, item_ct1);
             });

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<float, sycl::rounding_mode::automatic>();
+
+    // 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);
         }