server: Substitute rerank placeholder from the original template only. (#30864)

* server: realtime API: send complete after committed audio drains

* server: realtime API: complete a commit at its end row

Completion is triggered when the stream runs dry, which is a timing
proxy: audio appended while the tail drains leaks into the committed
item, a final commit before the drain merges both items, and back to
back commits collapse into one.

mtmd_audio_stream_flush() records the row count at the end of the
flushed audio (mtmd_audio_stream_get_n_rows_flushed), next to the rows
produced so far (mtmd_audio_stream_get_n_rows). The slot keeps a FIFO
of these end rows and completes each commit once every row up to it is
decoded.

---------

Co-authored-by: Pascal <admin@serveurperso.com>

conversion/__init__.py

@@ -196,6 +196,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
     "MiniCPMForCausalLM": "minicpm",
     "MiniCPMV4_6ForConditionalGeneration": "minicpm",
     "MiniCPMV4_7ForConditionalGeneration": "minicpm",
+    "MiniCPMOcrForConditionalGeneration": "minicpm",
     "MiniMaxText01ForCausalLM": "minimax",
     "MiniMaxM1ForCausalLM": "minimax",
     "MiniMaxM2ForCausalLM": "minimax",
@@ -361,6 +362,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
     "MiniMaxM3SparseForConditionalGeneration": "minimax",
     "MiniCPMV4_6ForConditionalGeneration": "minicpm",
     "MiniCPMV4_7ForConditionalGeneration": "minicpm",
+    "MiniCPMOcrForConditionalGeneration": "minicpm",
     "Mistral3ForConditionalGeneration": "llava",
     "Mistral3Model": "llava",
     "NemotronH_Omni_Reasoning_V3": "nemotron",

conversion/base.py

@@ -1837,6 +1837,9 @@ class TextModel(ModelBase):
         if chkhsh == "1431a23e583c97432bc230bff598d103ddb5a1f89960c8f1d1051aaa944d0b35":
             # ref: https://huggingface.co/sapienzanlp/Minerva-7B-base-v1.0
             res = "minerva-7b"
+        if chkhsh == "6e4d72c6e5260b90606e26a0931fa15e4b42ec6e97719d12919f2b1f93478262":
+            # ref: https://huggingface.co/domyn/Domyn-Small-v1.0
+            res = "domyn-small"
         if chkhsh == "7e57df22b1fe23a7b1e1c7f3dc4e3f96d43a4eb0836d0c6bdc3436d7b2f1c664":
             # ref: https://huggingface.co/tencent/Hunyuan-A13B-Instruct
             res = "hunyuan"

conversion/minicpm.py

@@ -272,3 +272,18 @@ class MiniCPMV4_7VisionModel(MiniCPMV4_6VisionModel):
             yield self.format_tensor_name(gguf.MODEL_TENSOR.V_TOK_EMBD_SEP, suffix=""), data_torch[ids]
             return
         yield from super().modify_tensors(data_torch, name, bid)
+
+
+# MiniCPM-OCR 1.0: the 4.6 stack with model_type minicpm_ocr.
+
+@ModelBase.register("MiniCPMOcrForConditionalGeneration")
+@ModelBase.example("openbmb/MiniCPM-OCR-1.0")
+class MiniCPMOCR1_0TextModel(MiniCPMV4_6TextModel):
+    model_arch = gguf.MODEL_ARCH.QWEN35
+    no_mtp = True
+
+
+@ModelBase.register("MiniCPMOcrForConditionalGeneration")
+@ModelBase.example("openbmb/MiniCPM-OCR-1.0")
+class MiniCPMOCR1_0VisionModel(MiniCPMV4_6VisionModel):
+    projector_type = gguf.VisionProjectorType.MINICPM_OCR_1_0

conversion/nemotron.py

@@ -1,5 +1,6 @@
 from __future__ import annotations
 
+import re
 from typing import Any, Callable, Iterable, Sequence, TYPE_CHECKING
 
 import numpy as np
@@ -367,9 +368,21 @@ class NemotronModel(TextModel):
     model_arch = gguf.MODEL_ARCH.NEMOTRON
 
     def set_vocab(self):
-        self._set_vocab_sentencepiece()
-        self.gguf_writer.add_pad_token_id(0)
-        self.gguf_writer.add_unk_token_id(1)
+        if (self.dir_model / "tokenizer.model").exists():
+            self._set_vocab_sentencepiece()
+            self.gguf_writer.add_pad_token_id(0)
+            self.gguf_writer.add_unk_token_id(1)
+        else:
+            # SPM-style BPE (Metaspace + byte fallback): mark <0xXX> tokens as BYTE so they decode to raw bytes
+            tokens, toktypes, tokpre = self.get_vocab_base()
+            toktypes = [gguf.TokenType.BYTE if re.fullmatch(r"<0x[0-9A-F]{2}>", tok) else typ for tok, typ in zip(tokens, toktypes)]
+            self.gguf_writer.add_tokenizer_model("gpt2")
+            self.gguf_writer.add_tokenizer_pre(tokpre)
+            self.gguf_writer.add_token_list(tokens)
+            self.gguf_writer.add_token_types(toktypes)
+
+            special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
+            special_vocab.add_to_gguf(self.gguf_writer)
 
     def set_gguf_parameters(self):
         super().set_gguf_parameters()
@@ -383,7 +396,10 @@ class NemotronModel(TextModel):
         rot_pct = self.rope_parameters["partial_rotary_factor"]
         n_embd = self.find_hparam(["hidden_size", "n_embd"])
         n_head = self.find_hparam(["num_attention_heads", "n_head"])
-        self.gguf_writer.add_rope_dimension_count(int(rot_pct * n_embd) // n_head)
+        # Use explicit head_dim when available; fallback keeps standard Nemotron working.
+        # Without this, models where n_head * head_dim != n_embd (e.g. 48 heads, head_dim=128) produce the wrong rope dimension count.
+        head_dim = self.hparams.get("head_dim") or n_embd // n_head
+        self.gguf_writer.add_rope_dimension_count(int(rot_pct * head_dim))
 
         # * RopeScaling for Nemotron
         factor = self.hparams.get("factor") or self.rope_parameters.get("factor")

convert_hf_to_gguf_update.py

@@ -178,6 +178,7 @@ pre_computed_hashes = [
     {"name": "glm4", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/zai-org/GLM-4.5-Air", "chkhsh": "9ca2dd618e8afaf09731a7cf6e2105b373ba6a1821559f258b272fe83e6eb902"},
     {"name": "glm4", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/zai-org/GLM-4.7-Flash", "chkhsh": "cdf5f35325780597efd76153d4d1c16778f766173908894c04afc20108536267"},
     {"name": "minerva-7b", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/sapienzanlp/Minerva-7B-base-v1.0", "chkhsh": "1431a23e583c97432bc230bff598d103ddb5a1f89960c8f1d1051aaa944d0b35"},
+    {"name": "domyn-small", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/domyn/Domyn-Small-v1.0", "chkhsh": "6e4d72c6e5260b90606e26a0931fa15e4b42ec6e97719d12919f2b1f93478262"},
     {"name": "hunyuan", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/tencent/Hunyuan-A13B-Instruct", "chkhsh": "7e57df22b1fe23a7b1e1c7f3dc4e3f96d43a4eb0836d0c6bdc3436d7b2f1c664"},
     {"name": "hunyuan-dense", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/tencent/Hunyuan-4B-Instruct", "chkhsh": "bba3b3366b646dbdded5dbc42d59598b849371afc42f7beafa914afaa5b70aa6"},
     {"name": "hy_v4", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/tencent/Hy4-preview", "chkhsh": "e6ddf9c6686791c12d698d34c31ab9be1fea9af5a3d9a6909783ab382198ae1c"},

docs/backend/SYCL.md

@@ -819,6 +819,9 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
 | GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. |
 | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute. Unsupported types and layouts fall back to the standalone op kernels. See `ggml_sycl_can_fuse()`. |
 | GGML_SYCL_ENABLE_ESIMD | 0 or 1 (default)| Enable ESIMD kernels when available. |
+| GGML_SYCL_ENABLE_XMX_ESIMD | 0 or 1 (default) | Use the XMX (DPAS) multi-column mat-vec kernels for reordered Q2_K to Q6_K and Q8_0 weights on Xe2 GPUs. Speeds up speculative decoding and short prompt batches. |
+| GGML_SYCL_XMX_MAX_COLS | non-negative integer, 80 (default) | Most activation columns (batch size) the `GGML_SYCL_ENABLE_XMX_ESIMD` kernels take; wider batches use the other mat-mul kernels. A larger value keeps XMX for bigger batches (as repeated launches), `0` never uses it. |
+| GGML_SYCL_XMX_GLU_MAX_COLS | 0 to 16, 16 (default) | Most columns the fused gate and up XMX kernel takes; wider batches run gate and up as separate mat-vecs. 16 is the most the kernel is built for, larger values are clamped to 16 and `0` turns the fused kernel off. |
 | GGML_SYCL_XMX_GATHER_TYPES | decimal bitmask, all bits set (default) | Weight formats that may use the XMX dequant-GEMM paths, which dequantize weights straight into the XMX tiles. This speeds up prompt processing of MoE models on GPUs with XMX units (Arc A- and B-series, Arc Pro, Data Center GPU Max), for example pp512 of Qwen3-30B-A3B UD-IQ3_XXS by about 50% on an Arc Pro B60. Bits:<br>* 1: IQ4_NL, 2: IQ3_S, 4: IQ4_XS, 8: IQ3_XXS, 16: IQ2_XXS, 32: IQ2_XS, 64: IQ2_S, 128: IQ1_S, 256: IQ1_M<br>* 512: Q8_0, 1024: Q4_K, 2048: Q5_K, 4096: Q6_K (MoE `MUL_MAT_ID` only)<br>Add values to combine them, for example `3` for IQ4_NL and IQ3_S; `0` disables the paths. A set bit does not force the path: batches of more than 64 tokens per expert or row lengths that are not a multiple of 256 (32 for IQ4_NL and Q8_0) use the library GEMM. |
 | GGML_SYCL_XMX_GATHER_SHAPES | decimal bitmask, 255 (default) | XMX `joint_matrix` combinations the paths of `GGML_SYCL_XMX_GATHER_TYPES` may use; the operand type comes from `GGML_SYCL_DYNAMIC_PRECISION` and the best supported combination is picked automatically (logged as `fg_pick_combo`). Bits:<br>* Xe2, Xe3, Xe-HPC: 1: f16 8x16x16, 2: f16 16x16x16, 4: f16 32x64x16, 8: f16 32x64x32, 32: tf32 8x16x8, 64: bf16 8x16x16<br>* Xe-HPG (Arc A770, ARL-H): 16: f16 8x8x16, 128: bf16 8x8x16<br>Clear a bit to exclude a combination, or set a single bit to force one for testing. |
 | GGML_SYCL_DYNAMIC_PRECISION | `F16` (default with `GGML_SYCL_F16=ON`), `BF16`, `TF32` or `F32` (default otherwise) | Operand type of the XMX dequant-GEMM paths (`GGML_SYCL_XMX_GATHER_TYPES`); accumulation is always f32. `F16` is the fastest, but activations above 65504 overflow. `BF16` keeps the f32 range at a 7-bit mantissa, `TF32` keeps the range and the f16 mantissa but is about 30% slower and needs Xe2, Xe3 or Xe-HPC, and `F32` turns the XMX paths off. Ops that request a higher src1 precision ([TAG_GGML_PREC]) get it regardless of this setting. |

ggml/src/ggml-opencl/CMakeLists.txt

@@ -241,6 +241,7 @@ set(GGML_OPENCL_KERNELS
     mul_mm_f16_f32_kq_kqv
     conv2d
     conv2d_f16_f32
+    conv2_dw
     flash_attn_repack
     flash_attn_pre_f16
     flash_attn_f32_f16

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -1005,6 +1005,7 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_conv_2d_f16;
     cl_kernel kernel_conv_2d_f32;
     cl_kernel kernel_conv_2d_f16_f32;
+    cl_kernel kernel_conv2d_dw_f32_f32, kernel_conv2d_dwf16_f32;
     cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4;
     // [size_idx][kda][tgpp] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1.
     // tgpp 0 = TG variant (COLS_PER_LANE_GROUP=1), tgpp 1 = prefill variant (COLS_PER_LANE_GROUP=4).
@@ -3761,6 +3762,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
                 }
     }
 
+    // conv2_dw
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "conv2_dw.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("conv2_dw.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_conv2d_dw_f32_f32 = clCreateKernel(prog, "kernel_conv2d_dw_f32_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_conv2d_dwf16_f32 = clCreateKernel(prog, "kernel_conv2d_dwf16_f32", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
     // ssm_conv
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -9572,6 +9590,11 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             return (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16) ||
                    (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) ||
                    (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
+        case GGML_OP_CONV_2D_DW:
+            return op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
+                   op->src[0]->nb[0] == ggml_type_size(op->src[0]->type) &&
+                   op->src[1]->nb[0] == sizeof(float) && op->nb[0] == sizeof(float) &&
+                   (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16);
         case GGML_OP_SSM_CONV:
             return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
         case GGML_OP_SSM_SCAN: {
@@ -19715,6 +19738,90 @@ static void ggml_cl_conv_2d(ggml_backend_t backend, const ggml_tensor * src0, co
     backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size, local_work_size, dst);
 }
 
+static void ggml_cl_conv2_dw(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    GGML_TENSOR_BINARY_OP_LOCALS;
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *) src1->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    const cl_long offset0 = extra0->offset + src0->view_offs;
+    const cl_long offset1 = extra1->offset + src1->view_offs;
+    const cl_long offsetd = extrad->offset + dst->view_offs;
+
+    const cl_long C  = ne12;
+    const cl_long N  = ne13;
+    const cl_long KW = ne00;
+    const cl_long KH = ne01;
+    const cl_long W  = ne10;
+    const cl_long H  = ne11;
+    const cl_long OW = ne0;
+    const cl_long OH = ne1;
+
+    const cl_int s0 = dst->op_params[0];
+    const cl_int s1 = dst->op_params[1];
+    const cl_int p0 = dst->op_params[2];
+    const cl_int p1 = dst->op_params[3];
+    const cl_int d0 = dst->op_params[4];
+    const cl_int d1 = dst->op_params[5];
+
+    const cl_long cl_nb01 = nb01/ggml_type_size(src0->type);
+    const cl_long cl_nb03 = nb03/ggml_type_size(src0->type);
+    const cl_long cl_nb11 = nb11/ggml_type_size(src1->type);
+    const cl_long cl_nb12 = nb12/ggml_type_size(src1->type);
+    const cl_long cl_nb13 = nb13/ggml_type_size(src1->type);
+    const cl_long cl_nb1  = nb1/ggml_type_size(dst->type);
+    const cl_long cl_nb2  = nb2/ggml_type_size(dst->type);
+    const cl_long cl_nb3  = nb3/ggml_type_size(dst->type);
+
+    cl_kernel kernel;
+    if (src0->type == GGML_TYPE_F32) {
+        kernel = backend_ctx->kernel_conv2d_dw_f32_f32;
+    } else if (src0->type == GGML_TYPE_F16) {
+        kernel = backend_ctx->kernel_conv2d_dwf16_f32;
+    } else {
+        GGML_ASSERT(false && "Unsupported data type combination for conv2_dw");
+    }
+
+    cl_uint idx = 0;
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra0->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offset0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra1->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offset1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &C));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &N));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &KW));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &KH));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &W));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &H));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &OW));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &OH));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &s0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &s1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &p0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &p1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &d0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &d1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb01));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb03));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb11));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb12));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb13));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb2));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb3));
+
+    const size_t total = (size_t) N * C * OH * OW;
+    const size_t local_size = MIN((size_t) 256, backend_ctx->get_kernel_workgroup_size(kernel));
+    size_t global_work_size[] = { CEIL_DIV(total, local_size) * local_size };
+    size_t local_work_size[] = { local_size };
+
+    backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, dst);
+}
+
 // is_kq selects which of the two products this call is, and it is decided by the
 // CALLER -- the two admission arms in ggml_cl_mul_mat, each of which knows which
 // one it matched. It used to be re-derived here from nb01 > nb02, i.e. "K is
@@ -31030,6 +31137,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
             }
             func = ggml_cl_conv_2d;
             break;
+        case GGML_OP_CONV_2D_DW:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_conv2_dw;
+            break;
         case GGML_OP_SSM_CONV:
             if (!any_on_device) {
                 return false;

ggml/src/ggml-opencl/kernels/conv2_dw.cl

@@ -0,0 +1,139 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_conv2d_dw_f32_f32(
+        global const float * weights,
+        long                 weights_offset,
+        global const float * input,
+        long                 input_offset,
+        global       float * output,
+        long                 output_offset,
+        long C,
+        long N,
+        long KW,
+        long KH,
+        long W,
+        long H,
+        long OW,
+        long OH,
+        int s0,
+        int s1,
+        int p0,
+        int p1,
+        int d0,
+        int d1,
+        long nb01,
+        long nb03,
+        long nb11,
+        long nb12,
+        long nb13,
+        long nb1,
+        long nb2,
+        long nb3
+) {
+    weights = (global const float *) ((global const char *) weights + weights_offset);
+    input = (global const float *) ((global const char *) input + input_offset);
+    output = (global float *) ((global char *) output + output_offset);
+
+    const long gid = get_global_id(0);
+    const long OHW = OH * OW;
+    const long COHW = C * OHW;
+    const long total = N * COHW;
+    if (gid >= total) {
+        return;
+    }
+
+    const long n = gid / COHW;
+    const long c_rem = gid - n * COHW;
+    const long c = c_rem / OHW;
+    const long oh = (c_rem - c * OHW) / OW;
+    const long ow = c_rem - c * OHW - oh * OW;
+
+    float sum = 0.0f;
+    for (long kh = 0; kh < KH; ++kh) {
+        const long ih = oh * s1 + kh * d1 - p1;
+        if (ih < 0 || ih >= H) {
+            continue;
+        }
+        for (long kw = 0; kw < KW; ++kw) {
+            const long iw = ow * s0 + kw * d0 - p0;
+            if (iw < 0 || iw >= W) {
+                continue;
+            }
+            const long weight_idx = kw + kh * nb01 + c * nb03;
+            const long input_idx = iw + ih * nb11 + c * nb12 + n * nb13;
+            sum += weights[weight_idx] * input[input_idx];
+        }
+    }
+
+    const long output_idx = ow + oh * nb1 + c * nb2 + n * nb3;
+    output[output_idx] = sum;
+}
+
+kernel void kernel_conv2d_dwf16_f32(
+        global const half * weights,
+        long                weights_offset,
+        global const float * input,
+        long                 input_offset,
+        global       float * output,
+        long                 output_offset,
+        long C,
+        long N,
+        long KW,
+        long KH,
+        long W,
+        long H,
+        long OW,
+        long OH,
+        int s0,
+        int s1,
+        int p0,
+        int p1,
+        int d0,
+        int d1,
+        long nb01,
+        long nb03,
+        long nb11,
+        long nb12,
+        long nb13,
+        long nb1,
+        long nb2,
+        long nb3
+) {
+    weights = (global const half *) ((global const char *) weights + weights_offset);
+    input = (global const float *) ((global const char *) input + input_offset);
+    output = (global float *) ((global char *) output + output_offset);
+
+    const long gid = get_global_id(0);
+    const long OHW = OH * OW;
+    const long COHW = C * OHW;
+    const long total = N * COHW;
+    if (gid >= total) {
+        return;
+    }
+
+    const long n = gid / COHW;
+    const long c_rem = gid - n * COHW;
+    const long c = c_rem / OHW;
+    const long oh = (c_rem - c * OHW) / OW;
+    const long ow = c_rem - c * OHW - oh * OW;
+
+    float sum = 0.0f;
+    for (long kh = 0; kh < KH; ++kh) {
+        const long ih = oh * s1 + kh * d1 - p1;
+        if (ih < 0 || ih >= H) {
+            continue;
+        }
+        for (long kw = 0; kw < KW; ++kw) {
+            const long iw = ow * s0 + kw * d0 - p0;
+            if (iw < 0 || iw >= W) {
+                continue;
+            }
+            const long weight_idx = kw + kh * nb01 + c * nb03;
+            const long input_idx = iw + ih * nb11 + c * nb12 + n * nb13;
+            sum += (float) weights[weight_idx] * input[input_idx];
+        }
+    }
+
+    const long output_idx = ow + oh * nb1 + c * nb2 + n * nb3;
+    output[output_idx] = sum;
+}

ggml/src/ggml-sycl/common.hpp

@@ -68,6 +68,9 @@ extern int g_ggml_sycl_enable_optimize;
 extern int g_ggml_sycl_enable_fusion;
 extern int g_ggml_sycl_enable_esimd;
 extern int g_ggml_sycl_mmvq_wide;
+extern int g_ggml_sycl_enable_xmx_esimd;
+extern int g_ggml_sycl_xmx_max_cols;
+extern int g_ggml_sycl_xmx_glu_max_cols;
 extern int g_ggml_sycl_prioritize_dmmv;
 
 // Which quantized weight formats may take the XMX dequant-GEMM paths. A bitmask rather than one
@@ -227,6 +230,41 @@ typedef sycl::float2 dfloat2;
 
 #define MMVQ_MAX_BATCH_SIZE  8
 
+#if defined(__INTEL_LLVM_COMPILER) && __has_include(<sycl/ext/intel/esimd/xmx/dpas.hpp>)
+#    define GGML_SYCL_MMVQ_HAS_XMX
+#endif // __INTEL_LLVM_COMPILER
+
+// default of the most columns the XMX mul_mat_vec_q handles (GGML_SYCL_XMX_MAX_COLS), wider batches use other kernels
+#define GGML_SYCL_XMX_MAX_COLS_DEFAULT 80
+
+// most columns the fused gate and up XMX kernel is built for, GGML_SYCL_XMX_GLU_MAX_COLS can only lower this limit
+#define GGML_SYCL_XMX_GLU_MAX_COLS 16
+
+static inline bool ggml_sycl_xmx_supports_type(ggml_type type) {
+    switch (type) {
+        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_Q8_0:
+            return true;
+        default:
+            return false;
+    }
+}
+
+// the DPAS shapes and 2D block loads used by the XMX kernel are those of Xe2
+static inline bool ggml_sycl_xmx_supports_arch(gpu_arch arch) {
+    return arch == gpu_arch::intel_gpu_bmg_g21 || arch == gpu_arch::intel_gpu_bmg_g31 ||
+           arch == gpu_arch::intel_gpu_lnl_m;
+}
+
+// fewest columns for which the XMX kernel is faster than the other kernels
+static inline int ggml_sycl_xmx_min_cols(ggml_type type) {
+    return type == GGML_TYPE_Q2_K || type == GGML_TYPE_Q3_K || type == GGML_TYPE_Q4_K ? 1 : 2;
+}
+
 static int g_all_sycl_device_count = -1;
 static bool g_ggml_backend_sycl_buffer_type_initialized = false;
 
@@ -310,6 +348,10 @@ struct ggml_sycl_device_info {
 
 const ggml_sycl_device_info & ggml_sycl_info();
 
+static inline bool ggml_sycl_xmx_enabled(int device) {
+    return g_ggml_sycl_enable_xmx_esimd && ggml_sycl_xmx_supports_arch(ggml_sycl_info().devices[device].hw_info.arch);
+}
+
 static constexpr size_t SYCL_BUFFER_ALIGNMENT = 128;
 
 struct ggml_sycl_pool {

ggml/src/ggml-sycl/esimd.hpp

@@ -2,6 +2,8 @@
 #define GGML_SYCL_ESIMD_HPP
 
 #include <sycl/ext/intel/esimd.hpp>
+#include <sycl/ext/intel/esimd/xmx/dpas.hpp>
+#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
 
 #include "common.hpp"
 
@@ -584,6 +586,814 @@ template <> struct esimd_reorder_q_traits<GGML_TYPE_Q6_K> {
     }
 };
 
+//
+// Multi-column mul_mat_vec for reordered K-quant weights on the XMX engines (int8 DPAS).
+// A thread owns 16 weight rows (B operand, one dword is 4 weights of a row), the q8_1 activation
+// columns are the A operand. A super-block is 8 groups of 32, run as 2 parts of 4 groups. A quant
+// type supplies its loads and unpack as a traits struct.
+//
+using sycl::ext::intel::esimd::simd;
+constexpr int XMX_ROWS        = 16;  // weight rows per thread (DPAS N)
+constexpr int XMX_CHUNK       = 64;  // widest single launch, more columns run as repeated launches
+constexpr int XMX_FUSED_CHUNK = GGML_SYCL_XMX_GLU_MAX_COLS;  // widest launch of the fused gate and up kernel
+
+// threads that split K for one tile, sized so that the shared memory stays small
+template <int NC> constexpr int xmx_n_split() {
+    return NC <= 8 ? 32 : (NC == 16 ? 16 : 8);
+}
+
+inline bool xmx_supported(const void * vx, const int ncols) {
+    // 2D block loads need a 64 byte aligned base
+    return ncols % QK_K == 0 && (uintptr_t) vx % 64 == 0;
+}
+
+// transposed 2D block load of 16 rows, dword (kd, n) lands at [kd * 16 + n]
+// 2D block loads need rows of at least 64 bytes, shorter rows (SR) are read row by row
+template <int W, bool SR>
+ESIMD_INLINE simd<uint32_t, W * XMX_ROWS> xmx_load_t(const void * base, unsigned pitch, unsigned h, int x, int y) {
+    using namespace sycl::ext::intel::esimd;
+    if constexpr (SR) {
+        simd<uint32_t, W * XMX_ROWS> r;
+        for (int n = 0; n < XMX_ROWS; ++n) {
+            const unsigned row                = y + n < (int) h ? y + n : h;
+            const char *   p                  = (const char *) base + (size_t) row * (pitch + 1);
+            r.template select<W, XMX_ROWS>(n) = block_load<uint32_t, W>((const uint32_t *) p + x);
+        }
+        return r;
+    } else {
+        return load_2d<uint32_t, W, XMX_ROWS, 1, true, false>((const uint32_t *) base, pitch, h, pitch, x, y);
+    }
+}
+
+// per byte (q - zero) as int8 for q in 0..127, no borrow between bytes
+ESIMD_INLINE simd<uint32_t, 128> xmx_bytes_sub(simd<uint32_t, 128> x, uint32_t zero_x4) {
+    return ((x | 0x80808080u) - zero_x4) ^ 0x80808080u;
+}
+
+// How the scales of a quant type enter the DPAS epilogue of a group of 32 (see xmx_group)
+enum class xmx_scale_kind {
+    split,  // two scales per group, one per 16 wide half: the weights are masked to one half for each DPAS
+    min,    // one scale and one min per group: a second DPAS sums the activations for the min term
+    min16,  // a scale and a min per 16 wide half: the scales are folded into the DPAS operands
+    plain,  // one scale per group
+};
+
+// What every quant type states about itself. nc_wide is the tile width from which it takes the large register file.
+template <bool SR, int NC_WIDE, xmx_scale_kind KIND> struct xmx_traits_base {
+    static constexpr bool short_rows = SR;
+    static constexpr int  nc_wide    = NC_WIDE;
+    static constexpr bool has_min    = KIND == xmx_scale_kind::min;
+    static constexpr bool plain      = KIND == xmx_scale_kind::plain;
+    static constexpr bool min16      = KIND == xmx_scale_kind::min16;
+};
+
+// The part of a context every quant type shares: the last row (the clamp of 2D loads), this thread's first row and
+// the number of super-blocks per row. A type adds its array pointers and 2D load pitches.
+struct xmx_ctx_base {
+    unsigned h;
+    int      row0;
+    int      bpr;
+};
+
+ESIMD_INLINE void xmx_ctx_init(xmx_ctx_base & c, int nrows, int bpr, int row0) {
+    c.h    = nrows - 1;
+    c.row0 = row0;
+    c.bpr  = bpr;
+}
+
+// pitch of a 2D load over an array with `bytes` bytes per super-block
+constexpr unsigned xmx_pitch(int bpr, int bytes) {
+    return bpr * bytes - 1;
+}
+
+// per group row scales of the 16 rows, one vector per group
+template <int N> struct xmx_row_scales {
+    simd<float, 16> v[N];
+
+    ESIMD_INLINE simd<float, 16> get(int i) { return v[i]; }
+};
+
+// the 12 packed scale bytes of a super-block (Q3_K, Q4_K, Q5_K) of the 16 rows as three dwords
+ESIMD_INLINE void xmx_load_sc12(const uint8_t *            sc,
+                                const simd<uint32_t, 16> & bi,
+                                simd<uint32_t, 16> &       w0,
+                                simd<uint32_t, 16> &       w1,
+                                simd<uint32_t, 16> &       w2) {
+    using namespace sycl::ext::intel::esimd;
+    w0 = gather<uint32_t, 16>((const uint32_t *) sc, bi * 12);
+    w1 = gather<uint32_t, 16>((const uint32_t *) sc, bi * 12 + 4);
+    w2 = gather<uint32_t, 16>((const uint32_t *) sc, bi * 12 + 8);
+}
+
+// d and dmin (a half2 per super-block, Q2_K, Q4_K, Q5_K) of the 16 rows
+ESIMD_INLINE void xmx_load_dm(const sycl::half *          dm,
+                              const simd<uint32_t, 16> & bi,
+                              simd<float, 16> &          d,
+                              simd<float, 16> &          dmin) {
+    using namespace sycl::ext::intel::esimd;
+    d    = convert<float>(gather<sycl::half, 16>(dm, bi * (2 * (uint32_t) sizeof(sycl::half))));
+    dmin = convert<float>(
+        gather<sycl::half, 16>(dm, bi * (2 * (uint32_t) sizeof(sycl::half)) + (uint32_t) sizeof(sycl::half)));
+}
+
+// Q6_K reorder layout: [ql: nb*128] [qh: nb*64] [scales int8: nb*16] [d: nb*2]
+template <bool SR> struct xmx_traits_q6_k : xmx_traits_base<SR, 7, xmx_scale_kind::split> {
+    struct ctx : xmx_ctx_base {
+        const uint8_t *    ql;
+        const uint8_t *    qh;
+        const int8_t *     scales;
+        const sycl::half * dh;
+        unsigned           pl, ph, ps;
+    };
+
+    static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) {
+        const size_t nb = (size_t) nrows * bpr;
+        ctx          c;
+        xmx_ctx_init(c, nrows, bpr, row0);
+        c.ql     = (const uint8_t *) vx;
+        c.qh     = c.ql + nb * (QK_K / 2);
+        c.scales = (const int8_t *) (c.qh + nb * (QK_K / 4));
+        c.dh     = (const sycl::half *) (c.scales + nb * (QK_K / 16));
+        c.pl     = xmx_pitch(bpr, QK_K / 2);
+        c.ph     = xmx_pitch(bpr, QK_K / 4);
+        c.ps     = xmx_pitch(bpr, QK_K / 16);
+        return c;
+    }
+
+    // row scales of the 16 scale groups of a super-block
+    struct scales_t {
+        simd<int8_t, 256> sc8;
+        simd<float, 16>   dv;
+
+        ESIMD_INLINE simd<float, 16> get(int gr) {
+            simd<int8_t, 16> s = sc8.template select<16, 4>(64 * (gr / 4) + (gr % 4));
+            return sycl::ext::intel::esimd::convert<float>(s) * dv;
+        }
+    };
+
+    // rows: clamped row index of each of the 16 lanes
+    static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd<uint32_t, 16> & rows) {
+        using namespace sycl::ext::intel::esimd;
+        scales_t s;
+        s.dv = convert<float>(gather<sycl::half, 16>(c.dh, (rows * c.bpr + sb) * (uint32_t) sizeof(sycl::half)));
+        simd<uint32_t, 64> raw = xmx_load_t<4, SR>(c.scales, c.ps, c.h, sb * 4, c.row0);
+        s.sc8                  = raw.template bit_cast_view<int8_t>();
+        return s;
+    }
+
+    struct part_t {
+        simd<uint32_t, 128> a, b, h;
+    };
+
+    static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) {
+        part_t p;
+        p.a = xmx_load_t<8, SR>(c.ql, c.pl, c.h, sb * 32 + part * 16, c.row0);
+        p.b = xmx_load_t<8, SR>(c.ql, c.pl, c.h, sb * 32 + part * 16 + 8, c.row0);
+        p.h = xmx_load_t<8, SR>(c.qh, c.ph, c.h, sb * 16 + part * 8, c.row0);
+        return p;
+    }
+
+    static ESIMD_INLINE simd<uint32_t, 128> unpack(const part_t & p, int g, int /*part*/) {
+        simd<uint32_t, 128> qs = (g & 1) ? p.b : p.a;
+        simd<uint32_t, 128> x;
+        if (g == 0) {
+            x = (qs & 0x0F0F0F0Fu) | ((p.h << 4) & 0x30303030u);
+        } else if (g == 1) {
+            x = (qs & 0x0F0F0F0Fu) | ((p.h << 2) & 0x30303030u);
+        } else if (g == 2) {
+            x = ((qs >> 4) & 0x0F0F0F0Fu) | (p.h & 0x30303030u);
+        } else {
+            x = ((qs >> 4) & 0x0F0F0F0Fu) | ((p.h >> 2) & 0x30303030u);
+        }
+        return xmx_bytes_sub(x, 0x20202020u);
+    }
+};
+
+// Q3_K reorder layout: [qs: nb*64] [hmask: nb*32] [scales: nb*12] [d: nb*2]
+// chunk s of 32: 2 bits from qs byte 32*(s/4)+l at shift 2*(s%4), high bit is hmask[l] bit s, value is q - 4
+template <bool SR> struct xmx_traits_q3_k : xmx_traits_base<SR, 16, xmx_scale_kind::split> {
+    struct ctx : xmx_ctx_base {
+        const uint8_t *    qs;
+        const uint8_t *    hm;
+        const uint8_t *    sc;
+        const sycl::half * dh;
+        unsigned           pq, ph;
+    };
+
+    static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) {
+        const size_t nb = (size_t) nrows * bpr;
+        ctx          c;
+        xmx_ctx_init(c, nrows, bpr, row0);
+        c.qs = (const uint8_t *) vx;
+        c.hm = c.qs + nb * (QK_K / 4);
+        c.sc = c.hm + nb * (QK_K / 8);
+        c.dh = (const sycl::half *) (c.sc + nb * 12);
+        c.pq = xmx_pitch(bpr, QK_K / 4);
+        c.ph = xmx_pitch(bpr, QK_K / 8);
+        return c;
+    }
+
+    using scales_t = xmx_row_scales<16>;
+
+    static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd<uint32_t, 16> & rows) {
+        using namespace sycl::ext::intel::esimd;
+        scales_t           s;
+        simd<uint32_t, 16> bi = rows * c.bpr + sb;
+        simd<uint32_t, 16> w0, w1, w2;
+        xmx_load_sc12(c.sc, bi, w0, w1, w2);
+        simd<float, 16> dv = convert<float>(gather<sycl::half, 16>(c.dh, bi * (uint32_t) sizeof(sycl::half)));
+#pragma unroll
+        for (int i = 0; i < 16; ++i) {
+            const simd<uint32_t, 16> lo   = (i % 8) < 4 ? w0 : w1;
+            const int                bl   = i % 4;
+            simd<uint32_t, 16>       nib  = (lo >> (8 * bl + (i < 8 ? 0 : 4))) & 0xFu;
+            simd<uint32_t, 16>       hi   = ((w2 >> (8 * bl + 2 * (i / 4))) & 3u) << 4;
+            simd<int, 16>            code = nib | hi;
+            s.v[i]                        = convert<float>(code - 32) * dv;
+        }
+        return s;
+    }
+
+    struct part_t {
+        simd<uint32_t, 128> q, h;
+    };
+
+    static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) {
+        part_t p;
+        p.q = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 16 + part * 8, c.row0);
+        p.h = xmx_load_t<8, SR>(c.hm, c.ph, c.h, sb * 8, c.row0);
+        return p;
+    }
+
+    static ESIMD_INLINE simd<uint32_t, 128> unpack(const part_t & p, int g, int part) {
+        const int           s = 4 * part + g;
+        simd<uint32_t, 128> x = ((p.q >> (2 * g)) & 0x03030303u) | (((p.h >> s) & 0x01010101u) << 2);
+        return xmx_bytes_sub(x, 0x04040404u);
+    }
+};
+
+// Q4_K and Q5_K (QH) reorder layout: [qs: nb*128] ([qh: nb*32]) [scales: nb*12] [dm: nb*4]
+// group j of 32: nibble of qs byte 32*(j/2)+l (low for even j), Q5_K adds 16 * bit j of qh[l]
+// value is dall * sc_j * q - dmin * m_j, one scale and one min per group
+template <bool QH, bool SR> struct xmx_traits_q45_k : xmx_traits_base<SR, QH ? 16 : 7, xmx_scale_kind::min> {
+    struct ctx : xmx_ctx_base {
+        const uint8_t *    qs;
+        const uint8_t *    qh;
+        const uint8_t *    sc;
+        const sycl::half * dm;
+        unsigned           pq, ph;
+    };
+
+    static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) {
+        const size_t nb = (size_t) nrows * bpr;
+        ctx          c;
+        xmx_ctx_init(c, nrows, bpr, row0);
+        c.qs = (const uint8_t *) vx;
+        c.qh = c.qs + nb * (QK_K / 2);
+        c.sc = QH ? c.qh + nb * (QK_K / 8) : c.qh;
+        c.dm = (const sycl::half *) (c.sc + nb * 12);
+        c.pq = xmx_pitch(bpr, QK_K / 2);
+        c.ph = xmx_pitch(bpr, QK_K / 8);
+        return c;
+    }
+
+    struct scales_t {
+        xmx_row_scales<8> d;  // dall * scale of each group
+        xmx_row_scales<8> m;  // dmin * min of each group
+    };
+
+    static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd<uint32_t, 16> & rows) {
+        using namespace sycl::ext::intel::esimd;
+        scales_t           s;
+        simd<uint32_t, 16> bi = rows * c.bpr + sb;
+        simd<uint32_t, 16> w0, w1, w2;
+        xmx_load_sc12(c.sc, bi, w0, w1, w2);
+        simd<float, 16> d, dmin;
+        xmx_load_dm(c.dm, bi, d, dmin);
+#pragma unroll
+        for (int j = 0; j < 8; ++j) {
+            const int          jj = j % 4;
+            simd<uint32_t, 16> sc, m;
+            if (j < 4) {
+                sc = (w0 >> (8 * jj)) & 63u;
+                m  = (w1 >> (8 * jj)) & 63u;
+            } else {
+                sc = ((w2 >> (8 * jj)) & 0xFu) | (((w0 >> (8 * jj + 6)) & 3u) << 4);
+                m  = ((w2 >> (8 * jj + 4)) & 0xFu) | (((w1 >> (8 * jj + 6)) & 3u) << 4);
+            }
+            s.d.v[j] = convert<float>(simd<int, 16>(sc)) * d;
+            s.m.v[j] = convert<float>(simd<int, 16>(m)) * dmin;
+        }
+        return s;
+    }
+
+    struct part_t {
+        simd<uint32_t, 128> a, b, h;
+    };
+
+    static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) {
+        part_t p;
+        p.a = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 32 + part * 16, c.row0);
+        p.b = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 32 + part * 16 + 8, c.row0);
+        if constexpr (QH) {
+            p.h = xmx_load_t<8, SR>(c.qh, c.ph, c.h, sb * 8, c.row0);
+        }
+        return p;
+    }
+
+    static ESIMD_INLINE simd<uint32_t, 128> unpack(const part_t & p, int g, int part) {
+        simd<uint32_t, 128> qs = (g < 2) ? p.a : p.b;
+        simd<uint32_t, 128> x  = (qs >> (4 * (g & 1))) & 0x0F0F0F0Fu;
+        if constexpr (QH) {
+            x |= ((p.h >> (4 * part + g)) & 0x01010101u) << 4;
+        }
+        return x;
+    }
+};
+
+// Q2_K reorder layout: [qs: nb*64] [scales: nb*16] [dm: nb*4]
+// chunk s of 32: 2 bits from qs byte 32*(s/4)+l at shift 2*(s%4), each 16 wide half has a 4 bit scale and min
+// value is d * (code & 15) * q - dmin * (code >> 4)
+template <bool SR> struct xmx_traits_q2_k : xmx_traits_base<SR, 16, xmx_scale_kind::min16> {
+    struct ctx : xmx_ctx_base {
+        const uint8_t *    qs;
+        const uint8_t *    sc;
+        const sycl::half * dm;
+        unsigned           pq, ps;
+    };
+
+    static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) {
+        const size_t nb = (size_t) nrows * bpr;
+        ctx          c;
+        xmx_ctx_init(c, nrows, bpr, row0);
+        c.qs = (const uint8_t *) vx;
+        c.sc = c.qs + nb * (QK_K / 4);
+        c.dm = (const sycl::half *) (c.sc + nb * (QK_K / 16));
+        c.pq = xmx_pitch(bpr, QK_K / 4);
+        c.ps = xmx_pitch(bpr, QK_K / 16);
+        return c;
+    }
+
+    struct scales_t {
+        simd<uint8_t, 256> code;  // 4 bit scale (low) and min (high) of each 16 wide half
+        simd<float, 16>    d, dmin;
+
+        ESIMD_INLINE simd<uint8_t, 16> byte(int gr) { return code.template select<16, 4>(64 * (gr / 4) + (gr % 4)); }
+    };
+
+    static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd<uint32_t, 16> & rows) {
+        using namespace sycl::ext::intel::esimd;
+        scales_t           s;
+        simd<uint32_t, 16> bi  = rows * c.bpr + sb;
+        simd<uint32_t, 64> raw = xmx_load_t<4, SR>(c.sc, c.ps, c.h, sb * 4, c.row0);
+        s.code                 = raw.template bit_cast_view<uint8_t>();
+        xmx_load_dm(c.dm, bi, s.d, s.dmin);
+        return s;
+    }
+
+    struct part_t {
+        simd<uint32_t, 128> q;
+    };
+
+    static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) {
+        part_t p;
+        p.q = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 16 + part * 8, c.row0);
+        return p;
+    }
+
+    static ESIMD_INLINE simd<uint32_t, 128> unpack(const part_t & p, int g, int /*part*/) {
+        return (p.q >> (2 * g)) & 0x03030303u;
+    }
+};
+
+// Q8_0 reorder layout: [qs: nrows*K bytes] [d: one half per 32 weights]
+// the int8 weights feed DPAS as they are, a group is one block with one scale
+template <bool SR> struct xmx_traits_q8_0 : xmx_traits_base<SR, 16, xmx_scale_kind::plain> {
+    struct ctx : xmx_ctx_base {
+        const uint8_t *    qs;
+        const sycl::half * dh;
+        unsigned           pq;
+    };
+
+    static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) {
+        ctx c;
+        xmx_ctx_init(c, nrows, bpr, row0);
+        c.qs = (const uint8_t *) vx;
+        c.dh = (const sycl::half *) (c.qs + (size_t) nrows * bpr * QK_K);
+        c.pq = xmx_pitch(bpr, QK_K);
+        return c;
+    }
+
+    struct scales_t {
+        xmx_row_scales<8> d;  // scale of each of the 8 blocks of a super-block
+    };
+
+    static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd<uint32_t, 16> & rows) {
+        using namespace sycl::ext::intel::esimd;
+        scales_t           s;
+        simd<uint32_t, 16> blk = (rows * c.bpr + sb) * (QK_K / QK8_0);
+#pragma unroll
+        for (int j = 0; j < 8; ++j) {
+            s.d.v[j] = convert<float>(gather<sycl::half, 16>(c.dh, (blk + j) * (uint32_t) sizeof(sycl::half)));
+        }
+        return s;
+    }
+
+    struct part_t {
+        simd<uint32_t, 128> q[4];
+    };
+
+    static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) {
+        part_t p;
+#pragma unroll
+        for (int g = 0; g < 4; ++g) {
+            p.q[g] = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 64 + 8 * (4 * part + g), c.row0);
+        }
+        return p;
+    }
+
+    static ESIMD_INLINE simd<uint32_t, 128> unpack(const part_t & p, int g, int /*part*/) { return p.q[g]; }
+};
+
+template <bool SR> using xmx_traits_q4_k = xmx_traits_q45_k<false, SR>;
+template <bool SR> using xmx_traits_q5_k = xmx_traits_q45_k<true, SR>;
+
+// registers per thread, the large file halves the resident threads but avoids spills in wide tiles, the fused
+// kernel keeps two sets of weights and scales and always takes it
+template <typename T, int NC, bool FUSED> constexpr int xmx_grf() {
+    return FUSED || NC >= T::nc_wide ? 256 : 128;
+}
+
+// One DPAS of a block of RC columns: the int32 products of the weights B and the activations A, as float and scaled
+// per weight row. Every quant type sums one or two of these per group of 32.
+template <int RC, int NT>
+ESIMD_INLINE simd<float, RC * NT> xmx_term(const simd<int8_t, 512> &    B,
+                                           const simd<int8_t, RC * 32> & A,
+                                           const simd<float, NT> &       s) {
+    using namespace sycl::ext::intel::esimd;
+    namespace xmx = sycl::ext::intel::esimd::xmx;
+
+    simd<int, RC * NT> C = xmx::dpas<8, RC, int>(B, A);
+    return convert<float>(C) * s.template replicate<RC>();
+}
+
+// one group of 32 of one weight set, accumulated into Cf for all the columns of the tile
+template <typename T, int NC, typename A2T, typename D8T>
+ESIMD_INLINE void xmx_group(simd<float, NC * XMX_ROWS> & Cf,
+                            const typename T::part_t &   pd,
+                            typename T::scales_t &       sc,
+                            A2T &                        A2,
+                            D8T &                        d8f,
+                            simd<int8_t, 512> &          Bl,
+                            simd<int8_t, 512> &          Bh,
+                            const simd<int8_t, 512> &    Ol,
+                            const int                    g,
+                            const int                    part) {
+    using namespace sycl::ext::intel::esimd;
+
+    constexpr int  NT    = XMX_ROWS;
+    constexpr int  RC    = NC < 8 ? NC : 8;
+    constexpr int  NB    = NC / RC;
+    constexpr bool split = !T::has_min && !T::plain && !T::min16;  // two scales per group of 32
+    constexpr bool pair  = (RC * 32) % 64 == 0;                    // a block of RC rows fills whole registers
+
+    const int gi = 4 * part + g;
+
+    simd<uint32_t, 128> x = T::unpack(pd, g, part);
+    simd<int8_t, 512>   B = x.template bit_cast_view<int8_t>();
+    if constexpr (split) {
+        Bl.template select<256, 1>(0)   = B.template select<256, 1>(0);
+        Bh.template select<256, 1>(256) = B.template select<256, 1>(256);
+    }
+
+    simd<float, NT> rd;
+    simd<float, NT> rm;
+    simd<float, NT> rs0;
+    simd<float, NT> rs1;
+
+    // Q2_K: scale and min of each 16 wide half are 4 bit codes, the weight is d * scl * q - dmin * mc. scl * q
+    // is at most 45, so a dword multiply keeps the bytes apart and the scaled weights go to DPAS as int8. The
+    // min term sums a * mc over K, the codes mc take the place of the weights in a second DPAS.
+    simd<int8_t, 512> Bs;
+    simd<int8_t, 512> Bm;
+    if constexpr (T::min16) {
+        const simd<uint8_t, NT> c0  = sc.byte(2 * gi);
+        const simd<uint8_t, NT> c1  = sc.byte(2 * gi + 1);
+        simd<uint32_t, NT>      sl0 = convert<uint32_t>(simd<uint8_t, NT>(c0 & (uint8_t) 15));
+        simd<uint32_t, NT>      sl1 = convert<uint32_t>(simd<uint8_t, NT>(c1 & (uint8_t) 15));
+        simd<uint32_t, NT>      mc0 = convert<uint32_t>(simd<uint8_t, NT>(c0 >> (uint8_t) 4)) * 0x01010101u;
+        simd<uint32_t, NT>      mc1 = convert<uint32_t>(simd<uint8_t, NT>(c1 >> (uint8_t) 4)) * 0x01010101u;
+        simd<uint32_t, 128>     ws;
+        simd<uint32_t, 128>     wm;
+        ws.template select<64, 1>(0)  = x.template select<64, 1>(0) * sl0.template replicate<4>();
+        ws.template select<64, 1>(64) = x.template select<64, 1>(64) * sl1.template replicate<4>();
+        wm.template select<64, 1>(0)  = mc0.template replicate<4>();
+        wm.template select<64, 1>(64) = mc1.template replicate<4>();
+        Bs                            = ws.template bit_cast_view<int8_t>();
+        Bm                            = wm.template bit_cast_view<int8_t>();
+    }
+    if constexpr (T::plain) {
+        rd = sc.d.get(gi);
+    } else if constexpr (T::has_min) {
+        rd = sc.d.get(gi);
+        rm = sc.m.get(gi);
+    } else if constexpr (T::min16) {
+        // scales are applied through Bs and Bm
+    } else {
+        rs0 = sc.get(2 * gi);
+        rs1 = sc.get(2 * gi + 1);
+    }
+
+#pragma unroll
+    for (int b = 0; b < NB; ++b) {
+        simd<int8_t, RC * 32> A =
+            A2[b].template select<RC * 8, 1>(pair ? (g % 2) * RC * 8 : 0).template bit_cast_view<int8_t>();
+        simd<float, RC * NT> d8r =
+            d8f[b].template select<RC, RC == 1 ? 1 : 8>(gi).template replicate_vs_w_hs<RC, 1, NT, 0>(0);
+        auto Cb = Cf.template select<RC * NT, 1>(b * RC * NT);
+
+        simd<float, RC * NT> z;
+        if constexpr (T::plain) {
+            z = xmx_term<RC, NT>(B, A, rd);
+        } else if constexpr (T::has_min) {
+            // the second DPAS (all ones against the activations) is the exact integer sum of the activations
+            z = xmx_term<RC, NT>(B, A, rd) - xmx_term<RC, NT>(Ol, A, rm);
+        } else if constexpr (T::min16) {
+            z = xmx_term<RC, NT>(Bs, A, sc.d) - xmx_term<RC, NT>(Bm, A, sc.dmin);
+        } else {
+            z = xmx_term<RC, NT>(Bl, A, rs0) + xmx_term<RC, NT>(Bh, A, rs1);
+        }
+        Cb += z * d8r;
+    }
+}
+
+// activation of the GLU applied to the gate, uniform across the launch
+template <int N> ESIMD_INLINE simd<float, N> xmx_glu_act(const simd<float, N> & gate, const int glu_op) {
+    using namespace sycl::ext::intel::esimd;
+    if (glu_op == GGML_GLU_OP_SWIGLU) {
+        return gate / (1.0f + exp(-gate));
+    }
+    // GEGLU, the tanh is written with exp
+    simd<float, N> a = 0.79788456f * gate * (1.0f + 0.044715f * gate * gate);
+    simd<float, N> t = 1.0f - 2.0f / (exp(2.0f * a) + 1.0f);
+    return 0.5f * gate * (1.0f + t);
+}
+
+// activations come in through 2D block loads, one for the (d, s) pairs of a super-block (s is unused) and one per two
+// groups for the quants, the DPAS A operand needs no repacking. Groups with two scales (16 wide) mask the
+// weights instead, zeros in B drop the other half of K.
+// KS is the most threads that split K for one tile, the launch may use fewer.
+// FUSED computes up * act(gate) from the weights vx (up) and vg (gate), reading the activations once.
+template <typename T, int NC, int KS, bool FUSED>
+ESIMD_INLINE void xmx_mul_mat(const void * vx,
+                              const void * vg,
+                              const char * vy,
+                              float *      dst,
+                              const int    ncols,
+                              const int    nrows,
+                              const int    ncv,
+                              const int    stride_y_bytes,
+                              const int    ldd,
+                              const int    glu_op,
+                              const int    tile,
+                              const int    lid,
+                              const int    ks) {
+    using namespace sycl::ext::intel::esimd;
+
+    constexpr int  NT    = XMX_ROWS;
+    constexpr int  RC    = NC < 8 ? NC : 8;
+    constexpr int  NB    = NC / RC;
+    constexpr int  N_ACC = FUSED ? 2 : 1;
+    constexpr bool pair  = (RC * 32) % 64 == 0;  // a block of RC rows fills whole registers
+
+    if constexpr (KS > 1) {
+        slm_init<N_ACC * KS * NC * NT * sizeof(float)>();
+    }
+
+    const int  bpr  = ncols / QK_K;
+    const int  row0 = tile * NT;
+    const auto c    = T::make(vx, nrows, bpr, row0);
+    const auto cg   = T::make(FUSED ? vg : vx, nrows, bpr, row0);
+
+    simd<uint32_t, NT> rows;
+#pragma unroll
+    for (int n = 0; n < NT; ++n) {
+        rows[n] = row0 + n < nrows ? row0 + n : nrows - 1;
+    }
+
+    const uint32_t * yq  = (const uint32_t *) vy;
+    const uint32_t * yd  = (const uint32_t *) (vy + ncols);
+    const unsigned   yqw = ncols - 1;
+    const unsigned   ydw = ncols / QK8_1 * sizeof(sycl::half2) - 1;
+    const unsigned   yh  = ncv - 1;
+    const unsigned   yp  = stride_y_bytes - 1;
+
+    simd<float, NC * NT> Cf  = 0.0f;
+    simd<float, NC * NT> Cfg = 0.0f;  // gate, only used when fused
+    simd<int8_t, 512>    Bl  = int8_t(0);
+    simd<int8_t, 512>    Bh  = int8_t(0);
+    simd<int8_t, 512>    Ol  = int8_t(0);
+    if constexpr (T::has_min) {
+        Ol = int8_t(1);
+    }
+
+    for (int sb = lid; sb < bpr; sb += ks) {
+        auto                 sc = T::load_scales(c, sb, rows);
+        typename T::scales_t scg;
+        if constexpr (FUSED) {
+            scg = T::load_scales(cg, sb, rows);
+        }
+
+        // d of the 8 groups of the super-block, [column * 8 + group]
+        simd<float, 8 * RC> d8f[NB];
+#pragma unroll
+        for (int b = 0; b < NB; ++b) {
+            simd<uint32_t, 8 * RC>    ds = load_2d<uint32_t, 8, RC, 1, false, false>(yd, ydw, yh, yp, sb * 8, b * RC);
+            simd<sycl::half, 16 * RC> dh = ds.template bit_cast_view<sycl::half>();
+            d8f[b]                       = convert<float>(simd<sycl::half, 8 * RC>(dh.template select<8 * RC, 2>(0)));
+        }
+
+#pragma unroll
+        for (int part = 0; part < 2; ++part) {
+            const auto         pd = T::load_part(c, sb, part);
+            typename T::part_t pdg;
+            if constexpr (FUSED) {
+                pdg = T::load_part(cg, sb, part);
+            }
+
+            simd<uint32_t, 16 * RC> A2[NB];  // quants of two groups
+
+#pragma unroll
+            for (int g = 0; g < 4; ++g) {
+                const int gi = 4 * part + g;
+                if constexpr (pair) {
+                    if (g % 2 == 0) {
+#pragma unroll
+                        for (int b = 0; b < NB; ++b) {
+                            A2[b] =
+                                load_2d<uint32_t, 8, RC, 2, false, false>(yq, yqw, yh, yp, sb * 64 + gi * 8, b * RC);
+                        }
+                    }
+                } else {
+#pragma unroll
+                    for (int b = 0; b < NB; ++b) {
+                        A2[b].template select<RC * 8, 1>(0) =
+                            load_2d<uint32_t, 8, RC, 1, false, false>(yq, yqw, yh, yp, sb * 64 + gi * 8, b * RC);
+                    }
+                }
+
+                xmx_group<T, NC>(Cf, pd, sc, A2, d8f, Bl, Bh, Ol, g, part);
+                if constexpr (FUSED) {
+                    xmx_group<T, NC>(Cfg, pdg, scg, A2, d8f, Bl, Bh, Ol, g, part);
+                }
+            }
+        }
+    }
+
+    if constexpr (KS > 1) {
+        slm_block_store<float, NC * NT>(lid * NC * NT * sizeof(float), Cf);
+        if constexpr (FUSED) {
+            slm_block_store<float, NC * NT>((KS + lid) * NC * NT * sizeof(float), Cfg);
+        }
+        barrier();
+        if (lid != 0) {
+            return;
+        }
+        Cf  = 0.0f;
+        Cfg = 0.0f;
+#pragma unroll 4
+        for (int j = 0; j < ks; ++j) {
+            Cf += slm_block_load<float, NC * NT>(j * NC * NT * sizeof(float));
+            if constexpr (FUSED) {
+                Cfg += slm_block_load<float, NC * NT>((KS + j) * NC * NT * sizeof(float));
+            }
+        }
+    }
+
+    if constexpr (FUSED) {
+        Cf *= xmx_glu_act<NC * NT>(Cfg, glu_op);
+    }
+
+    if (row0 + NT <= nrows && ldd % 4 == 0 && (uintptr_t) dst % 16 == 0) {
+#pragma unroll
+        for (int m = 0; m < NC; ++m) {
+            if (NC <= 8 || m < ncv) {
+                block_store<float, NT>(dst + (size_t) m * ldd + row0, Cf.template select<NT, 1>(m * NT));
+            }
+        }
+    } else {
+        // tail tile or unaligned output
+        simd<uint32_t, NT> lane(0, 1);
+        simd_mask<NT>      in_rows = (lane + row0) < (uint32_t) nrows;
+#pragma unroll
+        for (int m = 0; m < NC; ++m) {
+            if (NC <= 8 || m < ncv) {
+                scatter<float, NT>(dst + (size_t) m * ldd + row0, lane * (uint32_t) sizeof(float),
+                                   Cf.template select<NT, 1>(m * NT), in_rows);
+            }
+        }
+    }
+}
+
+// KS_MAX bounds the threads that split K for one tile, the launch uses as few as keep every thread to the same
+// number of super-blocks as the busiest one
+template <typename T, int NC, int KS_MAX, bool FUSED>
+static void xmx_mul_mat_launch_ks(const void *    vx,
+                                  const void *    vg,
+                                  const char *    vy,
+                                  float *         dst,
+                                  const int       ncols,
+                                  const int       nrows,
+                                  const int       ncv,
+                                  const int       stride_y_bytes,
+                                  const int       ldd,
+                                  const int       glu_op,
+                                  dpct::queue_ptr stream) {
+    GGML_ASSERT(ncols % QK_K == 0);
+    const size_t n_tiles = (nrows + XMX_ROWS - 1) / XMX_ROWS;
+    const int    bpr     = ncols / QK_K;
+    const int    rounds  = (bpr + KS_MAX - 1) / KS_MAX;
+    const int    ks      = (bpr + rounds - 1) / rounds;
+    stream->submit([&](sycl::handler & cgh) {
+        cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(n_tiles * ks), sycl::range<1>(ks)),
+                         sycl::ext::oneapi::experimental::properties{
+                             sycl::ext::intel::experimental::grf_size<xmx_grf<T, NC, FUSED>()> },
+                         [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
+                             xmx_mul_mat<T, NC, KS_MAX, FUSED>(vx, vg, vy, dst, ncols, nrows, ncv, stride_y_bytes, ldd,
+                                                               glu_op, (int) it.get_group(0), (int) it.get_local_id(0),
+                                                               ks);
+                         });
+    });
+}
+
+template <typename T, int NC, bool FUSED>
+static void xmx_mul_mat_launch(const void *    vx,
+                               const void *    vg,
+                               const char *    vy,
+                               float *         dst,
+                               const int       ncols,
+                               const int       nrows,
+                               const int       ncv,
+                               const int       stride_y_bytes,
+                               const int       ldd,
+                               const int       glu_op,
+                               dpct::queue_ptr stream) {
+    // the fused kernel keeps two sets of accumulators in shared memory, half the threads for the same size
+    constexpr int ks_max = xmx_n_split<NC>() / (FUSED ? 2 : 1);
+    xmx_mul_mat_launch_ks<T, NC, ks_max, FUSED>(vx, vg, vy, dst, ncols, nrows, ncv, stride_y_bytes, ldd, glu_op,
+                                                stream);
+}
+
+// FUSED: vx and vg are the up and gate weights, dst gets up * act(gate) with glu_op a ggml_glu_op
+template <typename T, bool FUSED = false>
+static void xmx_mul_mat_ncols(const void *    vx,
+                              const void *    vg,
+                              const char *    vy,
+                              float *         dst,
+                              const int       ncols,
+                              const int       nrows,
+                              const int       ncols_dst,
+                              const int       stride_y_bytes,
+                              const int       ldd,
+                              const int       glu_op,
+                              dpct::queue_ptr stream) {
+    constexpr int chunk_max = FUSED ? XMX_FUSED_CHUNK : XMX_CHUNK;
+    for (int c0 = 0; c0 < ncols_dst; c0 += chunk_max) {
+        const int    n  = std::min(chunk_max, ncols_dst - c0);
+        const char * y  = vy + (size_t) c0 * stride_y_bytes;
+        float *      d  = dst + (size_t) c0 * ldd;
+        const int    nc = n > 32 ? 64 : (n > 16 ? 32 : (n > 8 ? 16 : n));
+        switch (nc) {
+#define XMX_CASE(N)                                                                                          \
+    case N:                                                                                                  \
+        xmx_mul_mat_launch<T, N, FUSED>(vx, vg, y, d, ncols, nrows, n, stride_y_bytes, ldd, glu_op, stream); \
+        break
+            XMX_CASE(1);
+            XMX_CASE(2);
+            XMX_CASE(3);
+            XMX_CASE(4);
+            XMX_CASE(5);
+            XMX_CASE(6);
+            XMX_CASE(7);
+            XMX_CASE(8);
+#undef XMX_CASE
+            // short row variants are for tests, not built for the wide tiles
+#define XMX_CASE_WIDE(N)                                                                                         \
+    case N:                                                                                                      \
+        if constexpr (!T::short_rows && (!FUSED || N <= XMX_FUSED_CHUNK)) {                                      \
+            xmx_mul_mat_launch<T, N, FUSED>(vx, vg, y, d, ncols, nrows, n, stride_y_bytes, ldd, glu_op, stream); \
+        }                                                                                                        \
+        break
+            XMX_CASE_WIDE(16);
+            XMX_CASE_WIDE(32);
+            XMX_CASE_WIDE(64);
+#undef XMX_CASE_WIDE
+            default:
+                GGML_ABORT("unsupported ncols_dst=%d for XMX MMV", n);
+        }
+    }
+}
+
 } // namespace ggml_sycl_esimd
 
 #endif // GGML_SYCL_ESIMD_HPP

ggml/src/ggml-sycl/fusion.cpp

@@ -37,6 +37,13 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg
     const bool reorder_pair = wu->type == GGML_TYPE_Q4_K && wg->type == GGML_TYPE_Q4_K;
     const bool plain_pair   = (wu->type == GGML_TYPE_Q5_K || wu->type == GGML_TYPE_IQ4_XS) &&
                             (wg->type == GGML_TYPE_Q5_K || wg->type == GGML_TYPE_IQ4_XS);
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+    // same-type pairs of the types of the XMX kernels, which also take a few more columns
+    // (the XMX kernels walk whole QK_K super-blocks, so they need rows of a multiple of QK_K)
+    const bool xmx_pair = wu->type == wg->type && ggml_sycl_xmx_supports_type(wu->type) && wu->ne[0] % QK_K == 0;
+#else
+    const bool xmx_pair = false;
+#endif // GGML_SYCL_MMVQ_HAS_XMX
     // the ESIMD kernel covers same-type q2_K..q6_K and q8_0, single column only (as unfused DMMV)
 #ifdef GGML_SYCL_DMMV_HAS_ESIMD
     const bool esimd_pair   = g_ggml_sycl_enable_esimd && wu->type == wg->type && act->ne[1] == 1 &&
@@ -47,7 +54,7 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg
 #endif
     // q8_0 rows are walked in 32-wide blocks, so they only need to be a multiple of QK8_0
     const int64_t row_align = esimd_pair && wu->type == GGML_TYPE_Q8_0 ? QK8_0 : QK_K;
-    if ((!reorder_pair && !plain_pair && !esimd_pair) || wu->ne[0] % row_align != 0) {
+    if ((!reorder_pair && !plain_pair && !xmx_pair && !esimd_pair) || wu->ne[0] % row_align != 0) {
         return false;
     }
 
@@ -67,7 +74,8 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg
         return false;
     }
     // mat-vec only: one column per decoded token, up to the batch the reorder kernels cover
-    if (act->ne[1] > MMVQ_MAX_BATCH_SIZE) {
+    const int max_cols = xmx_pair ? std::max(MMVQ_MAX_BATCH_SIZE, g_ggml_sycl_xmx_glu_max_cols) : MMVQ_MAX_BATCH_SIZE;
+    if (act->ne[1] > max_cols) {
         return false;
     }
 

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -103,6 +103,9 @@ int g_ggml_sycl_enable_vmm = 1;
 int g_ggml_sycl_enable_fusion = 1;
 int g_ggml_sycl_enable_esimd = 1;
 int g_ggml_sycl_mmvq_wide = 1;
+int g_ggml_sycl_enable_xmx_esimd = 1;
+int g_ggml_sycl_xmx_max_cols = GGML_SYCL_XMX_MAX_COLS_DEFAULT;
+int g_ggml_sycl_xmx_glu_max_cols = GGML_SYCL_XMX_GLU_MAX_COLS;
 int g_ggml_sycl_prioritize_dmmv = 0;
 int g_ggml_sycl_xmx_gather_types = GGML_SYCL_XMX_GATHER_TYPES_DEFAULT;
 int g_ggml_sycl_xmx_gather_shapes = GGML_SYCL_XMX_GATHER_SHAPES_DEFAULT;
@@ -424,6 +427,10 @@ static void ggml_check_sycl() try {
         g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1);
         g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
         g_ggml_sycl_mmvq_wide = ggml_sycl_get_env("GGML_SYCL_MMVQ_WIDE", 1);
+        g_ggml_sycl_enable_xmx_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_XMX_ESIMD", 1);
+        g_ggml_sycl_xmx_max_cols = std::max(0, ggml_sycl_get_env("GGML_SYCL_XMX_MAX_COLS", GGML_SYCL_XMX_MAX_COLS_DEFAULT));
+        g_ggml_sycl_xmx_glu_max_cols =
+            std::clamp(ggml_sycl_get_env("GGML_SYCL_XMX_GLU_MAX_COLS", GGML_SYCL_XMX_GLU_MAX_COLS), 0, GGML_SYCL_XMX_GLU_MAX_COLS);
         g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
         g_ggml_sycl_xmx_gather_types = ggml_sycl_get_env("GGML_SYCL_XMX_GATHER_TYPES", GGML_SYCL_XMX_GATHER_TYPES_DEFAULT);
         g_ggml_sycl_xmx_gather_shapes = ggml_sycl_get_env("GGML_SYCL_XMX_GATHER_SHAPES", GGML_SYCL_XMX_GATHER_SHAPES_DEFAULT);
@@ -561,7 +568,13 @@ static void ggml_check_sycl() try {
 #else
         GGML_LOG_INFO("  GGML_SYCL_ENABLE_ESIMD: %d disabled by compile flag\n", g_ggml_sycl_enable_esimd);
 #endif
+
         GGML_LOG_INFO("  GGML_SYCL_MMVQ_WIDE: %d\n", g_ggml_sycl_mmvq_wide);
+
+        GGML_LOG_INFO("  GGML_SYCL_ENABLE_XMX_ESIMD: %d\n", g_ggml_sycl_enable_xmx_esimd);
+        GGML_LOG_INFO("  GGML_SYCL_XMX_MAX_COLS: %d\n", g_ggml_sycl_xmx_max_cols);
+        GGML_LOG_INFO("  GGML_SYCL_XMX_GLU_MAX_COLS: %d\n", g_ggml_sycl_xmx_glu_max_cols);
+
         GGML_LOG_INFO("  GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv);
 
         g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1);
@@ -5165,9 +5178,24 @@ static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_
            src0->ne[0] % dmmv_x_required == 0 && src1->ne[1] == 1;
 }
 
-static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+// reordered weights of the types in ggml_sycl_xmx_supports_type() take the XMX kernel for wider batches
+static bool can_use_xmx_batch(int device, const ggml_tensor * src0, const ggml_tensor * src1) {
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+    const auto * extra = static_cast<const ggml_tensor_extra_gpu *>(src0->extra);
+    return ggml_sycl_xmx_enabled(device) && ggml_sycl_xmx_supports_type(src0->type) && extra &&
+           extra->optimized_feature.reorder && src1->ne[1] <= g_ggml_sycl_xmx_max_cols && src1->ne[2] == 1 &&
+           src1->ne[3] == 1;
+#else
+    GGML_UNUSED(device);
+    GGML_UNUSED(src0);
+    GGML_UNUSED(src1);
+    return false;
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+}
+
+static bool can_use_mul_mat_vec_q(int device, 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 &&
-           src1->ne[1] <= MMVQ_MAX_BATCH_SIZE;
+           (src1->ne[1] <= MMVQ_MAX_BATCH_SIZE || can_use_xmx_batch(device, src0, src1));
 }
 
 static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
@@ -5207,12 +5235,15 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
 
     // check data types and tensor shapes for custom matrix multiplication kernels:
     bool use_dequantize_mul_mat_vec = can_use_dequantize_mul_mat_vec(src0, src1, dst);
+    if (src1->ne[1] == 1 && ggml_sycl_xmx_min_cols(src0->type) == 1 && can_use_xmx_batch(ctx.device, src0, src1)) {
+        use_dequantize_mul_mat_vec = false;
+    }
 #ifdef GGML_SYCL_F16
     // dmmv may convert src1 to f16 in this build [TAG_GGML_PREC]
     use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && ggml_sycl_src1_f16_ok(dst);
 #endif
 
-    bool use_mul_mat_vec_q = can_use_mul_mat_vec_q(src0, src1, dst);
+    bool use_mul_mat_vec_q = can_use_mul_mat_vec_q(ctx.device, src0, src1, dst);
 
     bool use_mul_mat_q =  ggml_sycl_supports_mmq(src0->type)
         && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32;
@@ -5321,6 +5352,35 @@ static bool ggml_sycl_mul_mat_glu_mmvq_plain(ggml_backend_sycl_context & ctx, gg
     GGML_UNUSED(gate);
 }
 
+// Gate and up weights of one type on the XMX engines: one launch reads the activations once and applies the GLU.
+// Both weights must already be in the reorder layout. Returns false if the kernel declined.
+static bool ggml_sycl_mul_mat_glu_xmx(ggml_backend_sycl_context & ctx,
+                                      ggml_tensor *               glu,
+                                      ggml_tensor *               up,
+                                      const ggml_tensor *         wu,
+                                      const ggml_tensor *         wg,
+                                      const ggml_tensor *         act) {
+    const int64_t ne00 = wu->ne[0];
+    const int64_t ne11 = act->ne[1];
+
+    const queue_ptr stream           = ctx.stream();
+    const int       src1_padded_cols = GGML_PAD((int) ne00, MATRIX_ROW_PADDING);
+    const int       stride_y_bytes   = src1_padded_cols * (int) sizeof(block_q8_1) / QK8_1;
+
+    // log the up mat-mul: glu's own srcs are the two intermediates the fusion never materialises
+    scope_op_debug_print scope_dbg_print(__func__, up, /*num_src=*/2, " : fused with gate + GLU (XMX)");
+
+    ggml_sycl_pool_alloc<char> src1_q8_alloc(ctx.pool(), (size_t) ne11 * stride_y_bytes);
+    char *                     src1_ddq = src1_q8_alloc.get();
+
+    quantize_row_q8_1_sycl<quantize_and_reorder_q8_1_soa>((const float *) act->data, src1_ddq, (int) ne00, (int) ne11,
+                                                          src1_padded_cols, stream);
+
+    return ggml_sycl_mul_mat_vec_q_glu_xmx(ctx.device, wu->type, ggml_get_glu_op(glu), wu->data, wg->data, src1_ddq,
+                                           (float *) glu->data, (int) ne00, (int) wu->ne[1], (int) ne11, stride_y_bytes,
+                                           /*stride_col_dst=*/(int) glu->ne[0], stream);
+}
+
 #ifdef GGML_SYCL_DMMV_HAS_ESIMD
 // {mul_mat(gate), mul_mat(up), GLU} on the reorder layout via the ESIMD DMMV kernel, for
 // same-type K-quant or Q8_0 pairs and a single activation column. It reads the F32 activation
@@ -5389,11 +5449,33 @@ static bool ggml_sycl_mul_mat_glu_fused(ggml_backend_sycl_context & ctx, ggml_cg
         return false;
     }
 
+    // gate and up of one type the XMX kernel handles, it needs the reorder layout like the unfused mmvq path
+    if (wg->type == wu->type && gate->src[1] == act && act->ne[1] >= ggml_sycl_xmx_min_cols(wu->type) &&
+        act->ne[1] <= g_ggml_sycl_xmx_glu_max_cols) {
+        opt_for_reorder(&ctx, wu, act, up, mul_mat_algo::MMVQ);
+        opt_for_reorder(&ctx, wg, act, gate, mul_mat_algo::MMVQ);
+        if (can_use_xmx_batch(ctx.device, wu, act) && can_use_xmx_batch(ctx.device, wg, act) &&
+            ggml_sycl_mul_mat_glu_xmx(ctx, glu, up, wu, wg, act)) {
+            return true;
+        }
+    }
+
+    // the kernels below take up to MMVQ_MAX_BATCH_SIZE columns, a wider batch got here for the XMX kernel
+    if (act->ne[1] > MMVQ_MAX_BATCH_SIZE) {
+        return false;
+    }
+
     // quant pairs the reorder kernel does not serve (mixed gate/up types, q5_K off BMG) take the
     // standard-layout fused path instead; same-type q4_K / q5_K keep the reorder path below
     const bool reorder_pair = wg->type == wu->type &&
         (wu->type == GGML_TYPE_Q4_K || (wu->type == GGML_TYPE_Q5_K && ggml_sycl_q5_k_mmvq_reuse(ctx.device)));
     if (!reorder_pair) {
+        const auto is_plain_type = [](ggml_type type) {
+            return type == GGML_TYPE_Q5_K || type == GGML_TYPE_IQ4_XS;
+        };
+        if (!is_plain_type(wu->type) || !is_plain_type(wg->type)) {
+            return false;
+        }
         return ggml_sycl_mul_mat_glu_mmvq_plain(ctx, glu, gate, up, wu, wg, act);
     }
 

ggml/src/ggml-sycl/mmvq.cpp

@@ -6,6 +6,10 @@
 #include "quants.hpp"
 #include "vecdotq.hpp"
 
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+#    include "esimd.hpp"
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+
 // vec_dot_q_sycl_t adapters for the IQ vec_dots that take their codebook tables as extra
 // arguments: bind the constant tables here (as vec_dot_iq2_s_q8_1 / vec_dot_iq1_m_q8_1 already do
 // internally) so they can be used as template arguments of mul_mat_vec_q_moe.
@@ -2557,6 +2561,158 @@ static void mul_mat_vec_iq4_xs_q8_1_sycl_switch_ncols(
     }
 }
 
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+// vgate is the gate weight of the fused gate and up kernel, nullptr for a plain mat-vec
+template <template <bool> class Traits>
+static void ggml_sycl_mul_mat_vec_q_xmx_run(bool            short_rows,
+                                            const void *    vx,
+                                            const void *    vgate,
+                                            const char *    vy,
+                                            float *         dst,
+                                            int             ne00,
+                                            int64_t         nrows,
+                                            int64_t         ncols_dst,
+                                            int             stride_y_bytes,
+                                            int             ldd,
+                                            int             glu_op,
+                                            dpct::queue_ptr stream) {
+    namespace esimd = ggml_sycl_esimd;
+    if (vgate) {
+        if (short_rows) {
+            esimd::xmx_mul_mat_ncols<Traits<true>, true>(vx, vgate, vy, dst, ne00, nrows, ncols_dst, stride_y_bytes,
+                                                         ldd, glu_op, stream);
+        } else {
+            esimd::xmx_mul_mat_ncols<Traits<false>, true>(vx, vgate, vy, dst, ne00, nrows, ncols_dst, stride_y_bytes,
+                                                          ldd, glu_op, stream);
+        }
+    } else if (short_rows) {
+        esimd::xmx_mul_mat_ncols<Traits<true>>(vx, nullptr, vy, dst, ne00, nrows, ncols_dst, stride_y_bytes, ldd, 0,
+                                               stream);
+    } else {
+        esimd::xmx_mul_mat_ncols<Traits<false>>(vx, nullptr, vy, dst, ne00, nrows, ncols_dst, stride_y_bytes, ldd, 0,
+                                                stream);
+    }
+}
+
+static bool ggml_sycl_mul_mat_vec_q_xmx_type(ggml_type       type,
+                                             bool            short_rows,
+                                             const void *    vx,
+                                             const void *    vgate,
+                                             const char *    vy,
+                                             float *         dst,
+                                             int             ne00,
+                                             int64_t         nrows,
+                                             int64_t         ncols_dst,
+                                             int             stride_y_bytes,
+                                             int             ldd,
+                                             int             glu_op,
+                                             dpct::queue_ptr stream) {
+#    define XMX_RUN(TRAITS)                                                                                       \
+        ggml_sycl_mul_mat_vec_q_xmx_run<ggml_sycl_esimd::TRAITS>(short_rows, vx, vgate, vy, dst, ne00, nrows,     \
+                                                                 ncols_dst, stride_y_bytes, ldd, glu_op, stream); \
+        return true
+    switch (type) {
+        case GGML_TYPE_Q2_K:
+            XMX_RUN(xmx_traits_q2_k);
+        case GGML_TYPE_Q3_K:
+            XMX_RUN(xmx_traits_q3_k);
+        case GGML_TYPE_Q4_K:
+            XMX_RUN(xmx_traits_q4_k);
+        case GGML_TYPE_Q5_K:
+            XMX_RUN(xmx_traits_q5_k);
+        case GGML_TYPE_Q6_K:
+            XMX_RUN(xmx_traits_q6_k);
+        case GGML_TYPE_Q8_0:
+            XMX_RUN(xmx_traits_q8_0);
+        default:
+            return false;
+    }
+#    undef XMX_RUN
+}
+
+// multi-column XMX path for reordered weights, returns false when it does not apply
+static bool ggml_sycl_mul_mat_vec_q_xmx(int                 device,
+                                        const ggml_tensor * src0,
+                                        const ggml_tensor * dst,
+                                        const char *        src0_dd_i,
+                                        const char *        src1_ddq_i,
+                                        float *             dst_dd_i,
+                                        int64_t             row_diff,
+                                        int64_t             src1_ncols,
+                                        int64_t             src1_padded_col_size,
+                                        dpct::queue_ptr     stream) {
+    const auto * extra = static_cast<const ggml_tensor_extra_gpu *>(src0->extra);
+    if (!ggml_sycl_xmx_enabled(device) || !ggml_sycl_xmx_supports_type(src0->type) || !extra ||
+        !extra->optimized_feature.reorder) {
+        return false;
+    }
+    if (src1_ncols < ggml_sycl_xmx_min_cols(src0->type) || src1_ncols > g_ggml_sycl_xmx_max_cols) {
+        return false;
+    }
+
+    const int ne00 = src0->ne[0];
+    const int ldd  = dst->ne[0];
+    if (!ggml_sycl_esimd::xmx_supported(src0_dd_i, ne00)) {
+        return false;
+    }
+
+    // rows under 64 bytes use the variant without 2D block loads, it is not built for the widest tile
+    const bool short_rows = ne00 < 4 * QK_K;
+    if (short_rows && src1_ncols > 8) {
+        return false;
+    }
+
+    const int stride_y_bytes = src1_padded_col_size * sizeof(block_q8_1) / QK8_1;
+    // the 2D block loads of the activations need a 64 byte aligned base and a pitch of 16 bytes
+    if ((uintptr_t) src1_ddq_i % 64 != 0 || stride_y_bytes % 16 != 0) {
+        return false;
+    }
+    return ggml_sycl_mul_mat_vec_q_xmx_type(src0->type, short_rows, src0_dd_i, nullptr, src1_ddq_i, dst_dd_i, ne00,
+                                            row_diff, src1_ncols, stride_y_bytes, ldd, 0, stream);
+}
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+
+bool ggml_sycl_mul_mat_vec_q_glu_xmx(int              device,
+                                     enum ggml_type   src0_type,
+                                     enum ggml_glu_op glu_op,
+                                     const void *     vx,
+                                     const void *     vgate,
+                                     const void *     vy,
+                                     float *          dst,
+                                     int              ncols,
+                                     int              nrows,
+                                     int              ncols_dst,
+                                     int              stride_col_y_bytes,
+                                     int              stride_col_dst,
+                                     dpct::queue_ptr  stream) {
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+    if (!ggml_sycl_xmx_enabled(device) || !ggml_sycl_xmx_supports_type(src0_type)) {
+        return false;
+    }
+    if (glu_op != GGML_GLU_OP_SWIGLU && glu_op != GGML_GLU_OP_GEGLU) {
+        return false;
+    }
+    if (ncols_dst < ggml_sycl_xmx_min_cols(src0_type) || ncols_dst > g_ggml_sycl_xmx_glu_max_cols) {
+        return false;
+    }
+    if (!ggml_sycl_esimd::xmx_supported(vx, ncols) || !ggml_sycl_esimd::xmx_supported(vgate, ncols) ||
+        (uintptr_t) vy % 64 != 0 || stride_col_y_bytes % 16 != 0) {
+        return false;
+    }
+    // rows under 64 bytes use the variant without 2D block loads, it is not built for the wide tiles
+    const bool short_rows = ncols < 4 * QK_K;
+    if (short_rows && ncols_dst > 8) {
+        return false;
+    }
+    return ggml_sycl_mul_mat_vec_q_xmx_type(src0_type, short_rows, vx, vgate, (const char *) vy, dst, ncols, nrows,
+                                            ncols_dst, stride_col_y_bytes, stride_col_dst, (int) glu_op, stream);
+#else
+    GGML_UNUSED_VARS(device, src0_type, glu_op, vx, vgate, vy, dst, ncols, nrows, ncols_dst, stride_col_y_bytes,
+                     stride_col_dst, stream);
+    return false;
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+}
+
 void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1,
                                 ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i,
                                 const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low,
@@ -2568,6 +2724,13 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens
     const int64_t ne00     = src0->ne[0];
     const int64_t row_diff = row_high - row_low;
 
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+    if (ggml_sycl_mul_mat_vec_q_xmx(ctx.device, src0, dst, src0_dd_i, src1_ddq_i, dst_dd_i, row_diff, src1_ncols,
+                                    src1_padded_col_size, stream)) {
+        return;
+    }
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+
     int id;
     SYCL_CHECK(CHECK_TRY_ERROR(id = get_current_device_id()));
     const size_t q8_1_ts = sizeof(block_q8_1);

ggml/src/ggml-sycl/mmvq.hpp

@@ -80,6 +80,24 @@ bool ggml_sycl_mul_mat_vec_q_glu_reorder(
     dpct::queue_ptr    stream);
 
 
+// Fused dense-FFN GEMV on the XMX engines for reordered weights of one type, up * act(gate) for up to
+// GGML_SYCL_XMX_GLU_MAX_COLS (at most) activation columns. Returns false if it does not apply, the caller falls back.
+// vy: src1 quantized with quantize_and_reorder_q8_1_soa.
+bool ggml_sycl_mul_mat_vec_q_glu_xmx(
+    int                device,
+    enum ggml_type     src0_type,
+    enum ggml_glu_op   glu_op,
+    const void *       vx,                   // up weights
+    const void *       vgate,                // gate weights
+    const void *       vy,
+    float *            dst,
+    int                ncols,                // K, shared by both weights
+    int                nrows,                // output rows, i.e. weight ne[1]
+    int                ncols_dst,            // activation columns
+    int                stride_col_y_bytes,   // bytes between activation columns in vy
+    int                stride_col_dst,       // floats between output columns in dst
+    dpct::queue_ptr    stream);
+
 // Fused dense-FFN GEMV + GLU over the standard (non-reorder) layout; the gate and up
 // weights may carry different block types (q5_K / iq4_xs, mixed included).
 // vy: src1 quantized with plain quantize_q8_1 (padded rows). stride_col_y is in

ggml/src/ggml-vulkan/ggml-vulkan-buffers.cpp

@@ -9,7 +9,11 @@ static void * portable_aligned_malloc(size_t align, size_t size) {
 #ifdef _WIN32
     return _aligned_malloc(size, align);
 #else
-    return std::aligned_alloc(align, size);
+    void * ptr = nullptr;
+    if (posix_memalign(&ptr, align, size) != 0) {
+        return nullptr;
+    }
+    return ptr;
 #endif
 }
 

gguf-py/gguf/constants.py

@@ -383,6 +383,8 @@ class Keys:
         NORMALIZER_NFC           = "tokenizer.ggml.normalizer.nfc"
         NORMALIZER_LOWERCASE     = "tokenizer.ggml.normalizer.lowercase"
         NORMALIZER_STRIP_ACCENTS = "tokenizer.ggml.normalizer.strip_accents"
+        # Metaspace constants
+        METASPACE_PREPEND_SCHEME = "tokenizer.ggml.metaspace.prepend_scheme"
         # FIM/Infill special tokens constants
         FIM_PRE_ID           = "tokenizer.ggml.fim_pre_token_id"
         FIM_SUF_ID           = "tokenizer.ggml.fim_suf_token_id"
@@ -6450,6 +6452,7 @@ class VisionProjectorType:
     MINIMAXM3      = "minimax_m3"
     MINICPMV4_6    = "minicpmv4_6"
     MINICPMV4_7    = "minicpmv4_7"
+    MINICPM_OCR_1_0 = "minicpm_ocr_1_0"
     GRANITE_SPEECH = "granite_speech"  # audio
     MIMOVL         = "mimovl"
     MIMO_AUDIO     = "mimo_audio"

gguf-py/gguf/gguf_writer.py

@@ -1377,6 +1377,9 @@ class GGUFWriter:
     def add_normalizer_strip_accents(self, value: bool) -> None:
         self.add_bool(Keys.Tokenizer.NORMALIZER_STRIP_ACCENTS, value)
 
+    def add_metaspace_prepend_scheme(self, value: str) -> None:
+        self.add_string(Keys.Tokenizer.METASPACE_PREPEND_SCHEME, value)
+
     def add_eot_token_id(self, id: int) -> None:
         self.add_uint32(Keys.Tokenizer.EOT_ID, id)
 

gguf-py/gguf/vocab.py

@@ -56,6 +56,7 @@ class SpecialVocab:
     normalizer_lowercase: bool | None
     normalizer_strip_accents: bool | None
     clean_spaces: bool | None
+    metaspace_prepend_scheme: str | None
 
     def __init__(
         self, path: str | os.PathLike[str], load_merges: bool = False,
@@ -72,6 +73,7 @@ class SpecialVocab:
         self.normalizer_lowercase = None
         self.normalizer_strip_accents = None
         self.clean_spaces = None
+        self.metaspace_prepend_scheme = None
         if special_token_types is not None:
             self.special_token_types = special_token_types
         else:
@@ -126,6 +128,10 @@ class SpecialVocab:
             if not quiet:
                 logger.info(f'Setting clean_spaces to {self.clean_spaces}')
             gw.add_clean_spaces(self.clean_spaces)
+        if self.metaspace_prepend_scheme is not None:
+            if not quiet:
+                logger.info(f'Setting metaspace_prepend_scheme to {self.metaspace_prepend_scheme}')
+            gw.add_metaspace_prepend_scheme(self.metaspace_prepend_scheme)
 
     def _load(self, path: Path) -> None:
         self._try_load_from_tokenizer_json(path)
@@ -194,6 +200,16 @@ class SpecialVocab:
             for norm in normalizer.get('normalizers', []):
                 self._parse_normalizer(norm)
 
+    def _parse_pre_tokenizer(self, pre_tokenizer: dict) -> None:
+        # ref: https://huggingface.co/docs/tokenizers/api/pre-tokenizers
+        pre_tokenizer_type = pre_tokenizer.get('type')
+        if pre_tokenizer_type == 'Metaspace':
+            if 'prepend_scheme' in pre_tokenizer:
+                self.metaspace_prepend_scheme = pre_tokenizer['prepend_scheme']
+        elif pre_tokenizer_type == 'Sequence':
+            for pre in pre_tokenizer.get('pretokenizers', []):
+                self._parse_pre_tokenizer(pre)
+
     def _try_load_from_tokenizer_json(self, path: Path) -> bool:
         tokenizer = None
         tokenizer_file = path / 'tokenizer.json'
@@ -229,6 +245,8 @@ class SpecialVocab:
             # Parse normalizer configuration (e.g. Lowercase) into metadata
             if normalizer := tokenizer.get('normalizer'):
                 self._parse_normalizer(normalizer)
+            if pre_tokenizer := tokenizer.get('pre_tokenizer'):
+                self._parse_pre_tokenizer(pre_tokenizer)
             added_tokens = tokenizer.get('added_tokens', {})
         else:
             added_tokens = {}

src/llama-arch.cpp

@@ -445,6 +445,7 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
     { LLM_KV_TOKENIZER_NORMALIZER_NFC,           "tokenizer.ggml.normalizer.nfc"           },
     { LLM_KV_TOKENIZER_NORMALIZER_LOWERCASE,     "tokenizer.ggml.normalizer.lowercase"     },
     { LLM_KV_TOKENIZER_NORMALIZER_STRIP_ACCENTS, "tokenizer.ggml.normalizer.strip_accents" },
+    { LLM_KV_TOKENIZER_METASPACE_PREPEND_SCHEME, "tokenizer.ggml.metaspace.prepend_scheme" },
     { LLM_KV_TOKENIZER_FIM_PRE_ID,               "tokenizer.ggml.fim_pre_token_id"         },
     { LLM_KV_TOKENIZER_FIM_SUF_ID,               "tokenizer.ggml.fim_suf_token_id"         },
     { LLM_KV_TOKENIZER_FIM_MID_ID,               "tokenizer.ggml.fim_mid_token_id"         },

src/llama-arch.h

@@ -408,6 +408,7 @@ enum llm_kv {
     LLM_KV_TOKENIZER_NORMALIZER_NFC,
     LLM_KV_TOKENIZER_NORMALIZER_LOWERCASE,
     LLM_KV_TOKENIZER_NORMALIZER_STRIP_ACCENTS,
+    LLM_KV_TOKENIZER_METASPACE_PREPEND_SCHEME,
     LLM_KV_TOKENIZER_FIM_PRE_ID,
     LLM_KV_TOKENIZER_FIM_SUF_ID,
     LLM_KV_TOKENIZER_FIM_MID_ID,

src/llama-memory-hybrid-idx.cpp

@@ -758,7 +758,7 @@ void llama_memory_hybrid_idx_context::set_input_kpool(ggml_tensor * pool_cells,
     GGML_ASSERT(tail_idxs->ne[0] == (int64_t) kpool - 1 && tail_idxs->ne[1] == (int64_t) n_tokens);
     GGML_ASSERT(pool_idxs->ne[0] == (int64_t) kpool && pool_idxs->ne[1] == (int64_t) n_pool);
     GGML_ASSERT(ggml_backend_buffer_is_host(new_pool_idxs->buffer));
-    GGML_ASSERT(new_pool_idxs->ne[0] == (int64_t) kpool && new_pool_idxs->ne[1] == (int64_t) n_new_g);
+    GGML_ASSERT(ggml_nelements(new_pool_idxs) == (int64_t) kpool*n_new_g);
     // the graph always scatters the fresh pooled keys back into the cache, see build_qsa_sel
     GGML_ASSERT(new_pool_rep != nullptr && ggml_backend_buffer_is_host(new_pool_rep->buffer));
     GGML_ASSERT(new_pool_rep->ne[0] == (int64_t) n_new_g);

src/llama-model.cpp

@@ -1968,12 +1968,20 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
             if (buf == nullptr) {
                 throw std::runtime_error(format("unable to allocate %s buffer", ggml_backend_buft_name(buft)));
             }
-            // a multi_buffer has no base to lock, its parts are usually pinned memory (host buffer types with a max size)
-            if (use_mlock && ggml_backend_buffer_is_host(buf) && ggml_backend_buffer_get_base(buf) != nullptr) {
-                pimpl->mlock_bufs.emplace_back(new llama_mlock);
-                auto & mlock_buf = pimpl->mlock_bufs.back();
-                mlock_buf->init   (ggml_backend_buffer_get_base(buf));
-                mlock_buf->grow_to(ggml_backend_buffer_get_size(buf));
+            // lock the buffers that hold the tensors: a multi_buffer has no base, and its parts are pageable memory when pinning fails
+            if (use_mlock && !ml.no_alloc && ggml_backend_buffer_is_host(buf)) {
+                std::vector<ggml_backend_buffer_t> lock_bufs;
+                for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) {
+                    if (t->buffer != nullptr && std::find(lock_bufs.begin(), lock_bufs.end(), t->buffer) == lock_bufs.end()) {
+                        lock_bufs.push_back(t->buffer);
+                    }
+                }
+                for (ggml_backend_buffer_t lock_buf : lock_bufs) {
+                    pimpl->mlock_bufs.emplace_back(new llama_mlock);
+                    auto & mlock_buf = pimpl->mlock_bufs.back();
+                    mlock_buf->init   (ggml_backend_buffer_get_base(lock_buf));
+                    mlock_buf->grow_to(ggml_backend_buffer_get_size(lock_buf));
+                }
             }
             bufs.emplace_back(buf);
         }

src/llama-vocab.cpp

@@ -543,6 +543,7 @@ struct llm_tokenizer_bpe : llm_tokenizer {
                 byte_encode = false; // uses raw UTF-8, not GPT-2 byte encoding
                 break;
             case LLAMA_VOCAB_PRE_TYPE_SARVAM_MOE:
+            case LLAMA_VOCAB_PRE_TYPE_DOMYN_SMALL:
                 // Sarvam uses SPM-style BPE (same shape as Gemma4): spaces replaced with U+2581
                 // by the normalizer, BPE merges over the whole text on raw UTF-8.
                 regex_exprs = {
@@ -2097,6 +2098,8 @@ struct llama_vocab::impl {
     // Tokenizer normalizer options
     llama_vocab::normalizer_options normalizer_opts;
 
+    llama_metaspace_prepend_scheme metaspace_prepend_scheme = LLAMA_METASPACE_PREPEND_SCHEME_NEVER;
+
     std::unordered_map<std::string, llama_token> token_to_id;
     std::vector<token_data>                      id_to_token;
 
@@ -2612,6 +2615,10 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
             } else if (
                 tokenizer_pre == "minerva-7b") {
                 pre_type = LLAMA_VOCAB_PRE_TYPE_MINERVA;
+            } else if (
+                tokenizer_pre == "domyn-small") {
+                pre_type = LLAMA_VOCAB_PRE_TYPE_DOMYN_SMALL;
+                escape_whitespaces = true;
             } else if (
                 tokenizer_pre == "megrez") {
                 pre_type = LLAMA_VOCAB_PRE_TYPE_QWEN2;
@@ -2956,6 +2963,24 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
         normalizer_opts.strip_accents = normalizer_opts.lowercase;
         ml.get_key(LLM_KV_TOKENIZER_NORMALIZER_STRIP_ACCENTS, normalizer_opts.strip_accents, false);
 
+        {
+            std::string scheme;
+            if (ml.get_key(LLM_KV_TOKENIZER_METASPACE_PREPEND_SCHEME, scheme, false)) {
+                if (scheme == "never") {
+                    metaspace_prepend_scheme = LLAMA_METASPACE_PREPEND_SCHEME_NEVER;
+                } else if (scheme == "first") {
+                    metaspace_prepend_scheme = LLAMA_METASPACE_PREPEND_SCHEME_FIRST;
+                } else if (scheme == "always") {
+                    metaspace_prepend_scheme = LLAMA_METASPACE_PREPEND_SCHEME_ALWAYS;
+                } else {
+                    throw std::runtime_error(format("unknown metaspace prepend_scheme: '%s'", scheme.c_str()));
+                }
+                if (type == LLAMA_VOCAB_TYPE_BPE && metaspace_prepend_scheme == LLAMA_METASPACE_PREPEND_SCHEME_ALWAYS) {
+                    throw std::runtime_error("metaspace prepend_scheme 'always' is not supported for BPE vocab");
+                }
+            }
+        }
+
         // suppress tokens
         {
             const int suppress_idx = gguf_find_key(ctx, kv(LLM_KV_TOKENIZER_SUPPRESS_TOKENS).c_str());
@@ -3743,6 +3768,7 @@ std::vector<llama_token> llama_vocab::impl::tokenize(
                 // tokenizer.encode('', add_special_tokens=True)  returns [1]
                 // tokenizer.encode('', add_special_tokens=False) returns []
 
+                // TODO: honor metaspace_prepend_scheme (add_space_prefix currently behaves like "always")
                 bool is_prev_special = true;  // prefix with space if first token
 
                 if (add_special && add_bos) {
@@ -3813,6 +3839,13 @@ std::vector<llama_token> llama_vocab::impl::tokenize(
                             llama_escape_whitespace(text);
                         }
 
+                        // Metaspace prepend_scheme "first": add U+2581 only at the start of the input
+                        static const std::string escaped_space = "\xE2\x96\x81";
+                        if (metaspace_prepend_scheme == LLAMA_METASPACE_PREPEND_SCHEME_FIRST &&
+                            &fragment == &fragment_buffer.front() && text.rfind(escaped_space, 0) != 0) {
+                            text.insert(0, escaped_space);
+                        }
+
 #ifdef PRETOKENIZERDEBUG
                         LLAMA_LOG_WARN("TT: (%ld %ld %ld) '%s'\n", text.length(), fragment.offset, fragment.length, text.c_str());
 #endif
@@ -3860,6 +3893,7 @@ std::vector<llama_token> llama_vocab::impl::tokenize(
                     GGML_ASSERT(special_bos_id != LLAMA_TOKEN_NULL);
                     output.push_back(special_bos_id);
                 }
+                // TODO: honor metaspace_prepend_scheme
                 llm_tokenizer_ugm_session session(vocab, *static_cast<const llm_tokenizer_ugm *>(tokenizer.get()));
 
                 for (const auto & fragment : fragment_buffer) {
@@ -4143,10 +4177,17 @@ int32_t llama_vocab::impl::detokenize(
         }
     }
 
+    // Metaspace prepend_scheme "first": remove the leading space of the first rendered piece, so a skipped BOS does not consume it
+    // TODO: honor metaspace_prepend_scheme for the other vocab types
+    bool remove_first_space = type == LLAMA_VOCAB_TYPE_BPE && metaspace_prepend_scheme == LLAMA_METASPACE_PREPEND_SCHEME_FIRST;
+
     for (int32_t i = 0; i < n_tokens; ++i) {
         GGML_ASSERT(avail >= 0);
-        int32_t n_chars = token_to_piece(tokens[i], text, avail, remove_space, unparse_special);
+        int32_t n_chars = token_to_piece(tokens[i], text, avail, remove_space || remove_first_space, unparse_special);
         remove_space = false;
+        if (unparse_special || !(token_get_attr(tokens[i]) & (LLAMA_TOKEN_ATTR_UNKNOWN | LLAMA_TOKEN_ATTR_CONTROL))) {
+            remove_first_space = false;
+        }
         if (n_chars < 0) {
             avail = 0;
             total -= n_chars;

src/llama-vocab.h

@@ -71,6 +71,14 @@ enum llama_vocab_pre_type {
     LLAMA_VOCAB_PRE_TYPE_MMBERT            = 60,
     LLAMA_VOCAB_PRE_TYPE_K2_HORIZON        = 61,
     LLAMA_VOCAB_PRE_TYPE_FLINT             = 62,
+    LLAMA_VOCAB_PRE_TYPE_DOMYN_SMALL       = 63,
+};
+
+// HF Metaspace pre-tokenizer prepend_scheme
+enum llama_metaspace_prepend_scheme {
+    LLAMA_METASPACE_PREPEND_SCHEME_NEVER  = 0,
+    LLAMA_METASPACE_PREPEND_SCHEME_FIRST  = 1,
+    LLAMA_METASPACE_PREPEND_SCHEME_ALWAYS = 2,
 };
 
 struct LLM_KV;

src/models/clef.cpp

@@ -2,6 +2,7 @@
 
 #include "llama-ext.h"
 
+#include <algorithm>
 #include <cmath>
 
 void llama_model_clef::load_arch_hparams(llama_model_loader & ml) {

src/models/glm5-next.cpp

@@ -271,10 +271,8 @@ public:
     ggml_tensor * pool_mask     = nullptr; // F32/F16 [n_pool, n_tokens]
     ggml_tensor * tail_idxs     = nullptr; // I32     [kpool - 1, n_tokens]
     ggml_tensor * sel_mask      = nullptr; // F32     [n_sel, 1, 1, n_tokens] 0 for live selection slots, -inf for dead ones
-    ggml_tensor * sel_mask_2d   = nullptr; // F32     [n_sel, n_tokens] view of sel_mask, shared by the layers
     // n_new is never below 1, see build_inp_kpool
-    ggml_tensor * new_pool_idxs = nullptr; // I32     [kpool, n_new]   members of the pools completed this ubatch
-    ggml_tensor * new_pool_flat = nullptr; // I32     [kpool*n_new]    view of new_pool_idxs, shared by the layers
+    ggml_tensor * new_pool_idxs = nullptr; // I32     [kpool*n_new]    members of the pools completed this ubatch
     ggml_tensor * new_pool_rep  = nullptr; // I64     [n_new]          cell to write each new pooled key into
 
     const llama_memory_hybrid_idx_context * mctx;
@@ -329,16 +327,13 @@ llama_model_glm5_next::llm_graph_input_kpool * llama_model_glm5_next::graph::bui
             ggml_set_input(inp->sel_mask);
             // set_input_kpool always fills the mask, so it stays allocated in every graph
             ggml_build_forward_expand(gf, inp->sel_mask);
-            // one view for all layers: each view of a host input is a separate upload
-            inp->sel_mask_2d = ggml_reshape_2d(ctx0, inp->sel_mask, n_sel, n_tokens);
         }
     }
 
     inp->n_new = n_new;
 
-    inp->new_pool_idxs = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, kpool, n_new);
+    inp->new_pool_idxs = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, kpool*n_new);
     ggml_set_input(inp->new_pool_idxs);
-    inp->new_pool_flat = ggml_reshape_1d(ctx0, inp->new_pool_idxs, kpool*n_new);
     // one scatter row per new pool, each a distinct rep row (see kpool_build_state)
     inp->new_pool_rep = ggml_new_tensor_1d(ctx0, GGML_TYPE_I64, n_new);
     ggml_set_input(inp->new_pool_rep);
@@ -924,7 +919,7 @@ ggml_tensor * llama_model_glm5_next::graph::build_kpool_select(
     const int64_t n_kv = mctx_lid->get_n_kv();
 
     // Pool the entries completed by this ubatch. The last one is a dummy when the ubatch completes none.
-    ggml_tensor * rows = kpool_cache.gather_key_gate(inp_kpool->new_pool_flat);
+    ggml_tensor * rows = kpool_cache.gather_key_gate(inp_kpool->new_pool_idxs);
     rows = ggml_reshape_3d(ctx0, rows, 2*n_embd_indexer, kpool, n_new);
 
     ggml_tensor * pk = ggml_view_3d(ctx0, rows, n_embd_indexer, kpool, n_new, rows->nb[1], rows->nb[2], 0);
@@ -1023,7 +1018,7 @@ ggml_tensor * llama_model_glm5_next::graph::build_kpool_select(
     // Live slots (visible pools, real tail cells) address disjoint cells. Each dead slot writes its own dump row
     // n_kv + slot, so the scatter indices of a token are unique: idx = dump + live*(idx - dump), live = exp(mask).
     GGML_ASSERT(inp_kpool->sel_mask->ne[0] == n_sel && inp_kpool->sel_mask->ne[3] == n_tokens);
-    ggml_tensor * live  = ggml_exp(ctx0, inp_kpool->sel_mask_2d);
+    ggml_tensor * live  = ggml_reshape_2d(ctx0, ggml_exp(ctx0, inp_kpool->sel_mask), n_sel, n_tokens);
     ggml_tensor * dump  = ggml_arange(ctx0, (float) n_kv, (float) (n_kv + n_sel), 1.0f);
     ggml_tensor * idx_f = ggml_cast(ctx0, sel_idx, GGML_TYPE_F32);
     idx_f   = ggml_add(ctx0, ggml_mul(ctx0, ggml_sub(ctx0, idx_f, dump), live), dump);

src/models/nemotron.cpp

@@ -25,8 +25,11 @@ void llama_model_nemotron::load_arch_tensors(llama_model_loader &) {
         layer.attn_norm   = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
         layer.attn_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "bias", i), {n_embd}, 0);
 
-        create_tensor_qkv(layer, i, n_embd, n_embd, n_embd_gqa, n_embd_gqa, 0);
-        layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd, n_embd}, 0);
+        // n_embd_head_k * n_head equals n_embd for standard Nemotron (32 heads),
+        // but differs for models with non-standard GQA (e.g. 48 Q heads, head_dim=128).
+        const int64_t n_embd_q = n_embd_head_k * n_head;
+        create_tensor_qkv(layer, i, n_embd, n_embd_q, n_embd_gqa, n_embd_gqa, 0);
+        layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_q, n_embd}, 0);
 
         // optional bias tensors
         layer.wo_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);

src/models/qwen4exp.cpp

@@ -308,6 +308,20 @@ std::unique_ptr<llm_graph_context> llama_model_qwen4exp::build_arch_graph(const
     return std::make_unique<graph>(*this, params);
 }
 
+// DFlash drafters read the layer inputs as [n_embd, T] rows, so collapse the streams
+static ggml_tensor * hc_stream_mean(ggml_context * ctx, ggml_tensor * x) {
+    const int64_t n_embd = x->ne[0];
+    const int64_t hc     = x->ne[1];
+    const int64_t nt     = x->ne[2];
+
+    ggml_tensor * acc = ggml_view_2d(ctx, x, n_embd, nt, x->nb[2], 0);
+    for (int64_t s = 1; s < hc; ++s) {
+        acc = ggml_add(ctx, acc, ggml_view_2d(ctx, x, n_embd, nt, x->nb[2], s*x->nb[1]));
+    }
+
+    return ggml_scale(ctx, acc, 1.0f / (float) hc);
+}
+
 // Hyper-connections keep hc parallel residual streams [n_embd, hc, T] in place of layer norms.
 // Returns the mixed [n_embd, T] stream; `inject` gets the [hc, T] scatter weights.
 ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
@@ -447,7 +461,11 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
     ggml_build_forward_expand(gf, res_hc);
 
     for (int il = 0; il < n_layer; ++il) {
-        res->t_layer_inp[il] = res_hc;
+        if ((size_t) il < cparams.embeddings_layer_inp.size() && cparams.embeddings_layer_inp[il]) {
+            res->t_layer_inp[il] = hc_stream_mean(ctx0, res_hc);
+            cb(res->t_layer_inp[il], "layer_inp", il);
+            ggml_build_forward_expand(gf, res->t_layer_inp[il]);
+        }
 
         if (hparams.is_ple(il)) {
             res_hc = build_ple(inp->get_recr(), ple_emb, res_hc, il);

tests/test-backend-ops.cpp

@@ -12246,6 +12246,18 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
             true, 16, 8, b, false, true, false));
     }
 
+    // Gate and up on the SYCL XMX kernels: rows long enough for 2D block loads, a last tile of 8 rows, and
+    // columns on both sides of the 8 and 16 column tiles
+    for (ggml_type type : {GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K,
+                           GGML_TYPE_Q8_0}) {
+        for (ggml_glu_op glu_op : {GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU}) {
+            for (int64_t m_batch : {1, 2, 3, 5, 8, 9, 16, 17}) {
+                test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, m_batch, 1000, 1024,
+                    false, 16, 8, false, false, true, false, { 1, 1 }));
+            }
+        }
+    }
+
     for (int64_t k : {96, 896, 1024, 2880}) {
         for (ggml_glu_op glu_op : {GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU}) {
             test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, glu_op, 1, 32, k,

tests/test-server-component.cpp

@@ -435,6 +435,35 @@ static void test_decision_escape(testing & t) {
     });
 }
 
+static void test_rerank_template(testing & t) {
+    struct test_case {
+        const char * name;
+        std::string tmpl;
+        std::string query;
+        std::string doc;
+        std::string expected;
+    };
+    const test_case cases[] = {
+        {"plain inputs", "Q: {query}\nD: {document}", "query", "doc", "Q: query\nD: doc"},
+        {"document placeholder in query", "Q: {query}\nD: {document}", "Explain {document}.", "doc", "Q: Explain {document}.\nD: doc"},
+        {"query placeholder in query", "{query}|{document}", "{query}", "doc", "{query}|doc"},
+        {"placeholders in both inputs", "{query}|{document}", "{query} {document}", "{query} {document}", "{query} {document}|{query} {document}"},
+        {"repeated placeholders", "{query}|{document}|{query}|{document}", "{document}", "doc", "{document}|doc|{document}|doc"},
+        {"adjacent reversed placeholders", "{document}{query}", "query", "doc", "docquery"},
+        {"query only", "{query}", "{document}", "doc", "{document}"},
+        {"document only", "{document}", "query", "{query}", "{query}"},
+        {"no placeholders", "fixed", "{document}", "{query}", "fixed"},
+        {"unknown and incomplete placeholders", "{unknown}|{query}|{document}|{query", "query", "doc", "{unknown}|query|doc|{query"},
+        {"nested braces", "{{query}}|{{document}}", "query", "doc", "{query}|{doc}"},
+        {"placeholder formed at boundary", "{query}ument}", "{doc", "doc", "{document}"},
+        {"empty inputs", "Q: {query}\nD: {document}", "", "", "Q: \nD: "},
+        {"empty template", "", "query", "doc", ""},
+    };
+    for (const auto & test : cases) {
+        t.assert_equal(test.name, test.expected, format_rerank_template(test.tmpl, test.query, test.doc));
+    }
+}
+
 int main(int argc, char ** argv) {
     testing t;
 
@@ -454,6 +483,7 @@ int main(int argc, char ** argv) {
     t.test("lora_helpers",  test_lora_helpers);
     t.test("server_pipe",   test_server_pipe);
     t.test("decision_escape", test_decision_escape);
+    t.test("rerank_template", test_rerank_template);
 
     return t.summary();
 }

tools/mtmd/clip-impl.h

@@ -537,6 +537,7 @@ enum projector_type {
     PROJECTOR_TYPE_EXAONE4_5,
     PROJECTOR_TYPE_MINICPMV4_6,
     PROJECTOR_TYPE_MINICPMV4_7,
+    PROJECTOR_TYPE_MINICPM_OCR_1_0,
     PROJECTOR_TYPE_GRANITE_SPEECH,
     PROJECTOR_TYPE_MIMOVL,
     PROJECTOR_TYPE_MINIMAX_M3,
@@ -611,6 +612,7 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
     { PROJECTOR_TYPE_HUNYUANVL,         "hunyuanvl"},
     { PROJECTOR_TYPE_MINICPMV4_6,       "minicpmv4_6"},
     { PROJECTOR_TYPE_MINICPMV4_7,       "minicpmv4_7"},
+    { PROJECTOR_TYPE_MINICPM_OCR_1_0,   "minicpm_ocr_1_0"},
     { PROJECTOR_TYPE_GRANITE_SPEECH,    "granite_speech"},
     { PROJECTOR_TYPE_MIMOVL,            "mimovl"},
     { PROJECTOR_TYPE_MINIMAX_M3,        "minimax_m3"},

tools/mtmd/clip.cpp

@@ -1031,6 +1031,7 @@ static std::unique_ptr<clip_graph> clip_get_graph_builder(clip_ctx * ctx, const
             } break;
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             {
                 builder = std::make_unique<clip_graph_minicpmv4_6>(ctx, img);
             } break;
@@ -1496,6 +1497,7 @@ struct clip_model_loader {
                     } break;
                 case PROJECTOR_TYPE_MINICPMV4_6:
                 case PROJECTOR_TYPE_MINICPMV4_7:
+                case PROJECTOR_TYPE_MINICPM_OCR_1_0:
                     {
                         // MiniCPM-V 4.6/4.7 unified merger projector
                         // ViT merger 2x2 + final merger 2x2 = 4x spatial merge per dimension
@@ -2527,6 +2529,7 @@ struct clip_model_loader {
                     || model.proj_type == PROJECTOR_TYPE_MINICPMV
                     || model.proj_type == PROJECTOR_TYPE_MINICPMV4_6
                     || model.proj_type == PROJECTOR_TYPE_MINICPMV4_7
+                    || model.proj_type == PROJECTOR_TYPE_MINICPM_OCR_1_0
                 ) && layer.ff_up_w && layer.ff_down_w && layer.ff_down_w->ne[0] == hparams.n_embd;
             if (is_ffn_swapped) {
                 // swap up and down weights
@@ -2630,6 +2633,7 @@ struct clip_model_loader {
                 } break;
             case PROJECTOR_TYPE_MINICPMV4_6:
             case PROJECTOR_TYPE_MINICPMV4_7:
+            case PROJECTOR_TYPE_MINICPM_OCR_1_0:
                 {
                     const bool merger_required = hparams.n_merge == 4;
                     auto get_merger_tensor = [&](const std::string & name, bool required = true) {
@@ -4348,6 +4352,7 @@ int clip_n_output_tokens_x(const clip_ctx * ctx, const clip_image_f32 * img) {
         case PROJECTOR_TYPE_STEP3VL:
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             return img->nx() / (params.patch_size * params.n_merge);
         case PROJECTOR_TYPE_DEEPSEEKOCR:
         case PROJECTOR_TYPE_DEEPSEEKOCR2:
@@ -4379,6 +4384,7 @@ int clip_n_output_tokens_y(const clip_ctx * ctx, const clip_image_f32 * img) {
         case PROJECTOR_TYPE_STEP3VL:
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             return img->ny() / (params.patch_size * params.n_merge);
         default:
             break;
@@ -4445,6 +4451,7 @@ int clip_n_output_tokens(const clip_ctx * ctx, const clip_image_f32 * img) {
             } break;
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             {
                 n_patches /= params.n_merge * params.n_merge;
             } break;
@@ -5169,6 +5176,7 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
             } break;
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             {
                 const bool is_4x = hparams.n_merge == 2;
 
@@ -6437,6 +6445,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
             return ctx->model.mm_model_proj->ne[0];
         case PROJECTOR_TYPE_MINICPMV4_6:
         case PROJECTOR_TYPE_MINICPMV4_7:
+        case PROJECTOR_TYPE_MINICPM_OCR_1_0:
             return ctx->model.mm_ffn_down_w->ne[1];
         case PROJECTOR_TYPE_GLM_EDGE:
             return ctx->model.mm_model_mlp_3_w->ne[1];

tools/mtmd/mtmd-image.h

@@ -190,6 +190,17 @@ protected:
     }
 };
 
+// MiniCPM-OCR 1.0: find_best_resize aligns to patch_size * n_merge (28 in 4x).
+struct mtmd_image_preprocessor_minicpm_ocr_1_0 : mtmd_image_preprocessor_minicpmv {
+    using mtmd_image_preprocessor_minicpmv::mtmd_image_preprocessor_minicpmv;
+
+protected:
+    int get_slice_align() const override {
+        const int merge = hparams.n_merge > 0 ? hparams.n_merge : 2;
+        return hparams.patch_size * merge;
+    }
+};
+
 // custom llava-uhd slicing logic for LFM2
 // tiling params (min/max tiles, tile size, pixels tolerance) come from hparams
 // ref: https://github.com/huggingface/transformers/blob/v5.1.0/src/transformers/models/lfm2_vl/image_processing_lfm2_vl_fast.py

tools/mtmd/mtmd.cpp

@@ -722,6 +722,7 @@ struct mtmd_context {
                 } break;
             case PROJECTOR_TYPE_MINICPMV4_6:
             case PROJECTOR_TYPE_MINICPMV4_7:
+            case PROJECTOR_TYPE_MINICPM_OCR_1_0:
                 {
                     slice_tmpl        = MTMD_SLICE_TMPL_MINICPMV_2_6;
                     tok_ov_img_start  = {lookup_token("<image>")};
@@ -732,7 +733,11 @@ struct mtmd_context {
                     tok_row_end_trail = false; // no trailing end-of-row token
                     ov_img_first      = true;
                     use_image_id      = true;
-                    image_preproc     = std::make_unique<mtmd_image_preprocessor_minicpmv>(ctx_v);
+                    if (clip_get_projector_type(ctx_v) == PROJECTOR_TYPE_MINICPM_OCR_1_0) {
+                        image_preproc = std::make_unique<mtmd_image_preprocessor_minicpm_ocr_1_0>(ctx_v);
+                    } else {
+                        image_preproc = std::make_unique<mtmd_image_preprocessor_minicpmv>(ctx_v);
+                    }
                 } break;
             case PROJECTOR_TYPE_QWEN2VL:
             case PROJECTOR_TYPE_QWEN25VL:
@@ -3091,7 +3096,8 @@ struct mtmd_audio_stream {
     std::vector<float> out;     // rows of the last step
     std::vector<float> tok_embd;
 
-    int64_t n_rows = 0; // rows produced so far
+    int64_t n_rows = 0;         // rows produced so far
+    int64_t n_rows_flushed = 0; // rows produced once the audio up to the last flush is consumed
     bool pending_flush = false;
     bool ended = false;
 
@@ -3204,6 +3210,8 @@ int32_t mtmd_audio_stream_flush(mtmd_audio_stream * s, bool end) {
     // align to a whole row, then add the silence the model expects after the audio
     const int64_t n_align = (s->samples_per_tok - s->n_samples % s->samples_per_tok) % s->samples_per_tok;
     mtmd_audio_stream_add_silence(s, n_align + (int64_t) s->n_right_pad * s->samples_per_tok);
+    // one row per samples_per_tok, the left padding included
+    s->n_rows_flushed = s->n_samples / s->samples_per_tok;
     s->pending_flush = true;
     s->ended = end;
     return 0;
@@ -3353,3 +3361,11 @@ int32_t mtmd_audio_stream_step(mtmd_audio_stream * s, llama_token last_token) {
 const float * mtmd_audio_stream_get_output_embd(const mtmd_audio_stream * s) {
     return s->out.data();
 }
+
+int64_t mtmd_audio_stream_get_n_rows(const mtmd_audio_stream * s) {
+    return s->n_rows;
+}
+
+int64_t mtmd_audio_stream_get_n_rows_flushed(const mtmd_audio_stream * s) {
+    return s->n_rows_flushed;
+}

tools/mtmd/mtmd.h

@@ -412,6 +412,13 @@ MTMD_API int32_t mtmd_audio_stream_step(mtmd_audio_stream * stream, llama_token
 // rows of the last step, the reading size (in bytes) is n_rows * llama_model_n_embd_inp(model) * sizeof(float)
 MTMD_API const float * mtmd_audio_stream_get_output_embd(const mtmd_audio_stream * stream);
 
+// rows produced so far, the prefix rows included
+MTMD_API int64_t mtmd_audio_stream_get_n_rows(const mtmd_audio_stream * stream);
+
+// rows the stream produces up to the end of the last flush
+// once get_n_rows() reaches it, every row of the flushed audio has been produced
+MTMD_API int64_t mtmd_audio_stream_get_n_rows_flushed(const mtmd_audio_stream * stream);
+
 /////////////////////////////////////////
 // EXPERIMENTAL API for audio generation, subjected to breaking changes
 

tools/server/server-common.cpp

@@ -1945,6 +1945,26 @@ llama_tokens format_prompt_infill(
     return embd_inp;
 }
 
+std::string format_rerank_template(
+        const std::string & tmpl,
+        const std::string & query,
+        const std::string & doc) {
+    std::string prompt;
+    prompt.reserve(tmpl.size());
+    for (size_t pos = 0; pos < tmpl.size();) {
+        if (tmpl.compare(pos, 7, "{query}") == 0) {
+            prompt += query;
+            pos += 7;
+        } else if (tmpl.compare(pos, 10, "{document}") == 0) {
+            prompt += doc;
+            pos += 10;
+        } else {
+            prompt += tmpl[pos++];
+        }
+    }
+    return prompt;
+}
+
 server_tokens format_prompt_rerank(
         const struct llama_model * model,
         const struct llama_vocab * vocab,
@@ -1974,9 +1994,7 @@ server_tokens format_prompt_rerank(
             d += get_media_marker();
         }
 
-        std::string prompt = rerank_prompt;
-        string_replace_all(prompt, "{query}"   , q);
-        string_replace_all(prompt, "{document}", d  );
+        std::string prompt = format_rerank_template(rerank_prompt, q, d);
 
         if (query_files.empty() && doc_files.empty()) {
             server_tokens tokens = tokenize_input_subprompt(vocab, mctx, prompt, false, true, init_opt);

tools/server/server-common.h

@@ -592,6 +592,11 @@ llama_tokens format_prompt_infill(
         const bool spm_infill,
         const llama_tokens & tokens_prompt);
 
+std::string format_rerank_template(
+        const std::string & tmpl,
+        const std::string & query,
+        const std::string & doc);
+
 // format rerank task: [BOS]query[EOS][SEP]doc[EOS].
 // query_files/doc_files carry decoded media for multimodal (Qwen) inputs; when non-empty the
 // prompt is built with get_media_marker() sentinels and spliced via process_mtmd_prompt.

tools/server/server-context.cpp

@@ -25,6 +25,7 @@
 #include <algorithm>
 #include <cstddef>
 #include <cinttypes>
+#include <deque>
 #include <exception>
 #include <memory>
 #include <filesystem>
@@ -399,6 +400,7 @@ struct server_slot {
     // realtime transcription, see SERVER_TASK_TYPE_TRANSCRIBE_RT
     mtmd::audio_stream_ptr audio_stream;
     bool audio_wait = false; // the stream needs more audio before it can produce rows
+    std::deque<int64_t> audio_commit_rows; // end row of each non-final commit not yet completed
 
     server_slot_stats stats;
 
@@ -456,6 +458,7 @@ struct server_slot {
 
         audio_stream.reset();
         audio_wait = false;
+        audio_commit_rows.clear();
         inp_embd.clear();
     }
 
@@ -3071,7 +3074,11 @@ private:
                             if (task.params.control_action == SERVER_CONTROL_ACTION_AUDIO_APPEND) {
                                 mtmd_audio_stream_push(stream, task.params.control_audio.data(), task.params.control_audio.size());
                             } else {
-                                mtmd_audio_stream_flush(stream, task.params.control_action == SERVER_CONTROL_ACTION_AUDIO_END);
+                                const int32_t ret = mtmd_audio_stream_flush(stream, task.params.control_action == SERVER_CONTROL_ACTION_AUDIO_END);
+                                // a non-final commit completes its item once the rows of its audio are decoded
+                                if (ret == 0 && task.params.control_action == SERVER_CONTROL_ACTION_AUDIO_COMMIT) {
+                                    slot->audio_commit_rows.push_back(mtmd_audio_stream_get_n_rows_flushed(stream));
+                                }
                             }
                             slot->audio_wait = false;
                             res->success = true;
@@ -3402,6 +3409,22 @@ private:
             return true; // rows from a previous step, not decoded yet
         }
 
+        // every produced row is decoded and its delta queued, the commits ending there are complete
+        const int64_t n_rows = mtmd_audio_stream_get_n_rows(slot.audio_stream.get());
+        while (!slot.audio_commit_rows.empty() && slot.audio_commit_rows.front() <= n_rows) {
+            slot.audio_commit_rows.pop_front();
+            auto res = std::make_unique<server_task_result_cmpl_partial>();
+            res->id                   = slot.task->id;
+            res->index                = slot.task->index;
+            res->n_decoded            = slot.stats.n_gen;
+            res->n_prompt_tokens      = slot.task->n_tokens();
+            res->res_type             = slot.task->params.res_type;
+            res->oaicompat_model      = slot.task->params.oaicompat_model;
+            res->oaicompat_cmpl_id    = slot.task->params.oaicompat_cmpl_id;
+            res->audio_commit_drained = true;
+            queue_results.send(std::move(res));
+        }
+
         const int32_t n = mtmd_audio_stream_step(slot.audio_stream.get(), slot.sampled);
         if (n > 0) {
             const float * rows = mtmd_audio_stream_get_output_embd(slot.audio_stream.get());
@@ -6111,16 +6134,13 @@ void server_routes::init_routes() {
                     if (!post_control(SERVER_CONTROL_ACTION_AUDIO_COMMIT, {})) {
                         break;
                     }
-                    // the current item is the committed one, its transcript completes it and the next item opens
+                    // completed is sent once the rows of the committed audio are decoded (audio_commit_drained below)
                     std::lock_guard<std::mutex> lock(mutex_item);
                     send_event(json {
                         {"type",             "input_audio_buffer.committed"},
                         {"previous_item_id", item_prev.empty() ? json(nullptr) : json(item_prev)},
                         {"item_id",          item_id},
                     });
-                    send_completed();
-                    item_prev = item_id;
-                    item_id   = "item_" + random_string();
                 } else if (type == "input_audio_buffer.clear") {
                     send_event(json {{"type", "input_audio_buffer.cleared"}});
                 } else {
@@ -6144,7 +6164,13 @@ void server_routes::init_routes() {
                 break;
             }
             if (auto * partial = dynamic_cast<server_task_result_cmpl_partial *>(result.get())) {
-                if (!partial->content.text.empty()) {
+                if (partial->audio_commit_drained) {
+                    // complete the committed item and open the next one
+                    std::lock_guard<std::mutex> lock(mutex_item);
+                    send_completed();
+                    item_prev = item_id;
+                    item_id   = "item_" + random_string();
+                } else if (!partial->content.text.empty()) {
                     std::lock_guard<std::mutex> lock(mutex_item);
                     item_text += partial->content.text;
                     send_event(json {

tools/server/server-task.h

@@ -470,6 +470,7 @@ struct server_task_result_cmpl_partial : server_task_result {
     bool is_progress = false;
     bool is_begin = false; // whether to send 200 status to HTTP client (begin of SSE stream)
                            // ref: https://github.com/ggml-org/llama.cpp/pull/23884
+    bool audio_commit_drained = false; // realtime: non-final commit has fully decoded
     completion_token_output prob_output;
     server_slot_stats stats;
     result_prompt_progress progress;