mtmd : add MiniCPM-OCR 1.0 support (#30979)

Signed-off-by: tc-mb <tianchi_cai@icloud.com>

ci/run.sh

@@ -447,10 +447,14 @@ function gg_run_qwen3_0_6b {
 
     (time ./bin/llama-imatrix --model ${model_f16} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-imatrix.log
 
-    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa off --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa off                ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+    # TODO: these tests are extremely slow on WebGPU - re-enable when fixed
+    #       https://github.com/ggml-org/llama.cpp/actions/runs/39639985732
+    if [ -z ${GG_BUILD_WEBGPU} ]; then
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa off --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa off                ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+    fi
 
     function check_ppl {
         qnt="$1"

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-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,

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:

tools/server/README.md

@@ -771,37 +771,81 @@ This endpoint also supports multimodal embeddings. See the documentation for the
 
 ### POST `/reranking`: Rerank documents according to a given query
 
+> [!IMPORTANT]
+>
+> This endpoint is **not** OAI-compatible.
 Similar to https://jina.ai/reranker/ but might change in the future.
-Requires a reranker model (such as [bge-reranker-v2-m3](https://huggingface.co/BAAI/bge-reranker-v2-m3)) and the `--embedding --pooling rank` options.
+Requires a reranker model (such as [bge-reranker-v2-m3](https://huggingface.co/BAAI/bge-reranker-v2-m3)) and the `--embedding --pooling rank` or `--reranking` options.
 
-*Options:*
+`query` and each `documents`/`texts` element accept a plain string, or
+(for multimodal rerankers such as Qwen3-VL-Reranker) an object of the form
+`{"text": string | [strings], "image": url | data-uri | raw-base64 | [urls]}`.
+All keys are optional; a plain string is treated as `{"text": "..."}`.
 
-`query`: The query against which the documents will be ranked.
-
-`documents`: An array strings representing the documents to be ranked.
+Options: `top_n` (default: all), `model` (echoed in the Jina response), `return_text` (TEI only).
 
 *Aliases:*
   - `/rerank`
   - `/v1/rerank`
   - `/v1/reranking`
 
-*Examples:*
+The format is selected by which key carries the documents: `texts` (TEI, same input shape as Jina,
+set `return_text: true` to echo document text) or `documents` (Jina, multimodal supported).
+
+**Text only (TEI format)** - `query` + `texts`; set `return_text: true` to echo document text:
 
 ```shell
 curl http://127.0.0.1:8012/v1/rerank \
     -H "Content-Type: application/json" \
     -d '{
-        "model": "some-model",
-            "query": "What is panda?",
-            "top_n": 3,
-            "documents": [
-                "hi",
-            "it is a bear",
-            "The giant panda (Ailuropoda melanoleuca), sometimes called a panda bear or simply panda, is a bear species endemic to China."
-            ]
+        "query": "What is machine learning?",
+        "texts": [
+            "Machine learning is a subset of artificial intelligence.",
+            "A python is a large snake."
+        ],
+        "return_text": true
     }' | jq
 ```
 
+**Multimodal (Jina Format)** - mix text and images in `query` and/or `documents`:
+
+```shell
+curl http://127.0.0.1:8012/v1/rerank \
+    -H "Content-Type: application/json" \
+    -d '{
+        "model": "Qwen/Qwen3-VL-Reranker-2B",
+        "query": {
+            "text": "A city skyline with tall buildings and a river",
+            "image": "data:image/jpeg;base64,<base64>"
+        },
+        "documents": [
+            "a plain string document is also valid",
+            {
+                "text": "A photograph of the Boston skyline at dusk",
+                "image": "https://example.com/skyline.jpg"
+            },
+            {
+                "text": "caption for the two-figure spread",
+                "image": ["https://example.com/fig1.png", "data:image/jpeg;base64,<base64>"]
+            }
+        ],
+        "top_n": 2
+    }' | jq
+```
+
+Jina format returns a wrapped object with `model`, `object`, `usage` (`prompt_tokens` is the sum
+of `tokens_evaluated` across all scored pairs), and a `results` array of `{index, relevance_score}`.
+TEI format returns a bare array of `{index, score}` (+ `text` when `return_text: true`).
+Both are sorted descending by score and truncated to `top_n`; the score is the match probability in `[0, 1]`.
+
+Images never appear in any response shape.
+
+> **NOTE:** **Input size limit:** A bidirectional-encoder reranker (e.g. `bge-reranker-v2-m3`) must fit the whole prompt within
+`--ubatch-size`; if the input is too large to process, increase `--ubatch-size`.
+Causal decoder rerankers (e.g. `Qwen3-VL-Reranker`) have no such limitation.
+
+
+
 ### POST `/tts`: Generate speech audio from text
 
 Returns raw audio bytes (`audio/wav` by default) rather than JSON. For more info, see [tts/README.md](../tts/README.md)

tools/server/server-common.cpp

@@ -1265,7 +1265,6 @@ void handle_media(
         out_files.push_back(decoded_data);
     }
 }
-
 // load media files from an OAI content array, then replace each media part with a media marker text part
 static void oaicompat_content_load_media(json & content, const server_chat_params & opt, std::vector<raw_buffer> & out_files) {
     for (auto & p : content) {
@@ -1336,6 +1335,7 @@ server_tokens tokenize_oai_content_array(const llama_vocab * vocab, mtmd_context
     if (files.empty()) {
         return server_tokens(common_tokenize(vocab, prompt, add_special, parse_special), false);
     }
+
     return process_mtmd_prompt(mctx, prompt, files, init_opt);
 }
 
@@ -1951,18 +1951,49 @@ server_tokens format_prompt_rerank(
         mtmd_context * mctx,
         const std::string & query,
         const std::string & doc,
+        const std::vector<raw_buffer> & query_files,
+        const std::vector<raw_buffer> & doc_files,
         const mtmd_helper_init_opt & init_opt) {
     server_tokens result = {};
 
+    // media must be spliced through the mtmd pipeline
+    if ((!query_files.empty() || !doc_files.empty()) && mctx == nullptr) {
+        throw std::invalid_argument("multimodal rerank input requires --mmproj");
+    }
+
     const char * rerank_prompt = llama_model_chat_template(model, "rerank");
 
     if (rerank_prompt != nullptr) {
+        // multimodal: append one media marker per image (arrays already flattened by caller)
+        std::string q = query;
+        for (size_t i = 0; i < query_files.size(); i++) {
+            q += get_media_marker();
+        }
+        std::string d = doc;
+        for (size_t i = 0; i < doc_files.size(); i++) {
+            d += get_media_marker();
+        }
+
         std::string prompt = rerank_prompt;
-        string_replace_all(prompt, "{query}"   , query);
-        string_replace_all(prompt, "{document}", doc  );
-        server_tokens tokens = tokenize_input_subprompt(vocab, mctx, prompt, false, true, init_opt);
-        result.push_back(tokens);
+        string_replace_all(prompt, "{query}"   , q);
+        string_replace_all(prompt, "{document}", d  );
+
+        if (query_files.empty() && doc_files.empty()) {
+            server_tokens tokens = tokenize_input_subprompt(vocab, mctx, prompt, false, true, init_opt);
+            result.push_back(tokens);
+        } else {
+            // process_mtmd_prompt returns a complete mtmd server_tokens; return it directly
+            // (push_back into a non-mtmd result would hit the NULL placeholder tokens)
+            std::vector<raw_buffer> files;
+            files.insert(files.end(), query_files.begin(), query_files.end());
+            files.insert(files.end(), doc_files.begin(),   doc_files.end());
+            return process_mtmd_prompt(mctx, prompt, files, init_opt);
+        }
     } else {
+        if (!query_files.empty() || !doc_files.empty()) {
+            throw std::invalid_argument("multimodal rerank input requires a model with a rerank chat template");
+        }
+
         // Get EOS token - use SEP token as fallback if EOS is not available
         server_tokens query_tokens = tokenize_input_subprompt(vocab, mctx, query, false, false, init_opt);
         server_tokens doc_tokens   = tokenize_input_subprompt(vocab, mctx, doc,   false, false, init_opt);

tools/server/server-common.h

@@ -303,8 +303,11 @@ server_tokens process_mtmd_prompt(
                                         mtmd_context * mctx,
                                         const std::string & prompt,
                                         const std::vector<raw_buffer> & files,
-                                        const mtmd_helper_init_opt & init_opt,
-                                        bool is_placeholder = false);
+                                         const mtmd_helper_init_opt & init_opt,
+                                          bool is_placeholder = false);
+
+// decode one media url (http/file/data-uri) into out_files; throws on failure
+void handle_media(std::vector<raw_buffer> & out_files, const std::string & url, const std::string & media_path);
 
 /**
  * break the input "prompt" object into multiple prompt if needed, then tokenize them
@@ -590,12 +593,16 @@ llama_tokens format_prompt_infill(
         const llama_tokens & tokens_prompt);
 
 // 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.
 server_tokens format_prompt_rerank(
         const struct llama_model * model,
         const struct llama_vocab * vocab,
         mtmd_context * mctx,
         const std::string & query,
         const std::string & doc,
+        const std::vector<raw_buffer> & query_files,
+        const std::vector<raw_buffer> & doc_files,
         const mtmd_helper_init_opt & init_opt);
 
 // simple implementation of a pipe

tools/server/server-context.cpp

@@ -6284,83 +6284,7 @@ void server_routes::init_routes() {
     };
 
     this->post_rerank = [this](const server_http_req & req) {
-        auto res = create_response();
-        if (!params.embedding || params.pooling_type != LLAMA_POOLING_TYPE_RANK) {
-            res->error(format_error_response("This server does not support reranking. Start it with `--reranking`", ERROR_TYPE_NOT_SUPPORTED));
-            return res;
-        }
-
-        const json body = json::parse(req.body);
-
-        // if true, use TEI API format, otherwise use Jina API format
-        // Jina: https://jina.ai/reranker/
-        // TEI: https://huggingface.github.io/text-embeddings-inference/#/Text%20Embeddings%20Inference/rerank
-        bool is_tei_format = body.contains("texts");
-
-        json query;
-        if (body.count("query") == 1) {
-            query = body.at("query");
-            if (!query.is_string()) {
-                res->error(format_error_response("\"query\" must be a string", ERROR_TYPE_INVALID_REQUEST));
-                return res;
-            }
-        } else {
-            res->error(format_error_response("\"query\" must be provided", ERROR_TYPE_INVALID_REQUEST));
-            return res;
-        }
-
-        std::vector<std::string> documents = json_value(body, "documents",
-                                             json_value(body, "texts", std::vector<std::string>()));
-        if (documents.empty()) {
-            res->error(format_error_response("\"documents\" must be a non-empty string array", ERROR_TYPE_INVALID_REQUEST));
-            return res;
-        }
-
-        int top_n = json_value(body, "top_n", (int)documents.size());
-
-        // create and queue the task
-        json responses = json::array();
-        auto & rd = res->rd;
-        {
-            std::vector<server_task> tasks;
-            tasks.reserve(documents.size());
-            for (size_t i = 0; i < documents.size(); i++) {
-                auto tmp = format_prompt_rerank(ctx_server.model_tgt, ctx_server.vocab, ctx_server.mctx, query, documents[i], ctx_server.init_opt);
-                server_task task = server_task(SERVER_TASK_TYPE_RERANK);
-                task.id     = rd.get_new_id();
-                task.tokens = std::move(tmp);
-                tasks.push_back(std::move(task));
-            }
-            rd.post_tasks(std::move(tasks));
-        }
-
-        // wait for the results
-        auto all_results = rd.wait_for_all(req.should_stop);
-
-        // collect results
-        if (all_results.is_terminated) {
-            return res; // connection is closed
-        } else if (all_results.error) {
-            res->error(all_results.error->to_json());
-            return res;
-        } else {
-            for (auto & res : all_results.results) {
-                GGML_ASSERT(dynamic_cast<server_task_result_rerank*>(res.get()) != nullptr);
-                responses.push_back(res->to_json());
-            }
-        }
-
-        // write JSON response
-        json root = format_response_rerank(
-            body,
-            meta->model_name,
-            responses,
-            is_tei_format,
-            documents,
-            top_n);
-
-        res->ok(root);
-        return res;
+        return handle_rerank_impl(req);
     };
 
     this->post_tts = [this](const server_http_req & req) {
@@ -6822,6 +6746,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_embeddings_impl(cons
         return tokenize_input_subprompt(ctx_server.vocab, ctx_server.mctx, p, true, true, ctx_server.init_opt);
     };
 
+    // dispatch to the appropriate tokenizer based on the shape of the input
     std::vector<server_tokens> tokenized_prompts;
     if (prompt.is_array() && !json_is_array_and_contains_numbers(prompt)) {
         for (const auto & p : prompt) {
@@ -6895,6 +6820,154 @@ std::unique_ptr<server_res_generator> server_routes::handle_embeddings_impl(cons
     return res;
 }
 
+std::unique_ptr<server_res_generator> server_routes::handle_rerank_impl(const server_http_req & req) {
+    auto res = create_response();
+    if (!params.embedding || params.pooling_type != LLAMA_POOLING_TYPE_RANK) {
+        res->error(format_error_response("This server does not support reranking. Start it with `--reranking`", ERROR_TYPE_NOT_SUPPORTED));
+        return res;
+    }
+
+    const json body = json::parse(req.body);
+
+    // if true, use TEI API format, otherwise use Jina API format
+    // Jina: https://jina.ai/reranker/
+    // TEI: https://huggingface.github.io/text-embeddings-inference/#/Text%20Embeddings%20Inference/rerank
+    bool is_tei_format = body.contains("texts");
+
+    // expand one rerank input item: plain string or {"text": str|[str], "image": url|data-uri|raw-base64|[urls]}
+    auto parse_item = [&](const json & item, std::string & out_text, std::vector<raw_buffer> & out_files) {
+        if (item.is_string()) {
+            out_text = item.get<std::string>();
+            return;
+        }
+        if (!item.is_object()) {
+            throw std::invalid_argument("rerank input must be a string or an object of the form {\"text\": ..., \"image\": ...}");
+        }
+        if (item.contains("video") || item.contains("fps") || item.contains("max_frames")) {
+            throw std::invalid_argument("video/fps/max_frames are not supported for rerank input");
+        }
+        json text = item.value("text", json());
+        if (text.is_string()) {
+            out_text += text.get<std::string>();
+        } else if (text.is_array()) {
+            for (const auto & t : text) {
+                if (!t.is_string()) {
+                    throw std::invalid_argument("\"text\" array elements must be strings");
+                }
+                out_text += t.get<std::string>();
+            }
+        } else if (!text.is_null()) {
+            throw std::invalid_argument("\"text\" must be a string or an array of strings");
+        }
+        json image = item.value("image", json());
+        auto add_image = [&](const std::string & url) {
+            // marker is appended by format_prompt_rerank, one per decoded file
+            try {
+                handle_media(out_files, url, params.media_path);
+            } catch (const std::runtime_error & e) {
+                throw std::invalid_argument(e.what());
+            }
+        };
+        if (image.is_string()) {
+            add_image(image.get<std::string>());
+        } else if (image.is_array()) {
+            for (const auto & u : image) {
+                if (!u.is_string()) {
+                    throw std::invalid_argument("\"image\" array elements must be strings");
+                }
+                add_image(u.get<std::string>());
+            }
+        } else if (!image.is_null()) {
+            throw std::invalid_argument("\"image\" must be a string or an array of strings");
+        }
+    };
+
+    // Parse query
+    std::string query_str;
+    std::vector<raw_buffer> query_files;
+    if (body.count("query") == 1) {
+        try {
+            parse_item(body.at("query"), query_str, query_files);
+        } catch (const std::invalid_argument & e) {
+            res->error(format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST));
+            return res;
+        }
+    } else {
+        res->error(format_error_response("\"query\" must be provided", ERROR_TYPE_INVALID_REQUEST));
+        return res;
+    }
+
+    // Resolve Jina/TEI format ambiguity
+    json documents_json = json_value(body, "documents",
+        json_value(body, "texts", json::array()));
+    if (!documents_json.is_array() || documents_json.empty()) {
+        res->error(format_error_response("\"documents\" must be a non-empty array of strings or objects", ERROR_TYPE_INVALID_REQUEST));
+        return res;
+    }
+
+    // Parse documents
+    std::vector<std::string>          documents(documents_json.size());
+    std::vector<std::vector<raw_buffer>> document_files(documents_json.size());
+    for (size_t i = 0; i < documents_json.size(); i++) {
+        try {
+            parse_item(documents_json[i], documents[i], document_files[i]);
+        } catch (const std::invalid_argument & e) {
+            res->error(format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST));
+            return res;
+        }
+    }
+
+    int top_n = json_value(body, "top_n", (int)documents.size());
+    if (top_n < 0) {
+        res->error(format_error_response("\"top_n\" must be >= 0", ERROR_TYPE_INVALID_REQUEST));
+        return res;
+    }
+
+    // create and queue the task
+    json responses = json::array();
+    auto & rd = res->rd;
+    {
+        std::vector<server_task> tasks;
+        tasks.reserve(documents.size());
+        for (size_t i = 0; i < documents.size(); i++) {
+            auto tmp = format_prompt_rerank(ctx_server.model_tgt, ctx_server.vocab, ctx_server.mctx, query_str, documents[i], query_files, document_files[i], ctx_server.init_opt);
+            server_task task = server_task(SERVER_TASK_TYPE_RERANK);
+            task.id     = rd.get_new_id();
+            task.tokens = std::move(tmp);
+            tasks.push_back(std::move(task));
+        }
+        rd.post_tasks(std::move(tasks));
+    }
+
+    // wait for the results
+    auto all_results = rd.wait_for_all(req.should_stop);
+
+    // collect results
+    if (all_results.is_terminated) {
+        return res; // connection is closed
+    } else if (all_results.error) {
+        res->error(all_results.error->to_json());
+        return res;
+    } else {
+        for (auto & res : all_results.results) {
+            GGML_ASSERT(dynamic_cast<server_task_result_rerank*>(res.get()) != nullptr);
+            responses.push_back(res->to_json());
+        }
+    }
+
+    // write JSON response
+    json root = format_response_rerank(
+        body,
+        meta->model_name,
+        responses,
+        is_tei_format,
+        documents,
+        top_n);
+
+    res->ok(root);
+    return res;
+}
+
 std::unique_ptr<server_res_generator> server_routes::handle_count_tokens(const server_http_req & req, task_response_type res_type) {
     auto res = create_response();
     std::vector<raw_buffer> files;

tools/server/server-context.h

@@ -177,6 +177,7 @@ private:
     std::unique_ptr<server_res_generator> handle_slots_restore(const server_http_req & req, int id_slot);
     std::unique_ptr<server_res_generator> handle_slots_erase(const server_http_req &, int id_slot);
     std::unique_ptr<server_res_generator> handle_embeddings_impl(const server_http_req & req, task_response_type res_type);
+    std::unique_ptr<server_res_generator> handle_rerank_impl(const server_http_req & req);
     std::unique_ptr<server_res_generator> handle_count_tokens(const server_http_req & req, task_response_type res_type);
 
     // using unique_ptr to allow late initialization of const

tools/server/tests/conftest.py

@@ -7,9 +7,9 @@ from utils import *
 @pytest.fixture(scope="session", autouse=True)
 def configure_worker_port(request):
     worker_id = getattr(request.config, "workerinput", {}).get("workerid", "master")
-    if worker_id != "master":
-        worker_num = int(worker_id[2:])
-        os.environ["PORT"] = str(8080 + worker_num * 10)
+    worker_num = 0 if worker_id == "master" else int(worker_id[2:])
+    port_base = int(os.environ.get("LLAMA_SERVER_TEST_PORT_BASE", "8080"))
+    os.environ["PORT"] = str(port_base + worker_num * 10)
 
 
 # ref: https://stackoverflow.com/questions/22627659/run-code-before-and-after-each-test-in-py-test