bri-prism · Oct 14, 2026 · Sep 26, 2026 · Sep 27, 2026 · Sep 27, 2026 · Sep 27, 2026
diff --git a/common/speculative.cpp b/common/speculative.cpp
@@ -3490,6 +3490,15 @@ common_params common_base_params_to_speculative(const common_params & params) {
     result.cache_type_k  = params_spec.cache_type_k;
     result.cache_type_v  = params_spec.cache_type_v;
 
+    // ARC-LAB: the draft context inherits the target's physical batch (2048 on the Arc server = a 784 MiB compute
+    // buffer for an MTP layer that drafts 1-3 tokens); its prompt catch-up already runs in n_ubatch chunks.
+    if (const char * ub = getenv("LLAMA_ARG_SPEC_DRAFT_UBATCH")) {
+        const int32_t n = atoi(ub);
+        if (n > 0 && n < result.n_ubatch) {
+            result.n_ubatch = n;
+        }
+    }
+
     // the K-cache mean-center bias is calibrated for the target model
     // (per-head/channel K layout); the draft model has a different K
     // geometry, so it must not inherit the bias

diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt
@@ -105,6 +105,15 @@ endif()
 
 target_compile_options(ggml-sycl PRIVATE "-Wno-narrowing")
 
+# ARC-LAB: one device image per kernel. The Xe2 XMX kernels (DPAS) then only get JIT-compiled when they are first launched,
+# so a GPU without DPAS (Xe-LP, Gen9) never compiles them; with one image per source file, compiling any kernel of
+# fattn-dec / mmvq / ptq1-t2 built the DPAS ones too and the shader compiler crashed (UHD 630). OFF = compiler default.
+option(GGML_SYCL_DEVICE_CODE_SPLIT_PER_KERNEL "ggml: SYCL device code split per kernel" ON)
+if (GGML_SYCL_DEVICE_CODE_SPLIT_PER_KERNEL)
+    target_compile_options(ggml-sycl PRIVATE "-fsycl-device-code-split=per_kernel")
+    target_link_options(ggml-sycl PRIVATE "-fsycl-device-code-split=per_kernel")
+endif()
+
 message(STATUS "GGML_SYCL_SUPPORT_LEVEL_ZERO_API ${GGML_SYCL_SUPPORT_LEVEL_ZERO_API}")
 if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
     # Link against Level Zero loader for direct device memory allocation.

diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
@@ -60,6 +60,8 @@ void ggml_sycl_host_free(void* ptr);
 
 extern int g_ggml_sycl_debug;
 extern int g_ggml_sycl_enable_optimize;
+extern int g_ggml_sycl_quant_f16_gemm;
+extern int g_ggml_sycl_quant_s4_gemm;
 extern int g_ggml_sycl_enable_fusion;
 extern int g_ggml_sycl_enable_esimd;
 extern int g_ggml_sycl_prioritize_dmmv;
@@ -225,6 +227,7 @@ inline dpct::err0 ggml_sycl_set_device(const int device) try {
 struct optimize_feature {
     bool reorder=false;
     bool xmx_pq2=false; // PQ2_0 rewritten into the XMX layout (pq2_xmx.hpp); only that path can read it
+    bool t2=false;  // ARC-LAB: PTQ1_0 repacked to the TernSYCL 2-bit layout (ptq1-t2.cpp)
 };
 
 struct sycl_device_info {
@@ -1041,4 +1044,10 @@ static __dpct_inline__ float ggml_sycl_ue4m3_to_fp32(uint8_t x) {
     return raw * 0.5f;
 }
 
+// ARC-LAB lab knobs: integer tuning switches read through here can be changed while the process runs. With
+// GGML_SYCL_KNOBS_FILE=<path> set, the file (NAME=VALUE lines) is re-read at the start of a graph compute when its
+// mtime changes and its values override the environment; without it, each knob is read from the environment once.
+int  ggml_sycl_knob(const char * name, int def);
+void ggml_sycl_knobs_refresh();
+
 #endif // GGML_SYCL_COMMON_HPP
diff --git a/ggml/src/ggml-sycl/concat.cpp b/ggml/src/ggml-sycl/concat.cpp
@@ -94,6 +94,23 @@ template <typename T>
 static void concat_T_sycl(const T *x, const T *y, T *dst,
                             int ne00, int ne01, int ne02, int ne0, int ne1,
                             int ne2, int dim, queue_ptr stream) {
+  if (dim == 0 && ne0 < SYCL_CONCAT_BLOCK_SIZE) {
+      // ARC-LAB: narrow rows (delta-net conv state: 3 + 1 columns x 10240 channels) launched one work-group of 256 per
+      // row with 4 active lanes (12.5 us per call on the B580); one lane per element instead.
+      const int total = ne0 * ne1 * ne2;
+      const int nblk  = (total + SYCL_CONCAT_BLOCK_SIZE - 1) / SYCL_CONCAT_BLOCK_SIZE;
+      stream->parallel_for(sycl::nd_range<1>(nblk * SYCL_CONCAT_BLOCK_SIZE, SYCL_CONCAT_BLOCK_SIZE),
+                           [=](sycl::nd_item<1> it) {
+                               const int i = it.get_global_id(0);
+                               if (i >= total) {
+                                   return;
+                               }
+                               const int i0  = i % ne0;
+                               const int row = i / ne0;  // i1 + i2 * ne1
+                               dst[i] = i0 < ne00 ? x[row * ne00 + i0] : y[row * (ne0 - ne00) + i0 - ne00];
+                           });
+      return;
+  }
   int num_blocks = (ne0 + SYCL_CONCAT_BLOCK_SIZE - 1) / SYCL_CONCAT_BLOCK_SIZE;
   sycl::range<3> gridDim(ne2, ne1, num_blocks);
   switch (dim) {

diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp
@@ -170,6 +170,35 @@ static void dequantize_row_q4_0_sycl_reorder(const void *vx, dst_t *y, const int
 
 }
 
+template <typename dst_t>
+static void dequantize_row_pq2_0_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
+    GGML_ASSERT(k % QK_PQ2_0 == 0);
+    const uint8_t *    qs = (const uint8_t *) vx;
+    const sycl::half * d  = (const sycl::half *) (qs + k / 4);
+    stream->parallel_for(sycl::range<1>(k / 4), [=](sycl::id<1> i) {
+        const float   dv = d[i / (QK_PQ2_0 / 4)];
+        const uint8_t q  = qs[i];
+#pragma unroll
+        for (int j = 0; j < 4; ++j) {
+            y[i * 4 + j] = (float) (((q >> (2 * j)) & 3) - 1) * dv;
+        }
+    });
+}
+
+template <typename dst_t>
+static void dequantize_row_ptq1_0_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
+    GGML_ASSERT(k % QK_PTQ1_0 == 0);
+    const int64_t      nb = k / QK_PTQ1_0;
+    const uint8_t *    qs = (const uint8_t *) vx;
+    const uint8_t *    qh = qs + nb * 24;
+    const sycl::half * d  = (const sycl::half *) (qs + nb * 26);
+    stream->parallel_for(sycl::range<1>(k), [=](sycl::id<1> id) {
+        const int64_t i  = id[0];
+        const int64_t ib = i / QK_PTQ1_0;
+        y[i] = (float) (ptq1_0_trit(qs + ib * 24, qh + ib * 2, (int) (i % QK_PTQ1_0)) - 1) * (float) d[ib];
+    });
+}
+
 template <typename dst_t>
 static void dequantize_row_q8_0_sycl_reorder(const void *vx, dst_t *y, const int64_t k,
                                      dpct::queue_ptr stream) {
@@ -657,12 +686,20 @@ 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<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_PTQ1_0:
+            if (dst->src[0]->extra &&
+                ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
+                return dequantize_row_ptq1_0_sycl_reorder;
+            }
             return dequantize_block_sycl<QK_PTQ1_0, QR_PTQ1_0, dequantize_ptq1_0>;
         case GGML_TYPE_PQ2_0:
+            if (dst->src[0]->extra &&
+                ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
+                return dequantize_row_pq2_0_sycl_reorder;
+            }
             return dequantize_block_sycl<QK_PQ2_0, QR_PQ2_0, dequantize_pq2_0>;
-        case GGML_TYPE_Q2_0:
-            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             if (dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
@@ -751,12 +788,20 @@ 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<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_PTQ1_0:
+            if (dst->src[0]->extra &&
+                ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
+                return dequantize_row_ptq1_0_sycl_reorder;
+            }
             return dequantize_block_sycl<QK_PTQ1_0, QR_PTQ1_0, dequantize_ptq1_0>;
         case GGML_TYPE_PQ2_0:
+            if (dst->src[0]->extra &&
+                ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
+                return dequantize_row_pq2_0_sycl_reorder;
+            }
             return dequantize_block_sycl<QK_PQ2_0, QR_PQ2_0, dequantize_pq2_0>;
-        case GGML_TYPE_Q2_0:
-            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             if (dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
@@ -887,3 +932,68 @@ to_fp16_nc_sycl_t ggml_get_to_fp16_nc_sycl(ggml_type type) {
             return nullptr;
     }
 }
+
+void ggml_sycl_quant_to_s4(ggml_type type, bool reordered, const void * vx, uint8_t * w4, float * scales,
+                           int64_t nrows, int64_t ncols, dpct::queue_ptr stream) {
+    GGML_ASSERT(type == GGML_TYPE_PQ2_0 || type == GGML_TYPE_PTQ1_0);
+    GGML_ASSERT(ncols % 128 == 0);
+    const int64_t   nb   = nrows * ncols / 128;
+    const uint8_t * base = (const uint8_t *) vx;
+    const bool      ptq  = type == GGML_TYPE_PTQ1_0;
+    // one work item = 8 elements = one uint32 of s4
+    stream->parallel_for(sycl::range<1>(nb * 16), [=](sycl::id<1> id) {
+        const int64_t ib = id[0] / 16;
+        const int     c  = (int) (id[0] % 16);
+        int           v[8];
+        sycl::half    d;
+        if (ptq) {
+            const uint8_t * qs;
+            const uint8_t * qh;
+            if (reordered) {
+                qs = base + ib * 24;
+                qh = base + nb * 24 + ib * 2;
+                d  = ((const sycl::half *) (base + nb * 26))[ib];
+            } else {
+                const block_ptq1_0 * x = (const block_ptq1_0 *) vx + ib;
+                qs = x->qs;
+                qh = x->qh;
+                d  = x->d;
+            }
+            // trit t of byte b = ((b * 3^t) mod 256) * 3 >> 8
+            constexpr uint32_t pow3[5] = { 1, 3, 9, 27, 81 };
+            if (c < 10) {         // elements t*16 + m, 8 consecutive m
+                const uint32_t p = pow3[c / 2];
+                const int      m = (c % 2) * 8;
+#pragma unroll
+                for (int i = 0; i < 8; ++i) v[i] = (int) ((((qs[m + i] * p) & 0xFF) * 3) >> 8) - 1;
+            } else if (c < 15) {  // elements 80 + t*8 + m, m = 0..7
+                const uint32_t p = pow3[c - 10];
+#pragma unroll
+                for (int i = 0; i < 8; ++i) v[i] = (int) ((((qs[16 + i] * p) & 0xFF) * 3) >> 8) - 1;
+            } else {              // elements 120 + t*2 + h
+#pragma unroll
+                for (int i = 0; i < 8; ++i) v[i] = (int) ((((qh[i % 2] * pow3[i / 2]) & 0xFF) * 3) >> 8) - 1;
+            }
+        } else {
+            const uint8_t * qs;
+            if (reordered) {
+                qs = base + ib * 32;
+                d  = ((const sycl::half *) (base + nb * 32))[ib];
+            } else {
+                const block_pq2_0 * x = (const block_pq2_0 *) vx + ib;
+                qs = x->qs;
+                d  = x->d;
+            }
+#pragma unroll
+            for (int i = 0; i < 8; ++i) v[i] = ((qs[2 * c + i / 4] >> (2 * (i % 4))) & 3) - 1;
+        }
+        uint32_t out = 0;
+#pragma unroll
+        for (int i = 0; i < 8; ++i) out |= (uint32_t) (v[i] & 0xF) << (4 * i);
+        ((uint32_t *) w4)[id[0]] = out;
+        if (c == 0) {
+            const int64_t bpr = ncols / 128;
+            scales[(ib % bpr) * nrows + ib / bpr] = d;  // [group][row]
+        }
+    });
+}
diff --git a/ggml/src/ggml-sycl/convert.hpp b/ggml/src/ggml-sycl/convert.hpp
@@ -21,6 +21,11 @@ typedef to_t_sycl_t<float>      to_fp32_sycl_t;
 typedef to_t_sycl_t<sycl::half> to_fp16_sycl_t;
 
 to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst);
+
+// PQ2_0 / PTQ1_0 rows -> packed s4 (low nibble = even element) + f32 scales laid out [group][row].
+// All their values (-1..2) are exact in s4.
+void ggml_sycl_quant_to_s4(ggml_type type, bool reordered, const void * vx, uint8_t * w4, float * scales,
+                           int64_t nrows, int64_t ncols, dpct::queue_ptr stream);
 to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor * dst);
 
 #ifdef GGML_SYCL_HAS_BF16

diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp
@@ -47,6 +47,26 @@ static __dpct_inline__ void dequantize_q2_0(const void *vx, const int64_t ib,
 #endif // GGML_SYCL_F16
 }
 
+// Trit of element e (0..127) in a PTQ1_0 block: 16 bytes x 5, 8 bytes x 5, then 2 qh bytes x 4.
+static __dpct_inline__ int ptq1_0_trit(const uint8_t * qs, const uint8_t * qh, const int e) {
+    uint32_t v;
+    int      t;
+    if (e < 80) {
+        v = qs[e % 16];
+        t = e / 16;
+    } else if (e < 120) {
+        v = qs[16 + (e - 80) % 8];
+        t = (e - 80) / 8;
+    } else {
+        v = qh[(e - 120) % 2];
+        t = (e - 120) / 2;
+    }
+    for (int i = 0; i < t; ++i) {
+        v = (v * 3) & 0xFF;
+    }
+    return (int) ((v * 3) >> 8);
+}
+
 static __dpct_inline__ void dequantize_q4_0(const void *vx, const int64_t ib,
                                             const int iqs, dfloat2 &v) {
     const block_q4_0 * x = (const block_q4_0 *) vx;

diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp
@@ -238,6 +238,34 @@ static void convert_mul_mat_vec_f16_sycl(const void *vx, const dfloat *y,
 }
 
 #ifdef GGML_SYCL_DMMV_HAS_BF16
+// ARC-LAB: small-output bf16 mat-vec (Bonsai's delta-net gate projections: 48 x 5120, 96 calls per token).
+// The generic dequantize_mul_mat_vec gives each row one sub-group, so 48 rows used 48 sub-groups and reached
+// ~25 GB/s on the B580 (~20 us per call). Here each row gets a 256-wide work-group that splits K, loads 8 bf16
+// at a time and reduces over the group.
+static void bf16_gemv_rowsplit_sycl(const void * vx, const dfloat * y, float * dst, const int ncols,
+                                    const int nrows, dpct::queue_ptr stream) {
+    constexpr int WG = 256;
+    stream->parallel_for(
+        sycl::nd_range<1>(sycl::range<1>((size_t) nrows * WG), sycl::range<1>(WG)),
+        [=](sycl::nd_item<1> it) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+            const int row = it.get_group(0);
+            const int tid = it.get_local_id(0);
+            const uint16_t * x = (const uint16_t *) vx + (size_t) row * ncols;
+            float sum = 0.0f;
+            for (int i = tid * 8; i < ncols; i += WG * 8) {
+                const sycl::vec<uint16_t, 8> w = *(const sycl::vec<uint16_t, 8> *) (x + i);
+#pragma unroll
+                for (int j = 0; j < 8; ++j) {
+                    sum += sycl::bit_cast<float>((uint32_t) w[j] << 16) * (float) y[i + j];
+                }
+            }
+            sum = sycl::reduce_over_group(it.get_group(), sum, sycl::plus<float>());
+            if (tid == 0) {
+                dst[row] = sum;
+            }
+        });
+}
+
 static void convert_mul_mat_vec_bf16_sycl(const void *vx, const dfloat *y,
                                           float *dst, const int ncols,
                                           const int nrows,
@@ -2209,9 +2237,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
             convert_mul_mat_vec_f16_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             break;
 #ifdef GGML_SYCL_DMMV_HAS_BF16
-        case GGML_TYPE_BF16:
-            convert_mul_mat_vec_bf16_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+        case GGML_TYPE_BF16: {
+            static const bool rowsplit = getenv("GGML_SYCL_BF16_ROWSPLIT_OFF") == nullptr;
+            if (rowsplit && row_diff <= 512 && ne00 % 8 == 0 && ((uintptr_t) src0_dd_i % 16) == 0) {
+                bf16_gemv_rowsplit_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+            } else {
+                convert_mul_mat_vec_bf16_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+            }
             break;
+        }
 #endif
         default:
             printf("ggml_sycl_op_dequantize_mul_mat_vec unsupported GGML_TYPE %d\n", src0->type);