ggml : add aarch64 HWCAP fallbacks and fix fp16 variant detection (#25554)

* ggml : add fallback definitions for missing aarch64 HWCAP bits

* ggml : require HWCAP_ASIMDHP for the aarch64 fp16 cpu variants

Also rename has_fp16_va to has_fp16, the field gates the whole FEAT_FP16
extension, scalar and vector half-precision arithmetic together.

.github/workflows/build-self-hosted.yml

@@ -71,6 +71,26 @@ jobs:
           nvidia-smi
           GG_BUILD_CUDA=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
 
+  gpu-rocm:
+    runs-on: [self-hosted, Linux, AMD]
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+
+      - name: Test
+        id: ggml-ci
+        # HIP_LAUNCH_BLOCKING=1: workaround for an async-execution correctness
+        # issue on integrated RDNA3.5 (gfx1151) where batched inference returns
+        # incorrect output (perplexity ~88 vs ~9.4). Serializing kernel launches
+        # restores correctness. Remove once the underlying ROCm/HIP issue is fixed.
+        env:
+          HIP_LAUNCH_BLOCKING: "1"
+        run: |
+          rocminfo
+          GG_BUILD_ROCM=1 GG_BUILD_AMDGPU_TARGETS=gfx1151 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+
   gpu-vulkan-nvidia-cm:
     runs-on: [self-hosted, Linux, NVIDIA]
 

SECURITY.md

@@ -21,11 +21,18 @@ Please disclose it as a private [security advisory](https://github.com/ggml-org/
 
 A team of volunteers on a reasonable-effort basis maintains this project. As such, please give us at least 90 days to work on a fix before public exposure.
 
+### AI-powered code scan
+
+llama.cpp has an AI security scanner that scans the code periodically. The full prompts and tool set can be found in [ggml-org/security-scan-prompt](https://github.com/ggml-org/security-scan-prompt).
+
+We greatly appreciate reports that reflect genuine research effort, and we are happy to spend our time reviewing them. Findings that an autonomous AI agent can surface on its own add little on top of the scans we already run.
+
 ### Requirements
 
 Before submitting your report, ensure you meet the following requirements:
 
 - You have read this policy and fully understand it.
+- You have searched for existing discussions of the issue. If it has already been reported, your report will likely be rejected as a duplicate.
 - AI is only permitted in an assistive capacity as stated in [AGENTS.md](AGENTS.md). We do not accept reports that are written exclusively by AI.
 - Your report must include a working Proof-of-Concept in the form of a script and/or attached files.
 
@@ -46,6 +53,8 @@ Only vulnerabilities that fall within these parts of the project are considered
 
 Note that none of the topics under [Using llama.cpp securely](#using-llamacpp-securely) are considered vulnerabilities in LLaMA C++.
 
+Denial-of-Service (DoS) bugs are generally not treated as vulnerabilities. We don't reject them outright, but we look at them case-by-case and only accept those that are genuinely worth fixing.
+
 For vulnerabilities that fall within the `vendor` directory, please report them directly to the third-party project.
 
 ## Using llama.cpp securely

ci/run.sh

@@ -10,6 +10,9 @@
 # # with CUDA support
 # GG_BUILD_CUDA=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt
 #
+# # with ROCm support
+# GG_BUILD_ROCM=1 GG_BUILD_AMDGPU_TARGETS=gfx1151 bash ./ci/run.sh ./tmp/results ./tmp/mnt
+#
 # # with SYCL support
 # GG_BUILD_SYCL=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt
 #
@@ -89,7 +92,7 @@ if [ ! -z ${GG_BUILD_CUDA} ]; then
 fi
 
 if [ ! -z ${GG_BUILD_ROCM} ]; then
-    CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_HIP=ON"
+    CMAKE_EXTRA="${CMAKE_EXTRA} -DCMAKE_HIP_COMPILER=$(hipconfig -l)/clang -DGGML_HIP=ON -DGGML_HIP_ROCWMMA_FATTN=ON"
     if [ -z ${GG_BUILD_AMDGPU_TARGETS} ]; then
         echo "Missing GG_BUILD_AMDGPU_TARGETS, please set it to your GPU architecture (e.g. gfx90a, gfx1100, etc.)"
         exit 1
@@ -640,39 +643,52 @@ function gg_sum_rerank_tiny {
 
 function gg_check_build_requirements {
     if ! command -v git &> /dev/null; then
-        gg_printf 'git not found, please install'
+        gg_printf 'git not found, please install\n'
+        exit 1
     fi
 
     if ! command -v git-lfs &> /dev/null; then
-        gg_printf 'git-lfs not found, please install'
+        gg_printf 'git-lfs not found, please install\n'
+        exit 1
+    fi
+
+    if ! git config --get filter.lfs.clean &> /dev/null; then
+        gg_printf 'git-lfs not initialized, please run `git lfs install`\n'
+        exit 1
     fi
 
     if ! command -v wget &> /dev/null; then
-        gg_printf 'wget not found, please install'
+        gg_printf 'wget not found, please install\n'
+        exit 1
     fi
 
     if ! command -v python3 &> /dev/null; then
-        gg_printf 'python3 not found, please install'
+        gg_printf 'python3 not found, please install\n'
+        exit 1
     fi
 
     if ! command -v pip3 &> /dev/null; then
-        gg_printf 'pip3 not found, please install'
+        gg_printf 'pip3 not found, please install\n'
+        exit 1
     fi
 
     if ! python3 -m ensurepip --help &> /dev/null; then
-        gg_printf 'ensurepip not found, please install python3-venv package'
+        gg_printf 'ensurepip not found, please install python3-venv package\n'
+        exit 1
     fi
 
     if ! command -v cmake &> /dev/null; then
-        gg_printf 'cmake not found, please install'
+        gg_printf 'cmake not found, please install\n'
+        exit 1
     fi
 
     if ! command -v ccache &> /dev/null; then
-        gg_printf 'ccache not found, please consider installing for faster builds'
+        gg_printf 'ccache not found, please consider installing for faster builds\n'
     fi
 
     if ! command -v ctest &> /dev/null; then
-        gg_printf 'ctest not found, please install'
+        gg_printf 'ctest not found, please install\n'
+        exit 1
     fi
 }
 

common/fit.cpp

@@ -136,7 +136,10 @@ static std::vector<llama_device_memory_data> common_get_device_memory_data_impl(
         devs.push_back(llama_model_get_device(model, i));
     }
 
-    hp_ngl         = llama_model_n_layer(model) + llama_model_n_layer_nextn(model);
+    hp_ngl         = llama_model_n_layer(model);
+    if (mparams->load_mtp) {
+        hp_ngl    += llama_model_n_layer_nextn(model);
+    }
     hp_n_ctx_train = llama_model_n_ctx_train(model);
     hp_n_expert    = llama_model_n_expert(model);
 

conversion/__init__.py

@@ -70,6 +70,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
     "Exaone4ForCausalLM": "exaone",
     "ExaoneForCausalLM": "exaone",
     "ExaoneMoEForCausalLM": "exaone",
+    "ExaoneMoeForCausalLM": "exaone",
     "FalconForCausalLM": "falcon",
     "FalconH1ForCausalLM": "falcon_h1",
     "FalconMambaForCausalLM": "mamba",

conversion/deepseek.py

@@ -17,8 +17,11 @@ from .base import LazyTorchTensor, MmprojModel, ModelBase, TextModel, gguf, logg
 from .qwen import QwenModel
 
 
-@ModelBase.register("DeepseekOCRForCausalLM", "UnlimitedOCRForCausalLM")
+@ModelBase.register("DeepseekOCRForCausalLM")
 class DeepseekOCRVisionModel(MmprojModel):
+    # HF dynamic_preprocess() max_num, which differs per model
+    preproc_max_tiles = 9
+
     def __init__(self, *args, **kwargs):
         super().__init__(*args, **kwargs)
         self.clip_projector_type = gguf.VisionProjectorType.DEEPSEEKOCR
@@ -43,6 +46,9 @@ class DeepseekOCRVisionModel(MmprojModel):
         # @bluebread: there's no window_size in config but just add it here anyway
         self.gguf_writer.add_vision_window_size(self.hparams.get("window_size", 14))
 
+        self.gguf_writer.add_vision_preproc_min_tiles(2)
+        self.gguf_writer.add_vision_preproc_max_tiles(self.preproc_max_tiles)
+
         # SAM configuration
         sam_hparams = hparams['sam']
         self.gguf_writer.add_vision_sam_layers_count(sam_hparams['layers'])
@@ -93,8 +99,15 @@ class DeepseekOCRVisionModel(MmprojModel):
         return super().filter_tensors((name, gen))
 
 
+@ModelBase.register("UnlimitedOCRForCausalLM")
+class UnlimitedOCRVisionModel(DeepseekOCRVisionModel):
+    preproc_max_tiles = 32
+
+
 @ModelBase.register("DeepseekOCR2ForCausalLM")
 class DeepseekOCR2VisionModel(DeepseekOCRVisionModel):
+    preproc_max_tiles = 6
+
     def __init__(self, *args, **kwargs):
         super().__init__(*args, **kwargs)
         self.clip_projector_type = gguf.VisionProjectorType.DEEPSEEKOCR2
@@ -520,6 +533,13 @@ class DeepseekV4Model(TextModel):
         for key, value in raw_hparams.items():
             self.hparams.setdefault(key, value)
 
+        # workaround for special rope_parameters (main/compress) in transformers 5.x
+        if self.rope_parameters.get("full_attention", self.rope_parameters).get("rope_type") is None:
+            if (rope_scaling := raw_hparams.get("rope_scaling")) is not None:
+                if "rope_type" not in rope_scaling and (rope_type := rope_scaling.get("type")) is not None:
+                    rope_scaling["rope_type"] = rope_type
+                self.rope_parameters.update(**rope_scaling)
+
         self.block_count = self.hparams["num_hidden_layers"]
         if self.mtp_only:
             self.block_count += self.hparams.get("num_nextn_predict_layers", 0)

conversion/exaone.py

@@ -123,7 +123,9 @@ class Exaone4Model(TextModel):
                 yield (self.format_tensor_name(gguf.MODEL_TENSOR.ROPE_FREQS), torch.tensor(rope_factors, dtype=torch.float32))
 
 
-@ModelBase.register("ExaoneMoEForCausalLM")
+# note: transformers >= 5.1 renamed the class to "ExaoneMoeForCausalLM" (lowercase 'e'),
+#       so accept both spellings - LG AI have updated the configs of already-released models
+@ModelBase.register("ExaoneMoEForCausalLM", "ExaoneMoeForCausalLM")
 class ExaoneMoEModel(Exaone4Model):
     model_arch = gguf.MODEL_ARCH.EXAONE_MOE
 

docs/backend/SYCL.md

@@ -449,6 +449,8 @@ Or
 use 1 SYCL GPUs: [0] with Max compute units:512
 ```
 
+User can use the device management in [docs/multi-gpu.md](https://github.com/ggml-org/llama.cpp/blob/master/docs/multi-gpu.md), like parameter `--device SYCL0,SYCL1` to assign one or more devices.
+
 ## Windows
 
 ### Install GPU driver
@@ -763,6 +765,7 @@ Or
 use 1 SYCL GPUs: [0] with Max compute units:512
 ```
 
+User can use the device management in [docs/multi-gpu.md](https://github.com/ggml-org/llama.cpp/blob/master/docs/multi-gpu.md), like parameter `--device SYCL0,SYCL1` to assign one or more devices.
 
 ## Environment Variable
 
@@ -895,6 +898,45 @@ Pass these via `CXXFLAGS` or add a one-off `#define` to enable a flag on the spo
     set UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
   ```
 
+- When I set `SYCL_CACHE_PERSISTENT=1` in running time, I meet crash.
+
+  `SYCL_CACHE_PERSISTENT=1` is not recommended by llama.cpp SYCL backend.
+  When cache is enabled, SYCL runtime will try to cache and reuse JIT-compiled binaries.
+
+  We find some AI will tell user this cmd to speed up SYCL backend. It only speeds up the startup to skip the JIT process, instead of running speed.
+
+  It will bring negative impact when the SYCL binary file is changed frequently in your running environment. The new & old codes mix will lead to crash.
+
+  Compare to the benefit, it has brought more failed cases.
+  If you are not familiar with the SYCL compiler principle of JIT and AOT, please don't use it.
+
+  To restore, you need to remove the local cache: `~/.cache/libsycl_cache/` and execute `unset SYCL_CACHE_PERSISTENT` in running time.
+
+- How to use iGPU and dGPU in same time?
+
+  1. Detect the devices in your running time.
+  ```
+  source /opt/intel/oneapi/setvars.sh
+  ./build/bin/llama-server --list-devices
+
+  or
+  ./build/bin/llama-cli --list-devices
+  ./build/bin/llama-bench --list-devices
+  ./build/bin/llama-completion --list-devices
+
+  Available devices:
+    SYCL0: Intel(R) Arc(TM) A770 Graphics (15473 MiB, 15473 MiB free)
+    SYCL1: Intel(R) UHD Graphics 770 (59675 MiB, 44986 MiB free)
+  ```
+
+  The dGPU will be in the head of this list and iGPU will be the end.
+  If not all GPUs are listed, please check the env var: ONEAPI_DEVICE_SELECTOR and unset it.
+
+  2. Set the iGPU and dGPU
+
+  Set the iGPU and dGPU by `./build/bin/llama-server --device SYCL0,SYCL1,SYCLxxx`.
+
+
 ### **GitHub contribution**:
 Please add the `[SYCL]` prefix/tag in issues/PRs titles to help the SYCL contributors to check/address them without delay.
 

docs/ops.md

@@ -15,7 +15,7 @@ Legend:
 | Operation | BLAS | CANN | CPU | CUDA | ET | MTL | OpenCL | SYCL | Vulkan | WebGPU | ZenDNN | zDNN |
 |-----------|------|------|------|------|------|------|------|------|------|------|------|------|
 |                              ABS | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
-|                              ACC | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ |
+|                              ACC | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                              ADD | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                             ADD1 | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
 |                           ADD_ID | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -41,9 +41,9 @@ Legend:
 |                             DIAG | ❌ | ❌ | ✅ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                    DIAG_MASK_INF | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                              DIV | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
-|                     DSV4_HC_COMB | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
-|                     DSV4_HC_POST | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
-|                      DSV4_HC_PRE | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
+|                     DSV4_HC_COMB | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
+|                     DSV4_HC_POST | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
+|                      DSV4_HC_PRE | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
 |                              DUP | ❌ | ✅ | ✅ | 🟡 | ❌ | 🟡 | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                              ELU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                              EXP | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -59,7 +59,7 @@ Legend:
 |                             GELU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                         GELU_ERF | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                       GELU_QUICK | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
-|                         GET_ROWS | ❌ | 🟡 | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | ✅ | 🟡 | ❌ | ❌ |
+|                         GET_ROWS | ❌ | 🟡 | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | ❌ | ❌ |
 |                    GET_ROWS_BACK | ❌ | ❌ | 🟡 | 🟡 | ❌ | ❌ | ❌ | ❌ | 🟡 | ❌ | ❌ | ❌ |
 |                       GROUP_NORM | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                      HARDSIGMOID | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -68,7 +68,7 @@ Legend:
 |                        IM2COL_3D | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                          L2_NORM | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | 🟡 | ❌ | ❌ |
 |                       LEAKY_RELU | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
-|                LIGHTNING_INDEXER | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
+|                LIGHTNING_INDEXER | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
 |                              LOG | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                             MEAN | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                              MUL | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |

examples/model-conversion/scripts/causal/convert-model.sh

@@ -47,6 +47,7 @@ CMD_ARGS+=("../../convert_hf_to_gguf.py" "--verbose")
 CMD_ARGS+=("${MODEL_PATH}")
 CMD_ARGS+=("--outfile" "${CONVERTED_MODEL}")
 CMD_ARGS+=("--outtype" "${TYPE}")
+CMD_ARGS+=("--model-name" "${MODEL_NAME}")
 [[ -n "$METADATA_OVERRIDE" ]] && CMD_ARGS+=("--metadata" "${METADATA_OVERRIDE}")
 [[ -n "$MMPROJ" ]] && CMD_ARGS+=("${MMPROJ}")
 

examples/model-conversion/scripts/embedding/convert-model.sh

@@ -31,6 +31,7 @@ python ../../convert_hf_to_gguf.py --verbose \
     ${EMBEDDING_MODEL_PATH} \
     --outfile ${CONVERTED_MODEL} \
     --outtype ${TYPE} \
+    --model-name ${MODEL_NAME} \
     ${SENTENCE_TRANSFORMERS}
 
 echo ""

examples/sycl/start-svr.sh

@@ -12,6 +12,7 @@ This script processes files with specified options.
 
 Options:
   -h, --help    Display this help message and exit.
+  -d, --device  <value>    Set SYCL devices (default: SYCL0).
   -c, --context <value>    Set context length. Bigger need more memory.
   -p, --promote <value>    Prompt to start generation with.
   -m, --model   <value>    Full model file path.
@@ -41,10 +42,16 @@ MODEL_FILE=../models/Qwen3.5-4B-Q4_0.gguf
 NGL=99
 CONTEXT=4096
 GGML_SYCL_DEVICE=-1
+SYCL_DEVICES="SYCL0"
 SPLIT_MODE=layer
 LOG_VERBOSE=3
 while [[ $# -gt 0 ]]; do
     case "$1" in
+        -d|--device)
+            SYCL_DEVICES="$2"
+            shift
+            shift
+            ;;
         -c|--context)
             CONTEXT=$2
             # Shift twice to consume both the option flag and its value
@@ -95,8 +102,6 @@ while [[ $# -gt 0 ]]; do
     esac
 done
 
-
-
 source /opt/intel/oneapi/setvars.sh
 
 #export GGML_SYCL_DEBUG=1
@@ -107,17 +112,19 @@ source /opt/intel/oneapi/setvars.sh
 export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
 echo "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=${UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS}"
 
+echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
+
+
 if [ $GGML_SYCL_DEVICE -ne -1 ]; then
     echo "Use $GGML_SYCL_DEVICE as main GPU"
     #use signle GPU only
     GPUS_SETTING="-mg $GGML_SYCL_DEVICE -sm ${SPLIT_MODE}"
-    echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
 else
-    echo "Use all Intel GPUs, including iGPU & dGPU"
+    echo "Use Intel GPUs: ${SYCL_DEVICES}"
     GPUS_SETTING="-sm ${SPLIT_MODE}"
- fi
+fi
 
-echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE}  --mmap --host 0.0.0.0 --port 8000"
-ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap --host 0.0.0.0 --port 8000
+echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap --host 0.0.0.0 --port 8000"
+ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap --host 0.0.0.0 --port 8000
 
 

examples/sycl/test.sh

@@ -12,6 +12,7 @@ This script processes files with specified options.
 
 Options:
   -h, --help    Display this help message and exit.
+  -d, --device  <value>    Set SYCL devices (default: SYCL0).
   -c, --context <value>    Set context length. Bigger need more memory.
   -p, --promote <value>    Prompt to start generation with.
   -m, --model   <value>    Full model file path.
@@ -42,10 +43,16 @@ MODEL_FILE=../models/llama-2-7b.Q4_0.gguf
 NGL=99
 CONTEXT=4096
 GGML_SYCL_DEVICE=-1
+SYCL_DEVICES="SYCL0"
 SPLIT_MODE=layer
 LOG_VERBOSE=3
 while [[ $# -gt 0 ]]; do
     case "$1" in
+        -d|--device)
+            SYCL_DEVICES="$2"
+            shift
+            shift
+            ;;
         -c|--context)
             CONTEXT=$2
             # Shift twice to consume both the option flag and its value
@@ -115,16 +122,17 @@ source /opt/intel/oneapi/setvars.sh
 export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
 echo "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=${UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS}"
 
+echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
+
 if [ $GGML_SYCL_DEVICE -ne -1 ]; then
     echo "Use $GGML_SYCL_DEVICE as main GPU"
     #use signle GPU only
     GPUS_SETTING="-mg $GGML_SYCL_DEVICE -sm ${SPLIT_MODE}"
-    echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
 else
-    echo "Use all Intel GPUs, including iGPU & dGPU"
+    echo "Use Intel GPUs: ${SYCL_DEVICES}"
     GPUS_SETTING="-sm ${SPLIT_MODE}"
  fi
 
-echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE}  --mmap "
-ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap
+echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap "
+ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap
 

examples/sycl/win-start-svr.bat

@@ -13,6 +13,7 @@ set "MODEL_FILE=..\models\Qwen3.5-4B-Q4_0.gguf"
 set "NGL=99"
 set "CONTEXT=4096"
 set "GGML_SYCL_DEVICE=-1"
+set "SYCL_DEVICES=SYCL0"
 set "SPLIT_MODE=layer"
 set "LOG_VERBOSE=3"
 
@@ -36,6 +37,21 @@ if /I "%~1"=="--context" (
   goto parse_args
 )
 
+if /I "%~1"=="-d" (
+  if "%~2"=="" goto missing_value
+  set "SYCL_DEVICES=%~2"
+  shift
+  shift
+  goto parse_args
+)
+if /I "%~1"=="--device" (
+  if "%~2"=="" goto missing_value
+  set "SYCL_DEVICES=%~2"
+  shift
+  shift
+  goto parse_args
+)
+
 if /I "%~1"=="-m" (
   if "%~2"=="" goto missing_value
   set "MODEL_FILE=%~2"
@@ -130,6 +146,7 @@ echo This script processes files with specified options.
 echo.
 echo Options:
 echo   -h, --help    Display this help message and exit.
+echo   -d, --device ^<value^>    Set SYCL devices (default: SYCL0).
 echo   -c, --context ^<value^>    Set context length. Bigger need more memory.
 echo   -m, --model   ^<value^>    Full model file path.
 echo   -mg,--main-gpu ^<value^>   Set main GPU ID (0 - n) for single GPU mode.
@@ -160,19 +177,20 @@ REM Support malloc device memory more than 4GB.
 set "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1"
 echo UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=%UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS%
 
+echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
+
 if not "%GGML_SYCL_DEVICE%"=="-1" (
   echo Use %GGML_SYCL_DEVICE% as main GPU
   REM Use single GPU only.
   set "GPUS_SETTING=-mg %GGML_SYCL_DEVICE% -sm %SPLIT_MODE%"
-  echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
-) else (
-  echo Use all Intel GPUs, including iGPU ^& dGPU
+  ) else (
+  echo Use Intel GPUs: %SYCL_DEVICES%
   set "GPUS_SETTING=-sm %SPLIT_MODE%"
 )
 
-echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap --host 0.0.0.0 --port 8000
+echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device %SYCL_DEVICES% --mmap --host 0.0.0.0 --port 8000
 set "ZES_ENABLE_SYSMAN=1"
-%BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap --host 0.0.0.0 --port 8000
+%BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device "%SYCL_DEVICES%" --mmap --host 0.0.0.0 --port 8000
 
 endlocal
 

examples/sycl/win-test.bat

@@ -19,6 +19,7 @@ set "MODEL_FILE=..\models\llama-2-7b.Q4_0.gguf"
 set "NGL=99"
 set "CONTEXT=4096"
 set "GGML_SYCL_DEVICE=-1"
+set "SYCL_DEVICES=SYCL0"
 set "SPLIT_MODE=layer"
 set "LOG_VERBOSE=3"
 
@@ -42,6 +43,21 @@ if /I "%~1"=="--context" (
   goto parse_args
 )
 
+if /I "%~1"=="-d" (
+  if "%~2"=="" goto missing_value
+  set "SYCL_DEVICES=%~2"
+  shift
+  shift
+  goto parse_args
+)
+if /I "%~1"=="--device" (
+  if "%~2"=="" goto missing_value
+  set "SYCL_DEVICES=%~2"
+  shift
+  shift
+  goto parse_args
+)
+
 if /I "%~1"=="-p" (
   if "%~2"=="" goto missing_value
   set "INPUT_PROMPT=%~2"
@@ -151,6 +167,7 @@ echo This script processes files with specified options.
 echo.
 echo Options:
 echo   -h, --help    Display this help message and exit.
+echo   -d, --device ^<value^>    Set SYCL devices (default: SYCL0).
 echo   -c, --context ^<value^>    Set context length. Bigger need more memory.
 echo   -p, --promote ^<value^>    Prompt to start generation with.
 echo   -m, --model   ^<value^>    Full model file path.
@@ -182,19 +199,21 @@ REM Support malloc device memory more than 4GB.
 set "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1"
 echo UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=%UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS%
 
+echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
+
 if not "%GGML_SYCL_DEVICE%"=="-1" (
   echo Use %GGML_SYCL_DEVICE% as main GPU
   REM Use single GPU only.
   set "GPUS_SETTING=-mg %GGML_SYCL_DEVICE% -sm %SPLIT_MODE%"
-  echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
-) else (
-  echo Use all Intel GPUs, including iGPU ^& dGPU
+  )
+else (
+  echo Use Intel GPUs: %SYCL_DEVICES%
   set "GPUS_SETTING=-sm %SPLIT_MODE%"
 )
 
-echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m %MODEL_FILE% -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap
+echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m %MODEL_FILE% -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device %SYCL_DEVICES% --mmap
 set "ZES_ENABLE_SYSMAN=1"
-%BIN_FILE% -m "%MODEL_FILE%" -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap
+%BIN_FILE% -m "%MODEL_FILE%" -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device "%SYCL_DEVICES%" --mmap
 
 endlocal
 

ggml/include/ggml.h

@@ -2788,6 +2788,12 @@ extern "C" {
             struct ggml_cgraph * cgraph,
             struct ggml_tensor * tensor);
 
+    // add the tensor and its parents to the graph without marking them for compute
+    // the flag is set later, when the tensor is reached from a node that computes
+    GGML_API void ggml_build_forward_order(
+            struct ggml_cgraph * cgraph,
+            struct ggml_tensor * tensor);
+
     GGML_API void ggml_build_backward_expand(
         struct ggml_context *  ctx,        // context for gradient computation
         struct ggml_cgraph  *  cgraph,

ggml/src/ggml-cpu/arch/arm/cpu-feats.cpp

@@ -8,6 +8,22 @@
 #include <sys/sysctl.h>
 #endif
 
+#if !defined(HWCAP_FPHP)
+#define HWCAP_FPHP (1 << 9)
+#endif
+
+#if !defined(HWCAP_ASIMDHP)
+#define HWCAP_ASIMDHP (1 << 10)
+#endif
+
+#if !defined(HWCAP_ASIMDDP)
+#define HWCAP_ASIMDDP (1 << 20)
+#endif
+
+#if !defined(HWCAP_SVE)
+#define HWCAP_SVE (1 << 22)
+#endif
+
 #if !defined(HWCAP2_SVE2)
 #define HWCAP2_SVE2 (1 << 1)
 #endif
@@ -23,7 +39,7 @@
 struct aarch64_features {
     // has_neon not needed, aarch64 has NEON guaranteed
     bool has_dotprod     = false;
-    bool has_fp16_va     = false;
+    bool has_fp16        = false;
     bool has_sve         = false;
     bool has_sve2        = false;
     bool has_i8mm        = false;
@@ -36,7 +52,7 @@ struct aarch64_features {
         uint32_t hwcap2 = getauxval(AT_HWCAP2);
 
         has_dotprod = !!(hwcap & HWCAP_ASIMDDP);
-        has_fp16_va = !!(hwcap & HWCAP_FPHP);
+        has_fp16    = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP);
         has_sve     = !!(hwcap & HWCAP_SVE);
         has_sve2    = !!(hwcap2 & HWCAP2_SVE2);
         has_i8mm    = !!(hwcap2 & HWCAP2_I8MM);
@@ -75,7 +91,7 @@ static int ggml_backend_cpu_aarch64_score() {
     score += 1<<1;
 #endif
 #ifdef GGML_USE_FP16_VECTOR_ARITHMETIC
-    if (!af.has_fp16_va) { return 0; }
+    if (!af.has_fp16) { return 0; }
     score += 1<<2;
 #endif
 #ifdef GGML_USE_SVE

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -4033,7 +4033,11 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
                     continue;
                 }
 #ifndef NDEBUG
-                assert(node->buffer->buft == ggml_backend_cuda_buffer_type(cuda_ctx->device));
+                // On integrated GPUs (APUs, e.g. RDNA3.5) the scheduler may place a
+                // node's output on the host-visible buffer, which the compute path
+                // handles. Allow that here, mirroring the src-tensor check below.
+                assert(node->buffer->buft == ggml_backend_cuda_buffer_type(cuda_ctx->device) ||
+                       (integrated && ggml_backend_buft_is_cuda_host(node->buffer->buft)));
                 for (int j = 0; j < GGML_MAX_SRC; j++) {
                     if (node->src[j] != nullptr) {
                         assert(node->src[j]->buffer);
@@ -5205,6 +5209,7 @@ static bool ggml_backend_cuda_device_offload_op(ggml_backend_dev_t dev, const gg
 
 static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_t dev) {
 #ifdef GGML_CUDA_NO_PEER_COPY
+    GGML_UNUSED(dev);
     return nullptr;
 #else
     ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *)dev->context;

ggml/src/ggml-cuda/quantize.cu

@@ -8,7 +8,6 @@ struct __builtin_align__(32) float8 {
     float x; float y; float z; float w;
     float p; float q; float r; float s;
 };
-#endif
 
 #if CUDART_VERSION >= 12080
 static __device__ __forceinline__ float nvfp4_native_scale_error(
@@ -49,6 +48,7 @@ static __device__ __forceinline__ float nvfp4_native_scale_error(
     return err;
 }
 #endif // CUDART_VERSION >= 12080
+#endif // defined(BLACKWELL_MMA_AVAILABLE)
 
 __launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1)
 static __global__ void quantize_q8_1(

ggml/src/ggml-metal/ggml-metal.metal

@@ -11328,8 +11328,8 @@ kernel void kernel_lightning_indexer(
     const int i_kv_0   = tgpig.x*NK;            // first key of this threadgroup
     const int i_kv     = i_kv_0 + sgitg*NKPSG;  // first key of this simdgroup
 
-    threadgroup half4x4   sk4x4[NK*DK16];
-    threadgroup half    * sk = (threadgroup half *) sk4x4;
+    threadgroup half sk[NK * DK16 * 16];
+    threadgroup half4x4 * sk4x4 = (threadgroup half4x4 *) sk;
 
     for (short i = tiitg; i < NK*DK16; i += NTG) {
         const short ik  = i/DK16;

ggml/src/ggml-sycl/common.hpp

@@ -1022,9 +1022,20 @@ static T block_reduce(T val, T * shared_vals, int block_size_template) {
 }
 
 static __dpct_inline__ float ggml_sycl_ue4m3_to_fp32(uint8_t x) {
-    const uint32_t bits = x * (x != 0x7F && x != 0xFF);
-    const __nv_fp8_e4m3 xf = *reinterpret_cast<const __nv_fp8_e4m3 *>(&bits);
-    return static_cast<float>(xf) / 2;
+    // UE4M3 is unsigned: 4 exp bits (bias 7), 3 mantissa bits, no sign, no NaN.
+    // exp == 0xF is a valid exponent (256-448 range), not NaN.
+    if (x == 0 || x == 0x7F) {
+        return 0.0f;
+    }
+    const int exp = (x >> 3) & 0xF;
+    const int man = x & 0x7;
+    float raw;
+    if (exp == 0) {
+        raw = man * (1.0f / 8.0f) * sycl::pow(2.0f, -6.0f);
+    } else {
+        raw = (1.0f + man / 8.0f) * sycl::pow(2.0f, (float) exp - 7.0f);
+    }
+    return raw * 0.5f;
 }
 
 #endif // GGML_SYCL_COMMON_HPP

ggml/src/ggml-sycl/dsv4-hc.cpp

@@ -0,0 +1,280 @@
+#include "ggml-impl.h"
+#include "dsv4-hc.hpp"
+
+#include <cmath>
+
+static constexpr int DSV4_HC = 4;
+
+static void dsv4_hc_pre_f32_sycl(
+        const float * x, const float * weights, float * dst,
+        int64_t n_embd, int64_t hc, int64_t n_tokens,
+        int64_t sx0, int64_t sx1, int64_t sx2,
+        int64_t sw0, int64_t sw1,
+        int64_t sd0, int64_t sd1,
+        queue_ptr stream) {
+    const int64_t nr = n_embd * n_tokens;
+    const int64_t block_size = 256;
+    const int64_t num_blocks = (nr + block_size - 1) / block_size;
+
+    stream->parallel_for(
+        sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
+        [=](sycl::nd_item<1> item) {
+            const int64_t ir = item.get_global_id(0);
+            if (ir >= nr) {
+                return;
+            }
+
+            const int64_t i0 = ir % n_embd;
+            const int64_t it = ir / n_embd;
+
+            float sum = x[i0*sx0 + it*sx2] * weights[it*sw1];
+            for (int64_t ih = 1; ih < hc; ++ih) {
+                const float xv = x[i0*sx0 + ih*sx1 + it*sx2];
+                const float wv = weights[ih*sw0 + it*sw1];
+                sum += xv * wv;
+            }
+
+            dst[i0*sd0 + it*sd1] = sum;
+        });
+}
+
+static void dsv4_hc_comb_norm_cols(float * comb, float eps) {
+    for (int idst = 0; idst < DSV4_HC; ++idst) {
+        float sum = eps;
+        for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
+            sum += comb[idst + DSV4_HC*isrc];
+        }
+
+        const float inv_sum = 1.0f / sum;
+        for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
+            comb[idst + DSV4_HC*isrc] *= inv_sum;
+        }
+    }
+}
+
+static void dsv4_hc_comb_norm_rows(float * comb, float eps) {
+    for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
+        float sum = eps;
+        for (int idst = 0; idst < DSV4_HC; ++idst) {
+            sum += comb[idst + DSV4_HC*isrc];
+        }
+
+        const float inv_sum = 1.0f / sum;
+        for (int idst = 0; idst < DSV4_HC; ++idst) {
+            comb[idst + DSV4_HC*isrc] *= inv_sum;
+        }
+    }
+}
+
+static void dsv4_hc_comb_f32_sycl(
+        const float * mixes,
+        const float * scale,
+        const float * base,
+        float * dst,
+        int64_t n_tokens,
+        int64_t sm0,
+        int64_t sm1,
+        int64_t ss0,
+        int64_t sb0,
+        int64_t sd0,
+        int64_t sd1,
+        int64_t sd2,
+        float eps,
+        int32_t n_iter,
+        queue_ptr stream) {
+    constexpr int comb_offset = 2*DSV4_HC;
+
+    const int64_t block_size = 256;
+    const int64_t num_blocks = (n_tokens + block_size - 1) / block_size;
+
+    stream->parallel_for(
+        sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
+        [=](sycl::nd_item<1> item_ct1) {
+            const int64_t it = item_ct1.get_global_id(0);
+
+            if (it >= n_tokens) {
+                return;
+            }
+
+            const float scale_comb = scale[2*ss0];
+            float comb[DSV4_HC*DSV4_HC];
+
+            for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
+                float max = -INFINITY;
+                for (int idst = 0; idst < DSV4_HC; ++idst) {
+                    const int idx = idst + DSV4_HC*isrc;
+                    const float v = mixes[(comb_offset + idx)*sm0 + it*sm1] * scale_comb + base[(comb_offset + idx)*sb0];
+                    comb[idx] = v;
+                    max = fmaxf(max, v);
+                }
+
+                float sum = 0.0f;
+                for (int idst = 0; idst < DSV4_HC; ++idst) {
+                    const int idx = idst + DSV4_HC*isrc;
+                    const float v = expf(comb[idx] - max);
+                    comb[idx] = v;
+                    sum += v;
+                }
+
+                const float inv_sum = 1.0f / sum;
+                for (int idst = 0; idst < DSV4_HC; ++idst) {
+                    const int idx = idst + DSV4_HC*isrc;
+                    comb[idx] = comb[idx] * inv_sum + eps;
+                }
+            }
+
+            dsv4_hc_comb_norm_cols(comb, eps);
+            for (int32_t i = 1; i < n_iter; ++i) {
+                dsv4_hc_comb_norm_rows(comb, eps);
+                dsv4_hc_comb_norm_cols(comb, eps);
+            }
+
+            for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
+                for (int idst = 0; idst < DSV4_HC; ++idst) {
+                    const int idx = idst + DSV4_HC*isrc;
+                    dst[idst*sd0 + isrc*sd1 + it*sd2] = comb[idx];
+                }
+            }
+        });
+}
+
+static void dsv4_hc_post_f32_sycl(
+        const float * x, const float * residual, const float * post, const float * comb, float * dst,
+        int64_t n_embd, int64_t hc, int64_t n_tokens,
+        int64_t sx0, int64_t sx1,
+        int64_t sr0, int64_t sr1, int64_t sr2,
+        int64_t sp0, int64_t sp1,
+        int64_t sc0, int64_t sc1, int64_t sc2,
+        int64_t sd0, int64_t sd1, int64_t sd2,
+        queue_ptr stream) {
+    const int64_t nr = n_embd * hc * n_tokens;
+    const int64_t block_size = 256;
+    const int64_t num_blocks = (nr + block_size - 1) / block_size;
+
+    stream->parallel_for(
+        sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
+        [=](sycl::nd_item<1> item) {
+            const int64_t ir = item.get_global_id(0);
+            if (ir >= nr) {
+                return;
+            }
+
+            const int64_t i0   = ir % n_embd;
+            const int64_t idst = (ir / n_embd) % hc;
+            const int64_t it   = ir / (n_embd * hc);
+
+            float sum = x[i0*sx0 + it*sx1] * post[idst*sp0 + it*sp1];
+            for (int64_t isrc = 0; isrc < hc; ++isrc) {
+                sum += residual[i0*sr0 + isrc*sr1 + it*sr2] * comb[idst*sc0 + isrc*sc1 + it*sc2];
+            }
+
+            dst[i0*sd0 + idst*sd1 + it*sd2] = sum;
+        });
+}
+
+void ggml_sycl_op_dsv4_hc_pre(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * x       = dst->src[0];
+    const ggml_tensor * weights = dst->src[1];
+
+    GGML_ASSERT(x->type == GGML_TYPE_F32);
+    GGML_ASSERT(weights->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+    GGML_TENSOR_LOCALS(size_t, nbx, x,       nb);
+    GGML_TENSOR_LOCALS(size_t, nbw, weights, nb);
+    GGML_TENSOR_LOCALS(size_t, nbd, dst,     nb);
+
+    const int64_t n_embd   = x->ne[0];
+    const int64_t hc       = x->ne[1];
+    const int64_t n_tokens = x->ne[2];
+
+    queue_ptr stream = ctx.stream();
+
+    dsv4_hc_pre_f32_sycl(
+            (const float *) x->data, (const float *) weights->data, (float *) dst->data,
+            n_embd, hc, n_tokens,
+            nbx0 / sizeof(float), nbx1 / sizeof(float), nbx2 / sizeof(float),
+            nbw0 / sizeof(float), nbw1 / sizeof(float),
+            nbd0 / sizeof(float), nbd1 / sizeof(float),
+            stream);
+}
+
+void ggml_sycl_op_dsv4_hc_comb(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/3);
+
+    const ggml_tensor * mixes = dst->src[0];
+    const ggml_tensor * scale = dst->src[1];
+    const ggml_tensor * base  = dst->src[2];
+
+    GGML_ASSERT(mixes->type == GGML_TYPE_F32);
+    GGML_ASSERT(scale->type == GGML_TYPE_F32);
+    GGML_ASSERT(base->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+    constexpr int64_t hc_mix_dim = (2 + DSV4_HC)*DSV4_HC;
+
+    GGML_ASSERT(mixes->ne[0] == hc_mix_dim);
+    GGML_ASSERT(dst->ne[0] == DSV4_HC);
+    GGML_ASSERT(dst->ne[1] == DSV4_HC);
+    GGML_ASSERT(dst->ne[2] == mixes->ne[1]);
+    GGML_ASSERT(scale->ne[0] >= 3);
+    GGML_ASSERT(base->ne[0] == hc_mix_dim);
+
+    GGML_TENSOR_LOCALS(size_t, nbm, mixes, nb);
+    GGML_TENSOR_LOCALS(size_t, nbs, scale, nb);
+    GGML_TENSOR_LOCALS(size_t, nbb, base,  nb);
+    GGML_TENSOR_LOCALS(size_t, nbd, dst,   nb);
+
+    const int64_t n_tokens = mixes->ne[1];
+    const float eps = ggml_get_op_params_f32(dst, 0);
+    const int32_t n_iter = ggml_get_op_params_i32(dst, 1);
+
+    queue_ptr stream = ctx.stream();
+
+    dsv4_hc_comb_f32_sycl(
+            (const float *) mixes->data, (const float *) scale->data, (const float *) base->data, (float *) dst->data,
+            n_tokens,
+            nbm0 / sizeof(float), nbm1 / sizeof(float),
+            nbs0 / sizeof(float),
+            nbb0 / sizeof(float),
+            nbd0 / sizeof(float), nbd1 / sizeof(float), nbd2 / sizeof(float),
+            eps, n_iter, stream);
+}
+
+void ggml_sycl_op_dsv4_hc_post(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/4);
+    const ggml_tensor * x        = dst->src[0];
+    const ggml_tensor * residual = dst->src[1];
+    const ggml_tensor * post     = dst->src[2];
+    const ggml_tensor * comb     = dst->src[3];
+
+    GGML_ASSERT(x->type == GGML_TYPE_F32);
+    GGML_ASSERT(residual->type == GGML_TYPE_F32);
+    GGML_ASSERT(post->type == GGML_TYPE_F32);
+    GGML_ASSERT(comb->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+    GGML_TENSOR_LOCALS(size_t, nbx, x,        nb);
+    GGML_TENSOR_LOCALS(size_t, nbr, residual, nb);
+    GGML_TENSOR_LOCALS(size_t, nbp, post,     nb);
+    GGML_TENSOR_LOCALS(size_t, nbc, comb,     nb);
+    GGML_TENSOR_LOCALS(size_t, nbd, dst,      nb);
+
+    const int64_t n_embd   = x->ne[0];
+    const int64_t n_tokens = x->ne[1];
+    const int64_t hc       = residual->ne[1];
+
+    queue_ptr stream = ctx.stream();
+
+    dsv4_hc_post_f32_sycl(
+            (const float *) x->data, (const float *) residual->data,
+            (const float *) post->data, (const float *) comb->data, (float *) dst->data,
+            n_embd, hc, n_tokens,
+            nbx0 / sizeof(float), nbx1 / sizeof(float),
+            nbr0 / sizeof(float), nbr1 / sizeof(float), nbr2 / sizeof(float),
+            nbp0 / sizeof(float), nbp1 / sizeof(float),
+            nbc0 / sizeof(float), nbc1 / sizeof(float), nbc2 / sizeof(float),
+            nbd0 / sizeof(float), nbd1 / sizeof(float), nbd2 / sizeof(float),
+            stream);
+}

ggml/src/ggml-sycl/dsv4-hc.hpp

@@ -0,0 +1,10 @@
+#ifndef GGML_SYCL_DSV4_HC_HPP
+#define GGML_SYCL_DSV4_HC_HPP
+
+#include "common.hpp"
+
+void ggml_sycl_op_dsv4_hc_pre(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+void ggml_sycl_op_dsv4_hc_comb(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+void ggml_sycl_op_dsv4_hc_post(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+#endif // GGML_SYCL_DSV4_HC_HPP

ggml/src/ggml-sycl/element_wise.cpp

@@ -420,53 +420,31 @@ static void clamp(const T * x, T * dst, const float min, const float max, const
     }
 }
 
-template<typename T>
-static void gated_op_fused_geglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
+template<typename T, typename F>
+static void unary_gated_op_flat_kernel(const T * x, const T * g, T * dst, const uint64_t k, const sycl::nd_item<1> & item_ct1, F func) {
+    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
+        dst[i] = func(x[i]) * g[i];
+    }
+}
+
+template<typename T, typename F>
+static void unary_gated_op_generic_kernel(
+        const T * x,
+        const T * g,
+        T * dst,
+        const uint64_t k,
+        const sycl::uint3 n_fd,
+        const uint64_t o0,
+        const uint64_t o1,
+        const sycl::nd_item<1> & item_ct1,
+        F func) {
+
+    // rows of n columns at strides o0 and o1: two halves of one fused tensor, or two tensors
     SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
         const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
         const int64_t j0 = rc.x() * o0 + rc.y();
         const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
-        dst[i] = op_gelu(x[j0]) * g[j1];
-    }
-}
-
-template<typename T>
-static void gated_op_fused_reglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
-    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
-        const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
-        const int64_t j0 = rc.x() * o0 + rc.y();
-        const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
-        dst[i] = op_relu(x[j0]) * g[j1];
-    }
-}
-
-template<typename T>
-static void gated_op_fused_swiglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
-    SYCL_GLOBAL_ID_LOOP(k, item_ct1)  {
-        const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
-        const int64_t j0 = rc.x() * o0 + rc.y();
-        const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
-        dst[i] = op_silu(x[j0]) * g[j1];
-    }
-}
-
-template<typename T>
-static void gated_op_fused_geglu_erf(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
-    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
-        const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
-        const int64_t j0 = rc.x() * o0 + rc.y();
-        const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
-        dst[i] = op_gelu_erf(x[j0]) * g[j1];
-    }
-}
-
-template<typename T>
-static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
-    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
-        const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
-        const int64_t j0 = rc.x() * o0 + rc.y();
-        const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
-        dst[i] = op_gelu_quick(x[j0]) * g[j1];
+        dst[i] = func(x[j0]) * g[j1];
     }
 }
 
@@ -670,6 +648,35 @@ static inline void ggml_sycl_op_unary(
         });
 }
 
+template<typename F>
+static inline void ggml_sycl_op_unary_gated(
+        ggml_backend_sycl_context & ctx, ggml_tensor * dst, F func) {
+
+    dispatch_ggml_sycl_op_fused_glu(ctx, dst,
+        [func](const auto * x_ptr, const auto * g_ptr, auto * dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
+
+            const uint32_t num_blocks = (uint32_t) ceil_div(k, SYCL_GLU_BLOCK_SIZE);
+            const sycl::nd_range<1> launch_range(num_blocks * sycl::range<1>(SYCL_GLU_BLOCK_SIZE),
+                                                 sycl::range<1>(SYCL_GLU_BLOCK_SIZE));
+
+            // o0 == n and o1 == n make the index math the identity, so index flat
+            // note: not ggml_is_contiguous - a fused [gate|up] src0 is contiguous with o0 == 2n
+            if (o0 == n && o1 == n) {
+                main_stream->parallel_for(launch_range,
+                    [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                        unary_gated_op_flat_kernel(x_ptr, g_ptr, dst_ptr, k, item_ct1, func);
+                    });
+            } else {
+                // launch-invariant divisor, and only this path needs it
+                const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
+                main_stream->parallel_for(launch_range,
+                    [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                        unary_gated_op_generic_kernel(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1, func);
+                    });
+            }
+        });
+}
+
 
 static inline void ggml_sycl_op_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     GGML_ASSERT(dst->type == GGML_TYPE_F32);
@@ -967,42 +974,21 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
 }
 
 static inline void ggml_sycl_op_geglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
-        [](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
-            const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
-            const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
-            main_stream->parallel_for(
-                    sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
-                    sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                gated_op_fused_geglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
-            });
-        });
+    ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
+        return op_gelu(x);
+    });
 }
 
 static inline void ggml_sycl_op_reglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
-        [](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
-            const uint32_t num_blocks = ceil_div((uint32_t)k, SYCL_RELU_BLOCK_SIZE); // Using RELU block size for reglu
-            const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
-            main_stream->parallel_for(
-                    sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_RELU_BLOCK_SIZE)),
-                    sycl::range<1>(SYCL_RELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                gated_op_fused_reglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
-            });
-        });
+    ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
+        return op_relu(x);
+    });
 }
 
 static inline void ggml_sycl_op_swiglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
-        [](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
-            const uint32_t num_blocks = ceil_div((uint32_t)k, SYCL_SILU_BLOCK_SIZE); // Using SILU block size for swiglu
-            const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
-            main_stream->parallel_for(
-                    sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_SILU_BLOCK_SIZE)),
-                    sycl::range<1>(SYCL_SILU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                gated_op_fused_swiglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
-            });
-        });
+    ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
+        return op_silu(x);
+    });
 }
 
 __dpct_inline__ float ggml_sycl_op_swiglu_oai_single(float x, float g, float alpha = 1.702f, float limit = 7.0f) {
@@ -1097,29 +1083,15 @@ void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst)
 }
 
 static inline void ggml_sycl_op_geglu_erf(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
-        [](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
-            const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
-            const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
-            main_stream->parallel_for(
-                    sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
-                    sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                gated_op_fused_geglu_erf(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
-            });
-        });
+    ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
+        return op_gelu_erf(x);
+    });
 }
 
 static inline void ggml_sycl_op_geglu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
-        [](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
-            const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
-            const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
-            main_stream->parallel_for(
-                    sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
-                    sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                gated_op_fused_geglu_quick(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
-            });
-        });
+    ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
+        return op_gelu_quick(x);
+    });
 }
 
 

ggml/src/ggml-sycl/fattn-vec.hpp

@@ -73,6 +73,7 @@ static void flash_attn_ext_vec(const char* __restrict__ Q,
                         const int32_t nb31,
                         const int32_t nb32,
                         const int64_t nb33) {
+
 #ifdef SYCL_FLASH_ATTN
     // Skip unused kernel variants for faster compilation:
 
@@ -469,7 +470,6 @@ static void flash_attn_ext_vec(const char* __restrict__ Q,
         }
     }
 
-
     item_ct1.barrier(sycl::access::fence_space::local_space);
 
 #pragma unroll
@@ -591,22 +591,24 @@ void ggml_sycl_flash_attn_ext_vec_case_impl(ggml_backend_sycl_context & ctx, ggm
 
     const auto arch = ggml_sycl_info().devices[ctx.device].hw_info.arch;
     const int nthreads = ggml_sycl_fattn_vec_get_nthreads_device(arch);
-    // 256 threads would overflow the 64 KB work-group local memory at D == 512, so keep 128 there.
-    if (D <= 256 && nthreads == 256) {
-        constexpr int nthreads_hw = 256;
-        constexpr int nwarps = nthreads_hw / warp_size;
-        launch_fattn<D, cols_per_block, 1,
-                     flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
-                                        use_logit_softcap, warp_size, nthreads_hw>, warp_size>(
-            ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
-    } else {
-        constexpr int nthreads_hw = 128;
-        constexpr int nwarps = nthreads_hw / warp_size;
-        launch_fattn<D, cols_per_block, 1,
-                     flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
-                                        use_logit_softcap, warp_size, nthreads_hw>, warp_size>(
-            ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
+    if constexpr (D <= 256) {
+        if (nthreads == 256) {
+            constexpr int nthreads_hw = 256;
+            constexpr int nwarps = nthreads_hw / warp_size;
+            launch_fattn<D, cols_per_block, 1,
+                         flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
+                                            use_logit_softcap, warp_size, nthreads_hw>, warp_size>(
+                ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
+            return;
+        }
     }
+
+    constexpr int nthreads_hw = 128;
+    constexpr int nwarps = nthreads_hw / warp_size;
+    launch_fattn<D, cols_per_block, 1,
+                 flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
+                                    use_logit_softcap, warp_size, nthreads_hw>, warp_size>(
+        ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
 }
 
 template <int D, int type_K, int type_V>

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

@@ -62,6 +62,8 @@
 #include "ggml-sycl/repeat_back.hpp"
 #include "ggml-sycl/set_rows.hpp"
 #include "ggml-sycl/set.hpp"
+#include "ggml-sycl/dsv4-hc.hpp"
+#include "ggml-sycl/lightning-indexer.hpp"
 #include "ggml-sycl/conv2d.hpp"
 #include "ggml-sycl/conv2d-dw.hpp"
 #include "ggml-sycl/conv2d-transpose.hpp"
@@ -4942,6 +4944,18 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
         case GGML_OP_SET_ROWS:
             ggml_sycl_op_set_rows(ctx, dst);
             break;
+        case GGML_OP_DSV4_HC_PRE:
+            ggml_sycl_op_dsv4_hc_pre(ctx, dst);
+            break;
+        case GGML_OP_DSV4_HC_COMB:
+            ggml_sycl_op_dsv4_hc_comb(ctx, dst);
+            break;
+        case GGML_OP_DSV4_HC_POST:
+            ggml_sycl_op_dsv4_hc_post(ctx, dst);
+            break;
+        case GGML_OP_LIGHTNING_INDEXER:
+            ggml_sycl_op_lightning_indexer(ctx, dst);
+            break;
         case GGML_OP_DUP:
             ggml_sycl_dup(ctx, dst);
             break;
@@ -5795,17 +5809,33 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
 
         case GGML_OP_SET_ROWS:
             {
-
-                auto res = ((op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
-                         op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q5_0 ||
-                         op->type == GGML_TYPE_Q1_0 ||
-                         op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_IQ4_NL ||
-                         op->type == GGML_TYPE_MXFP4 || op->type == GGML_TYPE_NVFP4) &&
-                        op->src[0]->type == GGML_TYPE_F32 &&
-                        (op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32));
+                auto res = (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16 ||
+                            op->src[0]->type == GGML_TYPE_BF16) &&
+                           (op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32);
                 return res;
             }
             break;
+        case GGML_OP_DSV4_HC_PRE:
+            return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 &&
+                op->type == GGML_TYPE_F32;
+        case GGML_OP_DSV4_HC_COMB:
+            return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 &&
+                op->src[2]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
+        case GGML_OP_DSV4_HC_POST:
+            return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 &&
+                op->src[2]->type == GGML_TYPE_F32 && op->src[3]->type == GGML_TYPE_F32 &&
+                op->type == GGML_TYPE_F32;
+        case GGML_OP_LIGHTNING_INDEXER:
+            return op->src[0]->type == GGML_TYPE_F32 &&
+                (op->src[1]->type == GGML_TYPE_F16 || op->src[1]->type == GGML_TYPE_F32 ||
+                 op->src[1]->type == GGML_TYPE_BF16 || op->src[1]->type == GGML_TYPE_Q8_0 ||
+                 op->src[1]->type == GGML_TYPE_Q5_1 || op->src[1]->type == GGML_TYPE_Q5_0 ||
+                 op->src[1]->type == GGML_TYPE_Q4_1 || op->src[1]->type == GGML_TYPE_Q4_0 ||
+                 op->src[1]->type == GGML_TYPE_IQ4_NL) &&
+                op->src[2]->type == GGML_TYPE_F32 &&
+                op->src[3]->type == GGML_TYPE_F16 &&
+                op->type == GGML_TYPE_F32 &&
+                op->src[0]->ne[0] == WARP_SIZE * 8;
         case GGML_OP_CPY:
             {
                 ggml_type src0_type = op->src[0]->type;

ggml/src/ggml-sycl/lightning-indexer.cpp

@@ -0,0 +1,197 @@
+#include "lightning-indexer.hpp"
+#include "dequantize.hpp"
+
+static void lightning_indexer_f32_sycl(
+        const char * q, const char * k, const char * w, const char * m, float * dst,
+        int64_t n_embd, int64_t n_head, int64_t n_batch, int64_t n_stream, int64_t n_kv,
+        int64_t nem3,
+        int64_t nbq1, int64_t nbq2, int64_t nbq3,
+        int64_t nbk2, int64_t nbk3,
+        int64_t nbw1, int64_t nbw3,
+        int64_t nbm1, int64_t nbm3,
+        int64_t nb1, int64_t nb3,
+        ggml_type k_type,
+        queue_ptr stream) {
+
+    constexpr int64_t LANES = WARP_SIZE;
+    constexpr int64_t ELEMS_PER_LANE = 8;
+    constexpr int64_t ROWS_PER_BLOCK = 4;
+    constexpr int64_t BLOCK_SIZE = ROWS_PER_BLOCK * LANES;
+
+    const int64_t n_rows = n_batch * n_stream * n_kv;
+    const int64_t n_blocks = (n_rows + ROWS_PER_BLOCK - 1) / ROWS_PER_BLOCK;
+
+    stream->parallel_for(
+        sycl::nd_range<1>(
+            sycl::range<1>(n_blocks * BLOCK_SIZE),
+            sycl::range<1>(BLOCK_SIZE)),
+        [=](sycl::nd_item<1> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+            const int64_t ir   = item.get_global_id(0);
+            const int64_t lane = ir % LANES;
+            const int64_t row  = ir / LANES;
+            if (row >= n_rows) {
+                return;
+            }
+
+            const int64_t i_bs     = row / n_kv;
+            const int64_t i_kv     = row % n_kv;
+            const int64_t i_batch  = i_bs / n_stream;
+            const int64_t i_stream = i_bs % n_stream;
+
+            // load K row slice into registers (row is contiguous, nbk0 == type size)
+            const char * k_base = k + i_kv*nbk2 + i_stream*nbk3;
+            float k_local[ELEMS_PER_LANE];
+            if (k_type == GGML_TYPE_F16) {
+                const sycl::half * k_row = (const sycl::half *) k_base;
+#pragma unroll
+                for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                    k_local[j] = static_cast<float>(k_row[lane*ELEMS_PER_LANE + j]);
+                }
+            } else if (k_type == GGML_TYPE_F32) {
+                const float * k_row = (const float *) k_base;
+#pragma unroll
+                for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                    k_local[j] = k_row[lane*ELEMS_PER_LANE + j];
+                }
+            } else {
+                const int64_t lane_base = lane * ELEMS_PER_LANE;
+                switch (k_type) {
+                    case GGML_TYPE_BF16: {
+                        const sycl::ext::oneapi::bfloat16 * k_row = (const sycl::ext::oneapi::bfloat16 *) k_base;
+#pragma unroll
+                        for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                            k_local[j] = static_cast<float>(k_row[lane_base + j]);
+                        }
+                    } break;
+                    case GGML_TYPE_Q4_0:
+                    case GGML_TYPE_Q4_1:
+                    case GGML_TYPE_Q5_0:
+                    case GGML_TYPE_Q5_1: {
+#pragma unroll
+                        for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                            const int64_t idx = lane_base + j;
+                            const int64_t ib  = idx / QK4_0;
+                            const int iqs     = idx % (QK4_0/2);
+                            dfloat2 kv;
+                            if (k_type == GGML_TYPE_Q4_0) {
+                                dequantize_q4_0(k_base, ib, iqs, kv);
+                            } else if (k_type == GGML_TYPE_Q4_1) {
+                                dequantize_q4_1(k_base, ib, iqs, kv);
+                            } else if (k_type == GGML_TYPE_Q5_0) {
+                                dequantize_q5_0(k_base, ib, iqs, kv);
+                            } else {
+                                dequantize_q5_1(k_base, ib, iqs, kv);
+                            }
+                            k_local[j] = (idx % QK4_0) < (QK4_0/2) ? static_cast<float>(kv.x()) : static_cast<float>(kv.y());
+                        }
+                    } break;
+                    case GGML_TYPE_Q8_0: {
+#pragma unroll
+                        for (int64_t pair = 0; pair < ELEMS_PER_LANE / 2; ++pair) {
+                            const int64_t elem0 = lane_base + 2 * pair;
+                            dfloat2 kv;
+                            dequantize_q8_0(k_base, elem0 / QK8_0, elem0 % QK8_0, kv);
+                            k_local[2 * pair + 0] = static_cast<float>(kv.x());
+                            k_local[2 * pair + 1] = static_cast<float>(kv.y());
+                        }
+                    } break;
+                    case GGML_TYPE_IQ4_NL: {
+#pragma unroll
+                        for (int64_t pair = 0; pair < ELEMS_PER_LANE / 2; ++pair) {
+                            const int64_t elem0 = lane_base + 2 * pair;
+                            dfloat2 kv;
+                            dequantize_iq4_nl(k_base, elem0 / QK4_NL, elem0 % QK4_NL, kv);
+                            k_local[2 * pair + 0] = static_cast<float>(kv.x());
+                            k_local[2 * pair + 1] = static_cast<float>(kv.y());
+                        }
+                    } break;
+                    default:
+#pragma unroll
+                        for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                            k_local[j] = 0.0f;
+                        }
+                        break;
+                }
+            }
+
+            const char  * q_base = q + i_batch*nbq2 + i_stream*nbq3;
+            const float * w_base = (const float *) (w + i_batch*nbw1 + i_stream*nbw3);
+
+            float score = 0.0f;
+            for (int64_t h = 0; h < n_head; ++h) {
+                const float * q_row = (const float *) (q_base + h*nbq1);
+                float dot = 0.0f;
+#pragma unroll
+                for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
+                    const int64_t i = lane*ELEMS_PER_LANE + j;
+                    if (i < n_embd) {
+                        dot += q_row[i] * k_local[j];
+                    }
+                }
+                dot = sycl::reduce_over_group(item.get_sub_group(), dot, sycl::plus<float>());
+                if (lane == 0) {
+                    score += sycl::max(dot, 0.0f) * w_base[h];
+                }
+            }
+
+            if (lane == 0) {
+                const sycl::half * m_base = (const sycl::half *) (m + i_batch*nbm1 + (i_stream % nem3)*nbm3);
+                // flat-index store: storing through a strided base pointer
+                // hangs/misroutes writes on this stack when n_batch*n_stream > 1
+                const int64_t dst_idx = i_kv + i_batch*(nb1/sizeof(float)) + i_stream*(nb3/sizeof(float));
+                dst[dst_idx] = score + static_cast<float>(m_base[i_kv]);
+            }
+        });
+}
+
+void ggml_sycl_op_lightning_indexer(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/4);
+    const ggml_tensor * q = dst->src[0];
+    const ggml_tensor * k = dst->src[1];
+    const ggml_tensor * w = dst->src[2]; // weights
+    const ggml_tensor * m = dst->src[3]; // mask
+
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+    GGML_ASSERT(  q->type == GGML_TYPE_F32);
+    GGML_ASSERT(  w->type == GGML_TYPE_F32);
+    GGML_ASSERT(  m->type == GGML_TYPE_F16);
+    GGML_ASSERT(k->type == GGML_TYPE_F16 || k->type == GGML_TYPE_F32 || k->type == GGML_TYPE_BF16 ||
+                k->type == GGML_TYPE_Q8_0 || k->type == GGML_TYPE_Q5_1 || k->type == GGML_TYPE_Q5_0 ||
+                k->type == GGML_TYPE_Q4_1 || k->type == GGML_TYPE_Q4_0 || k->type == GGML_TYPE_IQ4_NL);
+
+    GGML_TENSOR_LOCALS(int64_t, neq, q, ne);
+    GGML_TENSOR_LOCALS(size_t,  nbq, q, nb);
+    GGML_TENSOR_LOCALS(int64_t, nek, k, ne);
+    GGML_TENSOR_LOCALS(size_t,  nbk, k, nb);
+    GGML_TENSOR_LOCALS(size_t,  nbw, w, nb);
+    GGML_TENSOR_LOCALS(int64_t, nem, m, ne);
+    GGML_TENSOR_LOCALS(size_t,  nbm, m, nb);
+    GGML_TENSOR_LOCALS(int64_t, ne, dst, ne);
+    GGML_TENSOR_LOCALS(size_t,  nb, dst, nb);
+
+    // input rows must be contiguous
+    GGML_ASSERT(nbq0 == ggml_type_size(q->type));
+    GGML_ASSERT(nbk0 == ggml_type_size(k->type));
+    GGML_ASSERT(nbm0 == ggml_type_size(m->type));
+    GGML_ASSERT(nb0  == ggml_type_size(dst->type));
+
+    const int64_t n_embd   = neq0;
+    const int64_t n_head   = neq1;
+    const int64_t n_batch  = neq2;
+    const int64_t n_stream = neq3;
+    const int64_t n_kv     = nek2;
+
+    GGML_ASSERT(n_embd == WARP_SIZE * 8);
+
+    lightning_indexer_f32_sycl(
+            (const char *) q->data, (const char *) k->data,
+            (const char *) w->data, (const char *) m->data, (float *) dst->data,
+            n_embd, n_head, n_batch, n_stream, n_kv, nem3,
+            nbq1, nbq2, nbq3,
+            nbk2, nbk3,
+            nbw1, nbw3,
+            nbm1, nbm3,
+            nb1, nb3,
+            k->type,
+            ctx.stream());
+}

ggml/src/ggml-sycl/lightning-indexer.hpp

@@ -0,0 +1,8 @@
+#ifndef GGML_SYCL_LIGHTNING_INDEXER_HPP
+#define GGML_SYCL_LIGHTNING_INDEXER_HPP
+
+#include "common.hpp"
+
+void ggml_sycl_op_lightning_indexer(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+#endif // GGML_SYCL_LIGHTNING_INDEXER_HPP

ggml/src/ggml-sycl/presets.hpp

@@ -20,8 +20,6 @@
 #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses
 
 #define SYCL_COL2IM_1D_BLOCK_SIZE 256
-#define SYCL_GELU_BLOCK_SIZE 256
-#define SYCL_SILU_BLOCK_SIZE 256
 #define SYCL_TANH_BLOCK_SIZE 256
 #define SYCL_RELU_BLOCK_SIZE 256
 #define SYCL_HARDSIGMOID_BLOCK_SIZE 256

ggml/src/ggml-sycl/set_rows.cpp

@@ -1,6 +1,10 @@
 #include "set_rows.hpp"
 #include "cpy.hpp"
 
+#include "ggml-quants.h"
+
+#include <vector>
+
 namespace utils {
 template<typename T>
 static constexpr bool is_arithmetic_v() {
@@ -20,7 +24,17 @@ convert (const char* src, char* dst) {
    *reinterpret_cast<TOut*>(dst) = dst_val;
 }
 
-template <typename TIdx, typename blockType, int qk, cpy_kernel_t cpyblck>
+#ifdef GGML_SYCL_HAS_BF16
+// sycl::vec::convert does not provide a half -> bfloat16 path, so route through float.
+template<>
+inline void convert<sycl::half, sycl::ext::oneapi::bfloat16>(const char* src, char* dst) {
+    const float tmp = sycl::vec<sycl::half, 1>(*reinterpret_cast<const sycl::half*>(src))
+                          .template convert<float, sycl::rounding_mode::automatic>()[0];
+    *reinterpret_cast<sycl::ext::oneapi::bfloat16*>(dst) = sycl::ext::oneapi::bfloat16(tmp);
+}
+#endif
+
+template <typename TIn, typename TIdx, typename blockType, int qk, cpy_kernel_t cpyblck>
 static void set_rows_sycl_q(const char * __restrict__ src0_d,
                             const TIdx * __restrict__ src1_d,
                             blockType * __restrict__ dst_d,
@@ -68,13 +82,22 @@ static void set_rows_sycl_q(const char * __restrict__ src0_d,
         const int64_t i11         = i02 % ne11;
         const int64_t i10         = i01;
         const size_t  src_offset  = calculate_offset<3>({ nb01, nb02, nb03 }, { i01, i02, i03 });
-        const char *  src_block   = src0_d + src_offset + i00 * sizeof(float);
+        const char *  src_block   = src0_d + src_offset + i00 * sizeof(TIn);
         const size_t  src1_offset = calculate_offset<3>({ nb10, nb11, nb12 }, { i10, i11, i12 });
         const int64_t dst_row     = src1_d[src1_offset / sizeof(TIdx)];
         const size_t  dst_offset =
             calculate_offset<3>({ nb1, nb2, nb3 }, { dst_row, i02, i03 }) + (i00 / qk) * sizeof(blockType);
         char * dst_block = reinterpret_cast<char *>(reinterpret_cast<char *>(dst_d) + dst_offset);
-        cpyblck(src_block, dst_block);
+        if constexpr (std::is_same_v<TIn, float>) {
+            cpyblck(src_block, dst_block);
+        } else {
+            float src_block_f32[qk];
+            const TIn * src_block_t = reinterpret_cast<const TIn *>(src_block);
+            for (int j = 0; j < qk; ++j) {
+                src_block_f32[j] = (float) src_block_t[j];
+            }
+            cpyblck(reinterpret_cast<const char *>(src_block_f32), dst_block);
+        }
     });
     GGML_UNUSED(ne10);
     GGML_UNUSED(ne13);
@@ -82,6 +105,139 @@ static void set_rows_sycl_q(const char * __restrict__ src0_d,
     GGML_UNUSED(nb13);
 }
 
+template<typename blockType>
+using quantize_row_qk_t = void (*)(const float *, blockType *, int64_t);
+
+using quantize_rows_f_t = size_t (*)(const float *, void *, int64_t, int64_t, const float *);
+
+template <typename TIn, typename TIdx, typename blockType, int qk, quantize_row_qk_t<blockType> quantize_row>
+static void set_rows_sycl_qk_host(
+        const ggml_tensor * src0,
+        const ggml_tensor * src1,
+        ggml_tensor * dst,
+        const int64_t ne00,
+        const int64_t ne01,
+        const int64_t ne02,
+        const int64_t ne03,
+        const int64_t ne11,
+        const int64_t ne12,
+        const size_t nb01,
+        const size_t nb02,
+        const size_t nb03,
+        const size_t nb10,
+        const size_t nb11,
+        const size_t nb12,
+        const size_t nb1,
+        const size_t nb2,
+        const size_t nb3,
+        queue_ptr stream) {
+    GGML_ASSERT(ne00 % qk == 0);
+
+    const size_t src0_bytes = ggml_nbytes(src0);
+    const size_t src1_bytes = ggml_nbytes(src1);
+
+    std::vector<char> src0_host(src0_bytes);
+    std::vector<char> src1_host(src1_bytes);
+
+    stream->memcpy(src0_host.data(), src0->data, src0_bytes);
+    stream->memcpy(src1_host.data(), src1->data, src1_bytes);
+    stream->wait();
+
+    std::vector<float> src_row_f32(ne00);
+    const int64_t nblocks = ne00 / qk;
+    std::vector<blockType> dst_row_q(nblocks);
+
+    for (int64_t i03 = 0; i03 < ne03; ++i03) {
+        for (int64_t i02 = 0; i02 < ne02; ++i02) {
+            for (int64_t i01 = 0; i01 < ne01; ++i01) {
+                const int64_t i12 = i03 % ne12;
+                const int64_t i11 = i02 % ne11;
+                const int64_t i10 = i01;
+
+                const size_t src1_offset = calculate_offset<3>({ nb10, nb11, nb12 }, { i10, i11, i12 });
+                const int64_t dst_row = *(const TIdx *) (src1_host.data() + src1_offset);
+
+                const size_t src0_row_offset = calculate_offset<3>({ nb01, nb02, nb03 }, { i01, i02, i03 });
+                const TIn * src_row = reinterpret_cast<const TIn *>(src0_host.data() + src0_row_offset);
+
+                for (int64_t i00 = 0; i00 < ne00; ++i00) {
+                    src_row_f32[i00] = (float) src_row[i00];
+                }
+
+                quantize_row(src_row_f32.data(), dst_row_q.data(), ne00);
+
+                const size_t dst_offset = calculate_offset<3>({ nb1, nb2, nb3 }, { dst_row, i02, i03 });
+                stream->memcpy((char *) dst->data + dst_offset, dst_row_q.data(), nblocks * sizeof(blockType));
+                stream->wait();
+            }
+        }
+    }
+}
+
+template <typename TIn, typename TIdx, typename blockType, int qk, quantize_rows_f_t quantize_rows>
+static void set_rows_sycl_iq_host(
+        const ggml_tensor * src0,
+        const ggml_tensor * src1,
+        ggml_tensor * dst,
+        const int64_t ne00,
+        const int64_t ne01,
+        const int64_t ne02,
+        const int64_t ne03,
+        const int64_t ne11,
+        const int64_t ne12,
+        const size_t nb01,
+        const size_t nb02,
+        const size_t nb03,
+        const size_t nb10,
+        const size_t nb11,
+        const size_t nb12,
+        const size_t nb1,
+        const size_t nb2,
+        const size_t nb3,
+        queue_ptr stream) {
+    GGML_ASSERT(ne00 % qk == 0);
+
+    const size_t src0_bytes = ggml_nbytes(src0);
+    const size_t src1_bytes = ggml_nbytes(src1);
+
+    std::vector<char> src0_host(src0_bytes);
+    std::vector<char> src1_host(src1_bytes);
+
+    stream->memcpy(src0_host.data(), src0->data, src0_bytes);
+    stream->memcpy(src1_host.data(), src1->data, src1_bytes);
+    stream->wait();
+
+    std::vector<float> src_row_f32(ne00);
+    const int64_t nblocks = ne00 / qk;
+    std::vector<blockType> dst_row_q(nblocks);
+
+    for (int64_t i03 = 0; i03 < ne03; ++i03) {
+        for (int64_t i02 = 0; i02 < ne02; ++i02) {
+            for (int64_t i01 = 0; i01 < ne01; ++i01) {
+                const int64_t i12 = i03 % ne12;
+                const int64_t i11 = i02 % ne11;
+                const int64_t i10 = i01;
+
+                const size_t src1_offset = calculate_offset<3>({ nb10, nb11, nb12 }, { i10, i11, i12 });
+                const int64_t dst_row = *(const TIdx *) (src1_host.data() + src1_offset);
+
+                const size_t src0_row_offset = calculate_offset<3>({ nb01, nb02, nb03 }, { i01, i02, i03 });
+                const TIn * src_row = reinterpret_cast<const TIn *>(src0_host.data() + src0_row_offset);
+
+                for (int64_t i00 = 0; i00 < ne00; ++i00) {
+                    src_row_f32[i00] = (float) src_row[i00];
+                }
+
+                quantize_rows(src_row_f32.data(), dst_row_q.data(), 1, ne00, nullptr);
+
+                const size_t dst_offset = calculate_offset<3>({ nb1, nb2, nb3 }, { dst_row, i02, i03 });
+                stream->memcpy((char *) dst->data + dst_offset, dst_row_q.data(), nblocks * sizeof(blockType));
+                stream->wait();
+            }
+        }
+    }
+}
+
 template<typename TIn, typename TIdx, typename TOut>
 static void k_set_rows(
         const char * __restrict__ src0, const TIdx * __restrict__ src1, char * __restrict__ dst,
@@ -200,31 +356,194 @@ static void set_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor * s
             break;
 #endif
         case GGML_TYPE_Q8_0:
-            set_rows_sycl_q<TIdx, block_q8_0, QK8_0, cpy_blck_f32_q8_0>(src0_d, src1_d, (block_q8_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q8_0, QK8_0, cpy_blck_f32_q8_0>(
+                src0_d, src1_d, (block_q8_0 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_Q1_0:
-            set_rows_sycl_q<TIdx, block_q1_0, QK1_0, cpy_blck_f32_q1_0>(src0_d, src1_d, (block_q1_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q1_0, QK1_0, cpy_blck_f32_q1_0>(
+                src0_d, src1_d, (block_q1_0 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            break;
+        case GGML_TYPE_Q2_0:
+            set_rows_sycl_q<TIn, TIdx, block_q2_0, QK2_0, cpy_blck_f32_q2_0>(
+                src0_d, src1_d, (block_q2_0 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_Q5_1:
-            set_rows_sycl_q<TIdx, block_q5_1, QK5_1, cpy_blck_f32_q5_1>(src0_d, src1_d, (block_q5_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q5_1, QK5_1, cpy_blck_f32_q5_1>(
+                src0_d, src1_d, (block_q5_1 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_Q5_0:
-            set_rows_sycl_q<TIdx, block_q5_0, QK5_0, cpy_blck_f32_q5_0>(src0_d, src1_d, (block_q5_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q5_0, QK5_0, cpy_blck_f32_q5_0>(
+                src0_d, src1_d, (block_q5_0 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_Q4_1:
-            set_rows_sycl_q<TIdx, block_q4_1, QK4_1, cpy_blck_f32_q4_1>(src0_d, src1_d, (block_q4_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q4_1, QK4_1, cpy_blck_f32_q4_1>(
+                src0_d, src1_d, (block_q4_1 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_Q4_0:
-            set_rows_sycl_q<TIdx, block_q4_0, QK4_0, cpy_blck_f32_q4_0>(src0_d, src1_d, (block_q4_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_q4_0, QK4_0, cpy_blck_f32_q4_0>(
+                src0_d, src1_d, (block_q4_0 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_IQ4_NL:
-            set_rows_sycl_q<TIdx, block_iq4_nl, QK4_NL, cpy_blck_f32_iq4_nl>(src0_d, src1_d, (block_iq4_nl *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_iq4_nl, QK4_NL, cpy_blck_f32_iq4_nl>(
+                src0_d, src1_d, (block_iq4_nl *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_MXFP4:
-            set_rows_sycl_q<TIdx, block_mxfp4, QK_MXFP4, cpy_blck_f32_mxfp4>(src0_d, src1_d, (block_mxfp4 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_mxfp4, QK_MXFP4, cpy_blck_f32_mxfp4>(
+                src0_d, src1_d, (block_mxfp4 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
             break;
         case GGML_TYPE_NVFP4:
-            set_rows_sycl_q<TIdx, block_nvfp4, QK_NVFP4, cpy_blck_f32_nvfp4>(src0_d, src1_d, (block_nvfp4 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            set_rows_sycl_q<TIn, TIdx, block_nvfp4, QK_NVFP4, cpy_blck_f32_nvfp4>(
+                src0_d, src1_d, (block_nvfp4 *) dst->data, ne00, ne01, ne02, ne03,
+                ne10, ne11, ne12, ne13, nb00, nb01,
+                nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
+            break;
+        case GGML_TYPE_Q2_K:
+            set_rows_sycl_qk_host<TIn, TIdx, block_q2_K, QK_K, quantize_row_q2_K_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_Q3_K:
+            set_rows_sycl_qk_host<TIn, TIdx, block_q3_K, QK_K, quantize_row_q3_K_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_Q4_K:
+            set_rows_sycl_qk_host<TIn, TIdx, block_q4_K, QK_K, quantize_row_q4_K_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_Q5_K:
+            set_rows_sycl_qk_host<TIn, TIdx, block_q5_K, QK_K, quantize_row_q5_K_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_Q6_K:
+            set_rows_sycl_qk_host<TIn, TIdx, block_q6_K, QK_K, quantize_row_q6_K_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ2_XXS:
+            set_rows_sycl_iq_host<TIn, TIdx, block_iq2_xxs, QK_K, quantize_iq2_xxs>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ2_XS:
+            set_rows_sycl_iq_host<TIn, TIdx, block_iq2_xs, QK_K, quantize_iq2_xs>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ2_S:
+            set_rows_sycl_iq_host<TIn, TIdx, block_iq2_s, QK_K, quantize_iq2_s>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ3_XXS:
+            set_rows_sycl_qk_host<TIn, TIdx, block_iq3_xxs, QK_K, quantize_row_iq3_xxs_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ3_S:
+            set_rows_sycl_qk_host<TIn, TIdx, block_iq3_s, QK_K, quantize_row_iq3_s_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ1_S:
+            set_rows_sycl_iq_host<TIn, TIdx, block_iq1_s, QK_K, quantize_iq1_s>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ1_M:
+            set_rows_sycl_iq_host<TIn, TIdx, block_iq1_m, QK_K, quantize_iq1_m>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
+            break;
+        case GGML_TYPE_IQ4_XS:
+            set_rows_sycl_qk_host<TIn, TIdx, block_iq4_xs, QK_K, quantize_row_iq4_xs_ref>(
+                src0, src1, dst,
+                ne00, ne01, ne02, ne03,
+                ne11, ne12,
+                nb01, nb02, nb03,
+                nb10, nb11, nb12,
+                nb1, nb2, nb3,
+                stream);
             break;
         default:
             GGML_ABORT("Unsupported tensor type!");
@@ -237,12 +556,21 @@ void ggml_sycl_op_set_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     const ggml_tensor * src0 = dst->src[0];
     const ggml_tensor * src1 = dst->src[1];
 
-    GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32 || dst->src[0]->type == GGML_TYPE_F16);
     GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I64 || dst->src[1]->type == GGML_TYPE_I32);
 
-    if (src1->type == GGML_TYPE_I64) {
-        set_rows_sycl<float, int64_t>(ctx, src0, src1, dst);
+    // dispatch on the index type (src1) and the source value type (src0)
+    if (src0->type == GGML_TYPE_F16) {
+        if (src1->type == GGML_TYPE_I64) {
+            set_rows_sycl<sycl::half, int64_t>(ctx, src0, src1, dst);
+        } else {
+            set_rows_sycl<sycl::half, int32_t>(ctx, src0, src1, dst);
+        }
     } else {
-        set_rows_sycl<float, int32_t>(ctx, src0, src1, dst);
+        if (src1->type == GGML_TYPE_I64) {
+            set_rows_sycl<float, int64_t>(ctx, src0, src1, dst);
+        } else {
+            set_rows_sycl<float, int32_t>(ctx, src0, src1, dst);
+        }
     }
 }

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

@@ -186,13 +186,22 @@ static bool is_pow2(uint32_t x) { return x > 1 && (x & (x-1)) == 0; }
 
 #define VK_DEVICE_DESCRIPTOR_POOL_SIZE 256
 
-#define VK_CHECK(err, msg)                                          \
+#define VK_CHECK(err, msg, dev)                                     \
     do {                                                            \
-        vk::Result err_ = (err);                                    \
+        vk::Result err_;                                            \
+        try {                                                       \
+            err_ = (err);                                           \
+        } catch (vk::DeviceLostError &) {                           \
+            ggml_vk_print_device_lost_info(dev);                    \
+            GGML_LOG_ERROR("ggml_vulkan: %s at %s:%d\n",            \
+                #err, __FILE__, __LINE__);                          \
+            throw;                                                  \
+        }                                                           \
         if (err_ != vk::Result::eSuccess) {                         \
-            fprintf(stderr, "ggml_vulkan: %s error %s at %s:%d\n",  \
+            GGML_LOG_ERROR("ggml_vulkan: %s error %s at %s:%d\n",   \
                 #err, to_string(err_).c_str(), __FILE__, __LINE__); \
-            exit(1);                                                \
+            throw vk::SystemError(vk::make_error_code(err_),        \
+                "ggml_vulkan: " msg);                               \
         }                                                           \
     } while (0)
 
@@ -302,9 +311,13 @@ struct vk_command_pool {
     }
 };
 
+static void ggml_vk_print_device_fault_info(const vk_device& device);
+static void ggml_vk_print_device_lost_info(const vk_device& device);
+
 // Prevent simultaneous submissions to the same queue.
 struct vk_queue_handle {
     vk::Queue queue;
+    vk_device_ref device;
     virtual void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) = 0;
     virtual void lock()   {}   // no-op by default (internally synchronized case)
     virtual void unlock() {}
@@ -315,7 +328,14 @@ struct vk_queue_handle_synchronized : vk_queue_handle {
     std::mutex mutex;
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
         std::lock_guard<std::mutex> guard(mutex);
-        queue.submit(submits, fence);
+        try {
+            queue.submit(submits, fence);
+        } catch (vk::DeviceLostError &) {
+            if (auto dev = device.lock()) {
+                ggml_vk_print_device_lost_info(dev);
+            }
+            throw;
+        }
     }
     void lock()   override { mutex.lock(); }
     void unlock() override { mutex.unlock(); }
@@ -324,7 +344,14 @@ struct vk_queue_handle_synchronized : vk_queue_handle {
 struct vk_queue_handle_unsynchronized : vk_queue_handle {
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
         // Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
-        queue.submit(submits, fence);
+        try {
+            queue.submit(submits, fence);
+        } catch (vk::DeviceLostError &) {
+            if (auto dev = device.lock()) {
+                ggml_vk_print_device_lost_info(dev);
+            }
+            throw;
+        }
     }
     // lock()/unlock() inherited no-ops
 };
@@ -835,6 +862,15 @@ struct vk_device_struct {
 
     bool pipeline_executable_properties_support {};
 
+    bool device_fault {};
+    PFN_vkGetDeviceFaultInfoEXT pfn_vkGetDeviceFaultInfoEXT {};
+
+    bool serialize_submissions {};
+
+    const ggml_cgraph * diag_cgraph {};
+    int diag_prev_start = -1;
+    int diag_prev_end = -1;
+
     size_t idx;
 
     bool mul_mat_l[GGML_TYPE_COUNT];
@@ -1118,6 +1154,57 @@ void vk_command_pool::destroy(vk::Device& device) {
     cmd_buffers.clear();
 }
 
+static void ggml_vk_print_device_fault_info(const vk_device& device) {
+    if (!device->device_fault || !device->pfn_vkGetDeviceFaultInfoEXT) {
+        return;
+    }
+
+    VkDeviceFaultCountsEXT fault_counts {};
+    fault_counts.sType = VK_STRUCTURE_TYPE_DEVICE_FAULT_COUNTS_EXT;
+    VkResult res = device->pfn_vkGetDeviceFaultInfoEXT(device->device, &fault_counts, nullptr);
+    if (res != VK_SUCCESS) {
+        GGML_LOG_ERROR("ggml_vulkan: vkGetDeviceFaultInfoEXT (counts) failed: %d\n", res);
+        return;
+    }
+
+    std::vector<VkDeviceFaultAddressInfoEXT> address_infos(fault_counts.addressInfoCount);
+    std::vector<VkDeviceFaultVendorInfoEXT> vendor_infos(fault_counts.vendorInfoCount);
+
+    VkDeviceFaultInfoEXT fault_info {};
+    fault_info.sType = VK_STRUCTURE_TYPE_DEVICE_FAULT_INFO_EXT;
+    fault_info.pAddressInfos = address_infos.data();
+    fault_info.pVendorInfos = vendor_infos.data();
+
+    res = device->pfn_vkGetDeviceFaultInfoEXT(device->device, &fault_counts, &fault_info);
+    if (res != VK_SUCCESS) {
+        GGML_LOG_ERROR("ggml_vulkan: vkGetDeviceFaultInfoEXT (info) failed: %d\n", res);
+        return;
+    }
+
+    if (fault_counts.addressInfoCount == 0 && fault_counts.vendorInfoCount == 0 && fault_info.description[0] == '\0') {
+        return;
+    }
+
+    if (fault_info.description[0] != '\0') {
+        GGML_LOG_ERROR("ggml_vulkan: device fault on %s: %s\n", device->name.c_str(), fault_info.description);
+    }
+
+    for (uint32_t i = 0; i < fault_counts.addressInfoCount; i++) {
+        const auto& info = address_infos[i];
+        GGML_LOG_CONT("  address fault %u: type=%d address=0x%llx precision=0x%llx\n",
+                i, (int)info.addressType,
+                (unsigned long long)info.reportedAddress,
+                (unsigned long long)info.addressPrecision);
+    }
+    for (uint32_t i = 0; i < fault_counts.vendorInfoCount; i++) {
+        const auto& info = vendor_infos[i];
+        GGML_LOG_CONT("  vendor fault %u: %s (code=0x%llx data=0x%llx)\n",
+                i, info.description,
+                (unsigned long long)info.vendorFaultCode,
+                (unsigned long long)info.vendorFaultData);
+    }
+}
+
 struct vk_buffer_struct {
     vk::Buffer buffer = VK_NULL_HANDLE;
     vk::DeviceMemory device_memory = VK_NULL_HANDLE;
@@ -2059,6 +2146,36 @@ static uint64_t ggml_vk_get_node_flops(const ggml_tensor * node) {
     return 0;
 }
 
+static void ggml_vk_print_node_list(const ggml_cgraph * cgraph, int start, int end) {
+    uint64_t total_flops = 0;
+    int n_ops = 0;
+    for (int j = start; j <= end && j < cgraph->n_nodes; j++) {
+        uint64_t flops = ggml_vk_get_node_flops(cgraph->nodes[j]);
+        total_flops += flops;
+        n_ops++;
+        if (flops > 0) {
+            GGML_LOG_CONT("  node %d: %s (%s) [%.2f GFLOP]\n",
+                    j, cgraph->nodes[j]->name, ggml_op_name(cgraph->nodes[j]->op),
+                    flops / 1e9);
+        } else {
+            GGML_LOG_CONT("  node %d: %s (%s)\n",
+                    j, cgraph->nodes[j]->name, ggml_op_name(cgraph->nodes[j]->op));
+        }
+    }
+    GGML_LOG_CONT("  total: %d ops, %.2f GFLOP\n", n_ops, total_flops / 1e9);
+}
+
+static void ggml_vk_print_device_lost_info(const vk_device& device) {
+    ggml_vk_print_device_fault_info(device);
+    if (device->serialize_submissions && device->diag_cgraph != nullptr && device->diag_prev_start >= 0) {
+        GGML_LOG_ERROR("ggml_vulkan: device lost on %s, likely caused by previous submission (nodes %d to %d):\n",
+                device->name.c_str(), device->diag_prev_start, device->diag_prev_end);
+        ggml_vk_print_node_list(device->diag_cgraph, device->diag_prev_start, device->diag_prev_end);
+    } else {
+        GGML_LOG_ERROR("ggml_vulkan: device lost on %s\n", device->name.c_str());
+    }
+}
+
 class vk_perf_logger {
   public:
     void print_timings(bool force = false) {
@@ -2471,17 +2588,27 @@ static void ggml_vk_wait_for_fence(ggml_backend_vk_context * ctx) {
     // Use waitForFences while most of the graph executes. Hopefully the CPU can sleep
     // during this wait.
     if (ctx->almost_ready_fence_pending) {
-        VK_CHECK(ctx->device->device.waitForFences({ ctx->almost_ready_fence }, true, UINT64_MAX), "almost_ready_fence");
+        VK_CHECK(ctx->device->device.waitForFences({ ctx->almost_ready_fence }, true, UINT64_MAX), "almost_ready_fence", ctx->device);
         ctx->device->device.resetFences({ ctx->almost_ready_fence });
         ctx->almost_ready_fence_pending = false;
     }
 
     // Spin (w/pause) waiting for the graph to finish executing.
     vk::Result result;
-    while ((result = ctx->device->device.getFenceStatus(ctx->fence)) != vk::Result::eSuccess) {
+    for (;;) {
+        try {
+            result = ctx->device->device.getFenceStatus(ctx->fence);
+        } catch (vk::DeviceLostError &) {
+            ggml_vk_print_device_lost_info(ctx->device);
+            GGML_LOG_ERROR("ggml_vulkan: getFenceStatus at %s:%d\n", __FILE__, __LINE__);
+            throw;
+        }
+        if (result == vk::Result::eSuccess) {
+            break;
+        }
         if (result != vk::Result::eNotReady) {
-            fprintf(stderr, "ggml_vulkan: error %s at %s:%d\n", to_string(result).c_str(), __FILE__, __LINE__);
-            exit(1);
+            GGML_LOG_ERROR("ggml_vulkan: error %s at %s:%d\n", to_string(result).c_str(), __FILE__, __LINE__);
+            throw vk::SystemError(vk::make_error_code(result), "ggml_vulkan: getFenceStatus");
         }
         for (uint32_t i = 0; i < 100; ++i) {
             YIELD();
@@ -3172,6 +3299,7 @@ static std::unique_ptr<vk_queue> ggml_vk_create_queue(vk_device& device, uint32_
     }
 
     h->queue = device->device.getQueue2(queue_info2);
+    h->device = device;
     q->handle = h;
 
     q->cmd_pool.init(device, q.get());
@@ -6117,6 +6245,8 @@ static vk_device ggml_vk_get_device(size_t idx) {
 #endif
             } else if (strcmp(VK_KHR_INTERNALLY_SYNCHRONIZED_QUEUES_EXTENSION_NAME, properties.extensionName) == 0) {
                 internally_sync_support = true;
+            } else if (strcmp("VK_EXT_device_fault", properties.extensionName) == 0) {
+                device->device_fault = true;
             }
         }
 
@@ -6471,8 +6601,18 @@ static vk_device ggml_vk_get_device(size_t idx) {
         }
 #endif
 
+        VkPhysicalDeviceFaultFeaturesEXT fault_features {};
+        fault_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FAULT_FEATURES_EXT;
+        if (device->device_fault) {
+            last_struct->pNext = (VkBaseOutStructure *)&fault_features;
+            last_struct = (VkBaseOutStructure *)&fault_features;
+            device_extensions.push_back("VK_EXT_device_fault");
+        }
+
         vkGetPhysicalDeviceFeatures2(device->physical_device, &device_features2);
 
+        device->device_fault = device->device_fault && fault_features.deviceFault;
+
         device->has_internally_synchronized_queues = internally_synchronized_queues_features.internallySynchronizedQueues;
 
         // Build queue create infos only after querying whether internally synchronized queues are enabled.
@@ -6771,6 +6911,11 @@ static vk_device ggml_vk_get_device(size_t idx) {
         device_create_info.setPNext(&device_features2);
         device->device = device->physical_device.createDevice(device_create_info);
 
+        if (device->device_fault) {
+            device->pfn_vkGetDeviceFaultInfoEXT = (PFN_vkGetDeviceFaultInfoEXT)
+                vkGetDeviceProcAddr(device->device, "vkGetDeviceFaultInfoEXT");
+        }
+
         // Queues
         device->compute_queue = ggml_vk_create_queue(device, compute_queue_family_index, 0, { vk::PipelineStageFlagBits::eComputeShader | vk::PipelineStageFlagBits::eTransfer }, false);
 
@@ -6893,6 +7038,8 @@ static vk_device ggml_vk_get_device(size_t idx) {
 
         device->idx = idx;
 
+        device->serialize_submissions = getenv("GGML_VK_SERIALIZE_SUBMISSIONS") != nullptr;
+
         device->disable_fusion = getenv("GGML_VK_DISABLE_FUSION") != nullptr;
 
         device->add_rms_fusion = !device->disable_fusion &&
@@ -8319,7 +8466,7 @@ static void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void *
         }
 
         ggml_vk_submit(subctx, dst->device->fence);
-        VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_buffer_write_2d waitForFences");
+        VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_buffer_write_2d waitForFences", dst->device);
         dst->device->device.resetFences({ dst->device->fence });
         ggml_vk_queue_command_pools_cleanup(dst->device);
     }
@@ -8431,7 +8578,7 @@ static void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, si
         ggml_vk_ctx_end(subctx);
         ggml_vk_submit(subctx, src->device->fence);
         VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX),
-                 "vk_buffer_read_2d uma waitForFences");
+                 "vk_buffer_read_2d uma waitForFences", src->device);
         src->device->device.resetFences({ src->device->fence });
         ggml_vk_queue_command_pools_cleanup(src->device);
 
@@ -8452,7 +8599,7 @@ static void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, si
         ggml_vk_ctx_end(subctx);
 
         ggml_vk_submit(subctx, src->device->fence);
-        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_read_2d waitForFences");
+        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_read_2d waitForFences", src->device);
         src->device->device.resetFences({ src->device->fence });
         ggml_vk_queue_command_pools_cleanup(src->device);
 
@@ -8487,7 +8634,7 @@ static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& sr
         ggml_vk_buffer_copy_async(subctx, dst, dst_offset, src, src_offset, size);
         ggml_vk_ctx_end(subctx);
         ggml_vk_submit(subctx, src->device->fence);
-        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_copy waitForFences");
+        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_copy waitForFences", src->device);
         src->device->device.resetFences({ src->device->fence });
         ggml_vk_queue_command_pools_cleanup(src->device);
     } else {
@@ -8531,7 +8678,7 @@ static void ggml_vk_buffer_memset(vk_buffer& dst, size_t offset, uint32_t c, siz
     ggml_vk_ctx_end(subctx);
 
     ggml_vk_submit(subctx, dst->device->fence);
-    VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_memset waitForFences");
+    VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_memset waitForFences", dst->device);
     dst->device->device.resetFences({ dst->device->fence });
     ggml_vk_queue_command_pools_cleanup(dst->device);
 }
@@ -14266,7 +14413,7 @@ static void ggml_vk_test_matmul(ggml_backend_vk_context * ctx, size_t m, size_t
 
     auto begin = std::chrono::high_resolution_clock::now();
     ggml_vk_submit(subctx, ctx->fence);
-    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_matmul waitForFences");
+    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_matmul waitForFences", ctx->device);
     ctx->device->device.resetFences({ ctx->fence });
     ggml_vk_queue_command_pools_cleanup(ctx->device);
 
@@ -14468,7 +14615,7 @@ static void ggml_vk_test_dequant(ggml_backend_vk_context * ctx, size_t ne, ggml_
     auto begin = std::chrono::high_resolution_clock::now();
 
     ggml_vk_submit(subctx, ctx->fence);
-    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_dequant waitForFences");
+    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_dequant waitForFences", ctx->device);
     ctx->device->device.resetFences({ ctx->fence });
     ggml_vk_queue_command_pools_cleanup(ctx->device);
 
@@ -14754,7 +14901,7 @@ static void ggml_vk_test_dequant_matmul(ggml_backend_vk_context * ctx, size_t m,
     auto begin = std::chrono::high_resolution_clock::now();
 
     ggml_vk_submit(subctx, ctx->fence);
-    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_dequant waitForFences");
+    VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "ggml_vk_test_dequant waitForFences", ctx->device);
     ctx->device->device.resetFences({ ctx->fence });
     ggml_vk_queue_command_pools_cleanup(ctx->device);
 
@@ -15553,7 +15700,9 @@ static void ggml_vk_compute_forward(ggml_backend_vk_context * ctx, ggml_cgraph *
             memset(mset.dst, mset.val, mset.n);
         }
 
-        if (almost_ready && !ctx->almost_ready_fence_pending) {
+        if (ctx->device->serialize_submissions) {
+            ggml_vk_submit(subctx, ctx->fence);
+        } else if (almost_ready && !ctx->almost_ready_fence_pending) {
             ggml_vk_submit(subctx, ctx->almost_ready_fence);
             ctx->almost_ready_fence_pending = true;
         } else {
@@ -16164,12 +16313,20 @@ static void ggml_vk_synchronize(ggml_backend_vk_context * ctx) {
             memcpy(cpy.dst, cpy.src, cpy.n);
         }
 
-        ggml_vk_submit(compute_ctx, {});
+        if (ctx->device->serialize_submissions) {
+            ggml_vk_submit(compute_ctx, ctx->fence);
+            VK_CHECK(ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX), "synchronize waitForFences", ctx->device);
+            ctx->device->device.resetFences({ ctx->fence });
+        } else {
+            ggml_vk_submit(compute_ctx, {});
+        }
         ctx->submit_pending = true;
     }
 
     if (ctx->submit_pending) {
-        if (ctx->device->async_use_transfer_queue && ctx->transfer_semaphore_last_submitted < ctx->transfer_semaphore.value) {
+        if (ctx->device->serialize_submissions) {
+            ctx->submit_pending = false;
+        } else if (ctx->device->async_use_transfer_queue && ctx->transfer_semaphore_last_submitted < ctx->transfer_semaphore.value) {
             vk::TimelineSemaphoreSubmitInfo tl_info{
                 1, &ctx->transfer_semaphore.value,
                 0, nullptr,
@@ -16186,7 +16343,9 @@ static void ggml_vk_synchronize(ggml_backend_vk_context * ctx) {
         } else {
             ctx->device->compute_queue->handle->submit({}, ctx->fence);
         }
-        ggml_vk_wait_for_fence(ctx);
+        if (!ctx->device->serialize_submissions) {
+            ggml_vk_wait_for_fence(ctx);
+        }
         ctx->submit_pending = false;
         if (cmd_buf) {
             cmd_buf->in_use = false;
@@ -16758,6 +16917,10 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
     VK_LOG_DEBUG("ggml_backend_vk_graph_compute(" << cgraph->n_nodes << " nodes)");
     ggml_backend_vk_context * ctx = (ggml_backend_vk_context *)backend->context;
 
+    ctx->device->diag_cgraph = nullptr;
+    ctx->device->diag_prev_start = -1;
+    ctx->device->diag_prev_end = -1;
+
     if (vk_instance.debug_utils_support) {
         vk::DebugUtilsLabelEXT dul = {};
         dul.pLabelName = "ggml_backend_vk_graph_compute";
@@ -16849,6 +17012,36 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
     }
     uint64_t flops_per_submit = std::min(flops_cap, ctx->last_total_flops / 40u);
 
+    auto const submit_after = [&](int start, int end) {
+        if (ctx->device->serialize_submissions) {
+            try {
+                auto res = ctx->device->device.waitForFences({ ctx->fence }, true, UINT64_MAX);
+                if (res != vk::Result::eSuccess) {
+                    GGML_LOG_ERROR("ggml_vulkan: waitForFences error during serialized submission\n");
+                    throw vk::SystemError(vk::make_error_code(res), "ggml_vulkan: waitForFences during serialized submission");
+                }
+            } catch (vk::DeviceLostError &) {
+                ggml_vk_print_device_fault_info(ctx->device);
+                GGML_LOG_ERROR("ggml_vulkan: device lost on %s waiting for submission (nodes %d to %d):\n",
+                        ctx->device->name.c_str(), start, end);
+                ggml_vk_print_node_list(cgraph, start, end);
+                throw;
+            }
+            ctx->device->device.resetFences({ ctx->fence });
+            ctx->submit_pending = false;
+            ctx->device->diag_cgraph = cgraph;
+            ctx->device->diag_prev_start = start;
+            ctx->device->diag_prev_end = end;
+        }
+        first_node_in_batch = true;
+        submitted_nodes = 0;
+        batch_flops = 0;
+        if (submit_count < 3) {
+            flops_per_submit *= 2;
+        }
+        submit_count++;
+    };
+
     for (int i = 0; i < cgraph->n_nodes; i++) {
         if (first_node_in_batch) {
             submit_node_idx = i;
@@ -16856,8 +17049,20 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
 
         {
             auto node_flops = ggml_vk_get_node_flops(cgraph->nodes[i]);
-            batch_flops += node_flops;
             total_flops += node_flops;
+
+            // Flush the current batch before recording a node that would push it over the flop threshold
+            if (flops_per_submit != 0 && submitted_nodes > 0 && batch_flops + node_flops >= flops_per_submit) {
+                vk_context flush_ctx = ggml_vk_get_compute_ctx(ctx);
+                ggml_vk_ctx_end(flush_ctx);
+                flush_ctx->exit_tensor_idx = -1;
+                ctx->compute_ctx.reset();
+                ggml_vk_compute_forward(ctx, cgraph, cgraph->nodes[submit_node_idx], submit_node_idx, false);
+                submit_after(submit_node_idx, i - 1);
+                submit_node_idx = i;
+            }
+
+            batch_flops += node_flops;
         }
 
         // op_srcs_fused_elementwise indicates whether an op's srcs all contribute to
@@ -17111,13 +17316,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
         }
 
         if (submit && enqueued) {
-            first_node_in_batch = true;
-            submitted_nodes = 0;
-            batch_flops = 0;
-            if (submit_count < 3) {
-                flops_per_submit *= 2;
-            }
-            submit_count++;
+            submit_after(submit_node_idx, i + (int)ctx->num_additional_fused_ops);
         }
         i += ctx->num_additional_fused_ops;
         ctx->num_additional_fused_ops = 0;
@@ -17133,13 +17332,13 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
         ggml_vk_ctx_end(compute_ctx);
 
         ggml_vk_submit(compute_ctx, ctx->device->fence);
-        VK_CHECK(ctx->device->device.waitForFences({ ctx->device->fence }, true, UINT64_MAX), "GGML_VULKAN_PERF waitForFences");
+        VK_CHECK(ctx->device->device.waitForFences({ ctx->device->fence }, true, UINT64_MAX), "GGML_VULKAN_PERF waitForFences", ctx->device);
         ctx->device->device.resetFences({ ctx->device->fence });
         ctx->compute_ctx.reset();
 
         // Get the results and pass them to the logger
         std::vector<uint64_t> timestamps(cgraph->n_nodes + 1);
-        VK_CHECK(ctx->device->device.getQueryPoolResults(ctx->query_pool, 0, ctx->query_idx, (cgraph->n_nodes + 1)*sizeof(uint64_t), timestamps.data(), sizeof(uint64_t), vk::QueryResultFlagBits::e64 | vk::QueryResultFlagBits::eWait), "get timestamp results");
+        VK_CHECK(ctx->device->device.getQueryPoolResults(ctx->query_pool, 0, ctx->query_idx, (cgraph->n_nodes + 1)*sizeof(uint64_t), timestamps.data(), sizeof(uint64_t), vk::QueryResultFlagBits::e64 | vk::QueryResultFlagBits::eWait), "get timestamp results", ctx->device);
         if (!vk_perf_logger_concurrent) {
             // Log each op separately
             for (int i = 1; i < ctx->query_idx; i++) {
@@ -18366,7 +18565,7 @@ static void ggml_backend_vk_device_event_synchronize(ggml_backend_dev_t dev, ggm
         vk::Semaphore sem = vkev->tl_semaphore.s;
         uint64_t val = vkev->tl_semaphore.value;
         vk::SemaphoreWaitInfo swi{vk::SemaphoreWaitFlags{}, sem, val};
-        VK_CHECK(device->device.waitSemaphores(swi, UINT64_MAX), "event_synchronize");
+        VK_CHECK(device->device.waitSemaphores(swi, UINT64_MAX), "event_synchronize", device);
 
         // Reset and move submitted events
         for (auto& event : vkev->events_submitted) {

ggml/src/ggml.c

@@ -7200,6 +7200,10 @@ void ggml_build_forward_expand(struct ggml_cgraph * cgraph, struct ggml_tensor *
     ggml_build_forward_impl(cgraph, tensor, true, true);
 }
 
+void ggml_build_forward_order(struct ggml_cgraph * cgraph, struct ggml_tensor * tensor) {
+    ggml_build_forward_impl(cgraph, tensor, true, false);
+}
+
 void ggml_build_backward_expand(
         struct ggml_context *  ctx,
         struct ggml_cgraph  *  cgraph,

src/llama-grammar.cpp

@@ -648,10 +648,12 @@ const char * llama_grammar_parser::parse_sequence(
             } else {
                 throw std::runtime_error(std::string("expecting ',' at ") + pos);
             }
-            bool has_max = max_times != UINT64_MAX;
-            if (min_times > MAX_REPETITION_THRESHOLD || (has_max && max_times > MAX_REPETITION_THRESHOLD)) {
+            if (min_times > MAX_REPETITION_THRESHOLD) {
                 throw std::runtime_error(std::string("number of repetitions exceeds sane defaults, please reduce the number of repetitions"));
             }
+            if (max_times != UINT64_MAX && max_times > MAX_REPETITION_THRESHOLD) {
+                max_times = UINT64_MAX;
+            }
             handle_repetitions(min_times, max_times);
         } else {
             break;

src/llama-model-loader.cpp

@@ -1249,7 +1249,13 @@ struct ggml_tensor * llama_model_loader::create_tensor(
         for (size_t dim = 0; dim < GGML_MAX_DIMS; dim++) {
             t_meta.ne[dim] = dim < ne.size() ? ne.begin()[dim] : 1;
             GGML_ASSERT(t_meta.ne[dim] >= 1);
-            t_meta.nb[dim] = dim == 0 ? ggml_type_size(type) : t_meta.ne[dim-1]*t_meta.nb[dim-1];
+            if (dim == 0) {
+                t_meta.nb[dim] = ggml_type_size(type);
+            } else if (dim == 1) {
+                t_meta.nb[dim] = ggml_row_size(type, t_meta.ne[dim-1]);
+            } else {
+                t_meta.nb[dim] = t_meta.nb[dim-1]*t_meta.ne[dim-1];
+            }
             GGML_ASSERT(t_meta.nb[dim] >= 1);
         }
         ggml_set_name(&t_meta, tn.str().c_str());
@@ -1272,10 +1278,18 @@ struct ggml_tensor * llama_model_loader::create_tensor(
     if (flags & TENSOR_ALLOW_RESHAPE) {
         for (size_t dim = 0; dim < GGML_MAX_DIMS; dim++) {
             t_meta.ne[dim] = dim < ne.size() ? ne.begin()[dim] : 1;
-            t_meta.nb[dim] = dim == 0 ? ggml_type_size(t_meta.type) : t_meta.ne[dim-1]*t_meta.nb[dim-1];
+            if (dim == 0) {
+                t_meta.nb[dim] = ggml_type_size(t_meta.type);
+            } else if (dim == 1) {
+                t_meta.nb[dim] = ggml_row_size(t_meta.type, t_meta.ne[dim-1]);
+            } else {
+                t_meta.nb[dim] = t_meta.ne[dim-1]*t_meta.nb[dim-1];
+            }
         }
     }
 
+    GGML_ASSERT(ggml_nbytes(&t_meta) == ggml_nbytes(cur));
+
     ggml_backend_buffer_type_t buft = buft_for_tensor(&t_meta);
     if (buft == nullptr) {
         return nullptr;

src/llama-model.cpp

@@ -2890,6 +2890,21 @@ void llama_model_base::create_tensor_qkv(llama_layer & layer, int bid,
         int64_t n_embd_, int64_t n_embd_q_, int64_t n_embd_k_, int64_t n_embd_v_,
         int flags) {
     const int64_t n_embd_qkv = n_embd_q_ + n_embd_k_ + n_embd_v_;
+
+    if (flags & TENSOR_SKIP) {
+        const int skip = TENSOR_NOT_REQUIRED | TENSOR_SKIP;
+
+        create_tensor(tn(LLM_TENSOR_ATTN_QKV, "weight", bid), {n_embd_, n_embd_qkv}, skip | TENSOR_SKIP_IF_VIRTUAL);
+        create_tensor(tn(LLM_TENSOR_ATTN_QKV, "bias",   bid), {n_embd_qkv},          skip | TENSOR_SKIP_IF_VIRTUAL);
+        create_tensor(tn(LLM_TENSOR_ATTN_Q,   "weight", bid), {n_embd_, n_embd_q_},  skip);
+        create_tensor(tn(LLM_TENSOR_ATTN_K,   "weight", bid), {n_embd_, n_embd_k_},  skip);
+        create_tensor(tn(LLM_TENSOR_ATTN_V,   "weight", bid), {n_embd_, n_embd_v_},  skip);
+        create_tensor(tn(LLM_TENSOR_ATTN_Q,   "bias",   bid), {n_embd_q_},           skip);
+        create_tensor(tn(LLM_TENSOR_ATTN_K,   "bias",   bid), {n_embd_k_},           skip);
+        create_tensor(tn(LLM_TENSOR_ATTN_V,   "bias",   bid), {n_embd_v_},           skip);
+        return;
+    }
+
     layer.wqkv = create_tensor(tn(LLM_TENSOR_ATTN_QKV, "weight", bid), {n_embd_, n_embd_qkv}, TENSOR_NOT_REQUIRED | TENSOR_SKIP_IF_VIRTUAL);
     if (layer.wqkv) {
         layer.wqkv_b = create_tensor(tn(LLM_TENSOR_ATTN_QKV, "bias", bid), {n_embd_qkv}, TENSOR_NOT_REQUIRED | TENSOR_SKIP_IF_VIRTUAL);

tests/test-backend-ops.cpp

@@ -9747,6 +9747,15 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
 static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
     std::vector<std::unique_ptr<test_case>> test_cases;
 
+    // SWIGLU at a 27B-class FFN width, fused [gate|up] vs split operands
+    // note: same bytes either way, so a backend that indexes them differently shows it here
+    for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
+        for (int64_t n_tokens : {512, 2048}) {
+            test_cases.emplace_back(new test_glu(GGML_GLU_OP_SWIGLU, type, { 2*17408, n_tokens, 1, 1 }, 0, false));
+            test_cases.emplace_back(new test_glu_split(GGML_GLU_OP_SWIGLU, type, { 17408, n_tokens, 1, 1 }, 0));
+        }
+    }
+
     // Conv2d: K=CRS=NPQ=4096 matmul performance
     uint32_t                        iwh_idx  = 0;
     uint32_t                        kwh_idx  = 1;

tests/test-backend-sampler.cpp

@@ -1668,9 +1668,18 @@ static std::vector<const backend_test_case *> collect_tests_to_run(const std::st
         }
     } else {
         for (const auto & test : BACKEND_TESTS) {
-            if (test.enabled_by_default) {
-                selected.push_back(&test);
+            if (!test.enabled_by_default) {
+                continue;
             }
+#ifdef GGML_USE_HIP
+            // TODO: remove this when https://github.com/ggml-org/llama.cpp/pull/26592 is merged
+            if (test.name == "penalties" || test.name == "set_sampler" ||
+                test.name == "mixed"     || test.name == "top_p") {
+                fprintf(stderr, "Skipping test '%s' on HIP backend (no backend TOP_K support)\n", test.name.c_str());
+                continue;
+            }
+#endif // GGML_USE_HIP
+            selected.push_back(&test);
         }
     }
 

tests/test-grammar-parser.cpp

@@ -153,6 +153,53 @@ int main()
         root ::= "a"{,10}"
     )""");
 
+    verify_failure(R"""(
+        root ::= "a"{5000}
+    )""");
+
+    verify_failure(R"""(
+        root ::= "a"{5000,}
+    )""");
+
+    verify_failure(R"""(
+        root ::= "a"{5000,6000}
+    )""");
+
+    verify_parsing(R"""(
+        root ::= "a"{0,5000}
+    )""", {
+        {"root", 0},
+        {"root_1", 1},
+    }, {
+        // root (index 0)
+        {LLAMA_GRETYPE_RULE_REF, /* root_1 */ 1},
+        {LLAMA_GRETYPE_END, 0},
+        // root_1 (index 1)
+        {LLAMA_GRETYPE_CHAR, 'a'},
+        {LLAMA_GRETYPE_RULE_REF, /* root_1 */ 1},
+        {LLAMA_GRETYPE_ALT, 0},
+        {LLAMA_GRETYPE_END, 0},
+    });
+
+    verify_parsing(R"""(
+        root ::= "a"{3,5000}
+    )""", {
+        {"root", 0},
+        {"root_1", 1},
+    }, {
+        // root (index 0)
+        {LLAMA_GRETYPE_CHAR, 'a'},
+        {LLAMA_GRETYPE_CHAR, 'a'},
+        {LLAMA_GRETYPE_CHAR, 'a'},
+        {LLAMA_GRETYPE_RULE_REF, /* root_1 */ 1},
+        {LLAMA_GRETYPE_END, 0},
+        // root_1 (index 1)
+        {LLAMA_GRETYPE_CHAR, 'a'},
+        {LLAMA_GRETYPE_RULE_REF, /* root_1 */ 1},
+        {LLAMA_GRETYPE_ALT, 0},
+        {LLAMA_GRETYPE_END, 0},
+    });
+
     verify_parsing(R"""(
         root  ::= "a"
     )""", {

tests/test-llama-archs.cpp

@@ -432,11 +432,19 @@ static bool arch_supported(const llm_arch arch) {
 
     // FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
 #ifdef GGML_USE_WEBGPU
-    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_MINIMAX_M3) {
+    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA) {
         return false;
     }
 #endif // GGML_USE_WEBGPU
 
+    // FIXME: jamba produces incorrect output (~0.55 NMSE vs CPU) on the HIP
+    // backend on RDNA3.5 (gfx1151); the SSM kernels need investigation.
+#ifdef GGML_USE_HIP
+    if (arch == LLM_ARCH_JAMBA) {
+        return false;
+    }
+#endif // GGML_USE_HIP
+
     return true;
 }
 

tests/test-model-resolution.cpp

@@ -195,7 +195,7 @@ static const std::vector<std::string> dspark_dflash = {
 
 struct plan_case {
     const char * name;
-    const std::vector<std::string> & files;
+    const std::vector<std::string> files;
     const char * hf_repo;
     const char * hf_file;
     bool sidecars;        // request mmproj + mtp + dflash + eagle3 + dspark

tests/test-mtmd-c-api.c

@@ -1,4 +1,6 @@
 #include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
 #include <assert.h>
 
 #include "mtmd.h"
@@ -62,6 +64,72 @@ int main(void) {
         }
     }
 
+    // test chunk save/load round-trip
+    for (size_t i = 0; i < n_chunks; i++) {
+        const mtmd_input_chunk * chunk = mtmd_input_chunks_get(chunks, i);
+        assert(chunk != NULL);
+        enum mtmd_input_chunk_type type = mtmd_input_chunk_get_type(chunk);
+
+        // query the required buffer size (out_buf == NULL)
+        size_t expected_len = 0;
+        int32_t rc = mtmd_input_chunk_save(chunk, NULL, 0, &expected_len);
+        printf("    Chunk %zu: save query rc = %d, expected_len = %zu\n", i, rc, expected_len);
+        assert(rc == 0);
+        assert(expected_len > 0);
+
+        // saving into a too-small buffer must fail, not crash
+        char tiny_buf[1];
+        rc = mtmd_input_chunk_save(chunk, tiny_buf, sizeof(tiny_buf), NULL);
+        printf("    Chunk %zu: save into too-small buffer rc = %d (expect non-zero)\n", i, rc);
+        assert(rc != 0);
+
+        // save into a properly-sized buffer
+        char * buf = (char *) malloc(expected_len);
+        assert(buf != NULL);
+        rc = mtmd_input_chunk_save(chunk, buf, expected_len, NULL);
+        assert(rc == 0);
+
+        // loading from a truncated buffer must fail gracefully, not crash
+        if (expected_len > 1) {
+            mtmd_input_chunk * bad = mtmd_input_chunk_load(buf, expected_len - 1);
+            printf("    Chunk %zu: load from truncated buffer = %p (expect NULL)\n", i, (void *) bad);
+            assert(bad == NULL);
+        }
+
+        // load it back
+        mtmd_input_chunk * loaded = mtmd_input_chunk_load(buf, expected_len);
+        assert(loaded != NULL);
+
+        // metadata must match the original chunk
+        assert(mtmd_input_chunk_get_type(loaded) == type);
+        assert(mtmd_input_chunk_get_n_tokens(loaded) == mtmd_input_chunk_get_n_tokens(chunk));
+        assert(mtmd_input_chunk_get_n_pos(loaded) == mtmd_input_chunk_get_n_pos(chunk));
+
+        if (type == MTMD_INPUT_CHUNK_TYPE_TEXT) {
+            size_t n_tok_orig, n_tok_loaded;
+            const llama_token * tok_orig   = mtmd_input_chunk_get_tokens_text(chunk, &n_tok_orig);
+            const llama_token * tok_loaded = mtmd_input_chunk_get_tokens_text(loaded, &n_tok_loaded);
+            printf("    Chunk %zu: loaded %zu text tokens (orig %zu), first token %d (orig %d)\n",
+                i, n_tok_loaded, n_tok_orig,
+                n_tok_loaded > 0 ? tok_loaded[0] : -1,
+                n_tok_orig   > 0 ? tok_orig[0]   : -1);
+            assert(n_tok_orig == n_tok_loaded);
+            for (size_t j = 0; j < n_tok_orig; j++) {
+                assert(tok_orig[j] == tok_loaded[j]);
+            }
+        } else if (type == MTMD_INPUT_CHUNK_TYPE_IMAGE || type == MTMD_INPUT_CHUNK_TYPE_AUDIO) {
+            const char * id_orig   = mtmd_input_chunk_get_id(chunk);
+            const char * id_loaded = mtmd_input_chunk_get_id(loaded);
+            printf("    Chunk %zu: loaded id '%s' (orig '%s')\n", i, id_loaded, id_orig);
+            assert(id_orig != NULL && id_loaded != NULL);
+            assert(strcmp(id_orig, id_loaded) == 0);
+        }
+
+        mtmd_input_chunk_free(loaded);
+        free(buf);
+    }
+    printf("Chunk save/load round-trip OK\n");
+
     // Free the chunks
     mtmd_input_chunks_free(chunks);
 

tools/mtmd/clip-impl.h

@@ -591,6 +591,8 @@ struct clip_image_u8 {
     }
 };
 
+struct mtmd_serialization; // forward declaration
+
 // For images, buf.size() == nx*ny*3
 //     Memory layout: RGBRGBRGB...
 // For seq, buf.size() == nx*ny*3*nt
@@ -671,6 +673,9 @@ struct clip_image_f32 {
         return buf.empty();
     }
 
+    void serialize(struct mtmd_serialization & ser) const;
+    void deserialize(struct mtmd_serialization & ser);
+
   private:
     std::vector<float> buf;
     int nx_ = 0;
@@ -752,6 +757,9 @@ struct clip_image_f32_batch {
         }
         return new_batch;
     }
+
+    void serialize(struct mtmd_serialization & ser) const;
+    void deserialize(struct mtmd_serialization & ser);
 };
 
 //

tools/mtmd/clip.cpp

@@ -708,9 +708,10 @@ ggml_tensor * clip_graph::build_attn(
         ggml_tensor * sinks) const {
     // these nodes are added to the graph together so that they are not reordered
     // by doing so, the number of splits in the graph is reduced
-    ggml_build_forward_expand(gf, q_cur);
-    ggml_build_forward_expand(gf, k_cur);
-    ggml_build_forward_expand(gf, v_cur);
+    // the order is fixed without the compute flag, so an unselected branch stays out of the compute set
+    ggml_build_forward_order(gf, q_cur);
+    ggml_build_forward_order(gf, k_cur);
+    ggml_build_forward_order(gf, v_cur);
 
     ggml_tensor * q = ggml_permute(ctx0, q_cur, 0, 2, 1, 3);
     //cb(q, "q", il);
@@ -1761,6 +1762,10 @@ struct clip_model_loader {
                             // qwen2 encoder is GQA, requires KEY_N_HEAD_KV
                             get_u32(string_format(KEY_N_HEAD_KV, "vision"), hparams.n_head_kv);
                         }
+                        // unlimited-ocr shares the v1 projector but tiles up to 32
+                        get_u32(KEY_PREPROC_MIN_TILES, hparams.preproc_min_tiles, false);
+                        get_u32(KEY_PREPROC_MAX_TILES, hparams.preproc_max_tiles, false);
+                        GGML_ASSERT(hparams.preproc_min_tiles <= hparams.preproc_max_tiles);
                      } break;
                 case PROJECTOR_TYPE_HUNYUANVL:
                     {
@@ -1909,6 +1914,9 @@ struct clip_model_loader {
                 if (hparams.image_max_pixels > 0) {
                     LOG_INF("%s: image_max_pixels:   %d%s\n", __func__, hparams.image_max_pixels, hparams.custom_image_max_tokens > 0 ? " (custom value)" : "");
                 }
+                if (hparams.preproc_max_tiles > 0) {
+                    LOG_INF("%s: preproc_tiles:      %d - %d\n", __func__, hparams.preproc_min_tiles, hparams.preproc_max_tiles);
+                }
             } else if (is_audio) {
                 LOG_INF("\n--- audio hparams ---\n");
                 LOG_INF("%s: n_mel_bins:         %d\n", __func__, hparams.n_mel_bins);

tools/mtmd/models/deepseekocr.cpp

@@ -253,6 +253,9 @@ ggml_cgraph * clip_graph_deepseekocr::build() {
 
     bool is_overview = img.add_viewsep;
     int n_tiles_per_row = 0;
+    // number of separate "row" images batched together in this graph call
+    // (captured now, before n_batch below gets repurposed as the SAM/ViT batch size)
+    const int n_rows_batch = n_batch;
 
     // note: we expect either a batch of rows or a batch of overviews, but not a mix of both
 
@@ -272,16 +275,18 @@ ggml_cgraph * clip_graph_deepseekocr::build() {
         GGML_ASSERT(img.ny() % img.nx() == 0);
         n_tiles_per_row = img.ny() / img.nx();
 
-        // input shape: [tile_size, tile_size * n_tiles_per_row, 3]
-        // we want to reshape it to [tile_size, tile_size, 3, n_tiles_per_row]
-        inp_raw = ggml_reshape_4d(ctx0, inp_raw, img.nx(), img.nx(), n_tiles_per_row, 3);
-        inp_raw = ggml_cont(ctx0, ggml_permute(ctx0, inp_raw, 0, 1, 3, 2));
+        // each entry is one "row" image of shape [tile_size, tile_size * n_tiles_per_row, 3];
+        // merge the tile axis into the batch axis, giving a combined SAM input of shape
+        // [tile_size, tile_size, 3, n_tiles_per_row * n_rows_batch] (tile fast, row slow)
+        inp_raw = ggml_reshape_4d(ctx0, inp_raw, img.nx() * img.nx(), n_tiles_per_row, 3, n_rows_batch);
+        inp_raw = ggml_cont(ctx0, ggml_permute(ctx0, inp_raw, 0, 2, 1, 3));
+        inp_raw = ggml_reshape_4d(ctx0, inp_raw, img.nx(), img.nx(), 3, n_tiles_per_row * n_rows_batch);
     }
 
     ggml_tensor * sam_out = build_sam(inp_raw);
 
     if (!is_overview) {
-        n_batch = n_tiles_per_row;
+        n_batch = n_tiles_per_row * n_rows_batch;
     }
 
     const int clip_n_patches = sam_out->ne[0] * sam_out->ne[1];
@@ -354,34 +359,36 @@ ggml_cgraph * clip_graph_deepseekocr::build() {
         const auto w     = h;
         const auto n_dim = cur->ne[0];
 
-        ggml_tensor * imgnl = ggml_repeat_4d(ctx0, model.image_newline, n_dim, 1, h, 1);
-        cur = ggml_reshape_3d(ctx0, cur, n_dim, w, h);
-        cur = ggml_reshape_2d(ctx0, ggml_concat(ctx0, cur, imgnl, 1), n_dim, (w + 1) * h);
-        cur = ggml_concat(ctx0, cur, model.view_seperator, 1);  // (n_dim, h*(w+1) + 1)
+        ggml_tensor * imgnl = ggml_repeat_4d(ctx0, model.image_newline, n_dim, 1, h, n_batch);
+        cur = ggml_reshape_4d(ctx0, cur, n_dim, w, h, n_batch);
+        cur = ggml_reshape_3d(ctx0, ggml_concat(ctx0, cur, imgnl, 1), n_dim, (w + 1) * h, n_batch);
+        ggml_tensor * vs = ggml_repeat_4d(ctx0, model.view_seperator, n_dim, 1, n_batch, 1);
+        cur = ggml_concat(ctx0, cur, vs, 1);  // (n_dim, h*(w+1) + 1, n_batch)
     } else {
         // tile row: interleave tiles within each row, add newline per row
-        const int grid_x      = static_cast<int>(std::sqrt(static_cast<float>(clip_n_patches)));
-        const int grid_y      = grid_x;
-        const auto n_dim      = cur->ne[0];
+        const int  grid_x = static_cast<int>(std::sqrt(static_cast<float>(clip_n_patches)));
+        const int  grid_y = grid_x;
+        const auto n_dim  = cur->ne[0];
 
-        // (n_dim, clip_n_patches, n_batch) -> (n_dim, grid_x, grid_y, n_batch)
-        cur = ggml_reshape_4d(ctx0, cur, n_dim, grid_x, grid_y, n_batch);
+        // merge n_dim into the grid_x axis, freeing the 4th axis for n_rows_batch
+        // (n_dim, clip_n_patches, n_tiles_per_row * n_rows_batch) -> (n_dim*grid_x, grid_y, n_tiles_per_row, n_rows_batch)
+        cur = ggml_reshape_4d(ctx0, cur, n_dim * grid_x, grid_y, n_tiles_per_row, n_rows_batch);
 
         // tiles: re-order from A.row0 A.row1 B.row0 B.row1 ...
         //        to A.row0 B.row0 A.row1 B.row1 ...
         //        then add nl: A.row0 B.row0 [nl] A.row1 B.row1 [nl] ...
-        // interleave tiles: (n_dim, grid_x, grid_y, n_batch) -> (n_dim, grid_x, n_batch, grid_y)
-        cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 0, 1, 3, 2));
+        // interleave tiles: -> (n_dim*grid_x, n_tiles_per_row, grid_y, n_rows_batch)
+        cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 0, 2, 1, 3));
 
-        // merge: (n_dim, grid_x, n_batch, grid_y) -> (n_dim, grid_x*n_batch, grid_y, 1)
-        cur = ggml_reshape_4d(ctx0, cur, n_dim, grid_x * n_batch, grid_y, 1);
+        // merge: -> (n_dim, grid_x*n_tiles_per_row, grid_y, n_rows_batch)
+        cur = ggml_reshape_4d(ctx0, cur, n_dim, grid_x * n_tiles_per_row, grid_y, n_rows_batch);
 
-        // append newline per row: (n_dim, grid_x*n_batch+1, grid_y, 1)
-        ggml_tensor * imgnl = ggml_repeat_4d(ctx0, model.image_newline, n_dim, 1, grid_y, 1);
+        // append newline per row: (n_dim, grid_x*n_tiles_per_row+1, grid_y, n_rows_batch)
+        ggml_tensor * imgnl = ggml_repeat_4d(ctx0, model.image_newline, n_dim, 1, grid_y, n_rows_batch);
         cur = ggml_concat(ctx0, cur, imgnl, 1);
 
-        // flatten: (n_dim, (grid_x*n_batch+1)*grid_y)
-        cur = ggml_reshape_2d(ctx0, cur, n_dim, (grid_x * n_batch + 1) * grid_y);
+        // flatten: (n_dim, (grid_x*n_tiles_per_row+1)*grid_y, n_rows_batch)
+        cur = ggml_reshape_3d(ctx0, cur, n_dim, (grid_x * n_tiles_per_row + 1) * grid_y, n_rows_batch);
     }
 
     cb(cur, "dsocr_output", -1);

tools/mtmd/models/deepseekocr2.cpp

@@ -14,8 +14,9 @@ ggml_cgraph * clip_graph_deepseekocr2::build() {
     {
         ggml_tensor * inp;
 
-        inp = ggml_reshape_2d(ctx0, sam_out, sam_out->ne[0] * sam_out->ne[1], sam_out->ne[2]); // H*W, C
-        inp = ggml_cont(ctx0, ggml_permute(ctx0, inp, 1, 0, 2, 3));
+        // H*W, C, B
+        inp = ggml_reshape_3d(ctx0, sam_out, sam_out->ne[0] * sam_out->ne[1], sam_out->ne[2], sam_out->ne[3]);
+        inp = ggml_cont(ctx0, ggml_permute(ctx0, inp, 1, 0, 2, 3)); // C, H*W, B
 
         auto num_image_tokens = inp->ne[1]; // H*W
         GGML_ASSERT(num_image_tokens == 144 || num_image_tokens == 256);
@@ -32,8 +33,10 @@ ggml_cgraph * clip_graph_deepseekocr2::build() {
             num_queries = 144;
         }
 
-        // (B, num_image_tokens + num_queries, C)
-        inp = ggml_concat(ctx0, inp, ggml_cast(ctx0, query_embed, inp->type), 1);
+        // repeat the query embedding per batch item, then append: (C, num_image_tokens + num_queries, B)
+        query_embed = ggml_cast(ctx0, query_embed, inp->type);
+        query_embed = ggml_repeat_4d(ctx0, query_embed, query_embed->ne[0], num_queries, inp->ne[2], 1);
+        inp = ggml_concat(ctx0, inp, query_embed, 1);
 
         auto seq_len = inp->ne[1];
 
@@ -57,7 +60,7 @@ ggml_cgraph * clip_graph_deepseekocr2::build() {
                                       /* learned_pos_embd */ nullptr, add_rope, vit_opts);
 
         cur = ggml_cont(ctx0,
-                        ggml_view_2d(ctx0, cur, cur->ne[0], num_queries, cur->nb[1],
+                        ggml_view_3d(ctx0, cur, cur->ne[0], num_queries, cur->ne[2], cur->nb[1], cur->nb[2],
                                      cur->nb[1] * (cur->ne[1] - num_queries))); // only take query tokens for output
 
         ggml_build_forward_expand(gf, cur);
@@ -71,7 +74,8 @@ ggml_cgraph * clip_graph_deepseekocr2::build() {
 
     // view_seperator only after the global view
     if (img.add_viewsep) {
-        cur = ggml_concat(ctx0, cur, model.view_seperator, 1); // (n_dim, 257)
+        ggml_tensor * vs = ggml_repeat_4d(ctx0, model.view_seperator, model.view_seperator->ne[0], 1, cur->ne[2], 1);
+        cur = ggml_concat(ctx0, cur, vs, 1); // (n_dim, 257, n_batch)
     }
 
     cb(cur, "dsocr2_output", -1);

tools/mtmd/models/models.h

@@ -138,12 +138,13 @@ struct clip_graph_deepseekocr : clip_graph {
     clip_graph_deepseekocr(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
     ggml_cgraph * build() override;
     ggml_tensor * build_sam(ggml_tensor * inp); // build the SAM model
-    // bool support_batch() const override { return true; } // TODO: support batch for DeepSeek-OCR v1
+    bool support_batch() const override { return true; }
 };
 
 struct clip_graph_deepseekocr2 : clip_graph_deepseekocr {
     clip_graph_deepseekocr2(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph_deepseekocr(ctx, img) {}
     ggml_cgraph * build() override; // reuses build_sam() from base
+    bool support_batch() const override { return true; }
 };
 
 struct clip_graph_conformer : clip_graph {

tools/mtmd/mtmd.cpp

@@ -22,8 +22,123 @@
 #include <cstdlib>
 #include <cstring>
 #include <climits>
+#include <type_traits>
 #include <vector>
 
+// remember to bump this if the serialization format changes
+#define MTMD_SERIALIZATION_VERSION 1
+
+struct mtmd_serialization {
+    // note: using 64-bit here for future-proofing
+    uint64_t version = MTMD_SERIALIZATION_VERSION;
+    std::vector<char> data;
+    size_t read_pos = 0; // cursor used when reading
+
+    // for writing
+    mtmd_serialization(uint64_t version) : version(version) {
+        write(version);
+    }
+
+    // for reading
+    mtmd_serialization(uint64_t version, const char * buf, size_t len) {
+        // copy buf to data
+        data.assign(buf, buf + len);
+        uint64_t ver_in = read<uint64_t>();
+        if (ver_in != version) {
+            throw std::runtime_error("version mismatch");
+        }
+        this->version = ver_in;
+    }
+
+    template <typename T>
+    void write(T value) {
+        static_assert(std::is_trivially_copyable<T>::value && !std::is_same<T, bool>::value,
+            "T must be trivially copyable and not bool");
+        const char * p = reinterpret_cast<const char *>(&value);
+        data.insert(data.end(), p, p + sizeof(T));
+    }
+
+    template <typename T>
+    T read() {
+        static_assert(std::is_trivially_copyable<T>::value && !std::is_same<T, bool>::value,
+            "T must be trivially copyable and not bool");
+        if (read_pos + sizeof(T) > data.size()) {
+            throw std::runtime_error("read OOB");
+        }
+        T value;
+        std::memcpy(&value, data.data() + read_pos, sizeof(T));
+        read_pos += sizeof(T);
+        return value;
+    }
+
+};
+
+template <>
+void mtmd_serialization::write<bool>(bool value) {
+    write<uint8_t>(value ? 1 : 0);
+}
+template <>
+bool mtmd_serialization::read<bool>() {
+    return read<uint8_t>() != 0;
+}
+
+template <>
+void mtmd_serialization::write<std::string>(std::string value) {
+    write<uint64_t>(value.size());
+    data.insert(data.end(), value.begin(), value.end());
+}
+template <>
+std::string mtmd_serialization::read<std::string>() {
+    uint64_t len = read<uint64_t>();
+    if (read_pos + len > data.size()) {
+        throw std::runtime_error("read_string OOB");
+    }
+    std::string str(data.data() + read_pos, len);
+    read_pos += len;
+    return str;
+}
+
+// only mtmd.cpp needs these, so they're implemented here rather than in clip-impl.h
+void clip_image_f32::serialize(mtmd_serialization & ser) const {
+    // remember to bump MTMD_SERIALIZATION_VERSION if this is changed
+    // note: buf is intentionally NOT serialized; the loaded clip_image_f32 will always be a placeholder
+    ser.write(add_viewsep);
+    ser.write(add_newline);
+    ser.write((int32_t)nx_);
+    ser.write((int32_t)ny_);
+}
+void clip_image_f32::deserialize(mtmd_serialization & ser) {
+    add_viewsep = ser.read<bool>();
+    add_newline = ser.read<bool>();
+    nx_ = ser.read<int32_t>();
+    ny_ = ser.read<int32_t>();
+    buf.clear(); // always a placeholder after loading
+}
+
+void clip_image_f32_batch::serialize(mtmd_serialization & ser) const {
+    // remember to bump MTMD_SERIALIZATION_VERSION if this is changed
+    ser.write(is_audio);
+    ser.write<uint64_t>(entries.size());
+    for (const auto & entry : entries) {
+        entry.serialize(ser);
+    }
+}
+void clip_image_f32_batch::deserialize(mtmd_serialization & ser) {
+    is_audio = ser.read<bool>();
+    uint64_t n = ser.read<uint64_t>();
+    constexpr size_t min_entry_bytes = sizeof(uint8_t) * 2 + sizeof(int32_t) * 2;
+    if (n > (ser.data.size() - ser.read_pos) / min_entry_bytes) {
+        throw std::runtime_error("entries count exceeds buffer size");
+    }
+    entries.clear();
+    entries.reserve(n);
+    for (uint64_t i = 0; i < n; i++) {
+        clip_image_f32 entry;
+        entry.deserialize(ser);
+        entries.push_back(std::move(entry));
+    }
+}
+
 // for still image data, layout is RGBRGBRGB...
 // length of data must be nx * ny * 3 bytes
 //
@@ -83,6 +198,7 @@ enum mtmd_pos_type {
     MTMD_POS_TYPE_NORMAL,    // number of positions equals to number of tokens
     MTMD_POS_TYPE_MROPE,     // qwen-vl mrope style, each image takes max(t,h,w) position indexes
     MTMD_POS_TYPE_HUNYUANVL, // HunyuanVL mrope + BOI/EOI/newline layout with XD-RoPE dim-3
+    MTMD_POS_TYPE_COUNT,     // for validation
 };
 
 struct mtmd_image_tokens {
@@ -136,6 +252,30 @@ struct mtmd_image_tokens {
             id
         };
     }
+
+    void serialize(mtmd_serialization & ser) const {
+        // remember to bump MTMD_SERIALIZATION_VERSION if this is changed
+        ser.write(nx);
+        ser.write(ny);
+        ser.write((uint32_t)pos);
+        ser.write(image_idx);
+        ser.write(n_temporal_merge);
+        ser.write(id);
+        batch_f32.serialize(ser);
+    }
+    void deserialize(mtmd_serialization & ser) {
+        nx = ser.read<uint32_t>();
+        ny = ser.read<uint32_t>();
+        uint32_t pos_raw = ser.read<uint32_t>();
+        if (pos_raw >= MTMD_POS_TYPE_COUNT) {
+            throw std::runtime_error("invalid pos type");
+        }
+        pos = (mtmd_pos_type)pos_raw;
+        image_idx = ser.read<uint32_t>();
+        n_temporal_merge = ser.read<uint32_t>();
+        id = ser.read<std::string>();
+        batch_f32.deserialize(ser);
+    }
 };
 using mtmd_image_tokens_ptr = std::unique_ptr<mtmd_image_tokens>;
 
@@ -161,6 +301,18 @@ struct mtmd_audio_tokens {
             id
         };
     }
+
+    void serialize(mtmd_serialization & ser) const {
+        // remember to bump MTMD_SERIALIZATION_VERSION if this is changed
+        ser.write(n_tokens);
+        ser.write(id);
+        batch_f32.serialize(ser);
+    }
+    void deserialize(mtmd_serialization & ser) {
+        n_tokens = ser.read<uint32_t>();
+        id = ser.read<std::string>();
+        batch_f32.deserialize(ser);
+    }
 };
 using mtmd_audio_tokens_ptr = std::unique_ptr<mtmd_audio_tokens>;
 
@@ -192,6 +344,66 @@ struct mtmd_input_chunk {
         }
         return false;
     }
+
+    void serialize(mtmd_serialization & ser) const {
+        // remember to bump MTMD_SERIALIZATION_VERSION if this is changed
+        ser.write((uint32_t)type);
+
+        ser.write<uint64_t>(tokens_text.size());
+        for (llama_token tok : tokens_text) {
+            ser.write((int32_t)tok);
+        }
+
+        ser.write(tokens_image != nullptr);
+        if (tokens_image) {
+            tokens_image->serialize(ser);
+        }
+
+        ser.write(tokens_audio != nullptr);
+        if (tokens_audio) {
+            tokens_audio->serialize(ser);
+        }
+    }
+    void deserialize(mtmd_serialization & ser) {
+        uint32_t type_raw = ser.read<uint32_t>();
+        if (type_raw >= MTMD_INPUT_CHUNK_TYPE_COUNT) {
+            throw std::runtime_error("invalid chunk type");
+        }
+        type = (mtmd_input_chunk_type)type_raw;
+
+        uint64_t n_tokens_text = ser.read<uint64_t>();
+        // reject before resize() so a tiny corrupted/malicious buffer can't force a huge allocation
+        if (n_tokens_text > (ser.data.size() - ser.read_pos) / sizeof(int32_t)) {
+            throw std::runtime_error("tokens_text length exceeds buffer size");
+        }
+        tokens_text.resize(n_tokens_text);
+        for (uint64_t i = 0; i < n_tokens_text; i++) {
+            tokens_text[i] = (llama_token)ser.read<int32_t>();
+        }
+
+        if (ser.read<bool>()) {
+            tokens_image = std::make_unique<mtmd_image_tokens>();
+            tokens_image->deserialize(ser);
+        } else {
+            tokens_image.reset();
+        }
+
+        if (ser.read<bool>()) {
+            tokens_audio = std::make_unique<mtmd_audio_tokens>();
+            tokens_audio->deserialize(ser);
+        } else {
+            tokens_audio.reset();
+        }
+
+        // catch buffers where the declared type doesn't match which payload is actually present,
+        // so a mismatched chunk can't slip through and null-deref/abort later in an accessor
+        if (type == MTMD_INPUT_CHUNK_TYPE_IMAGE && !tokens_image) {
+            throw std::runtime_error("type is IMAGE but tokens_image is missing");
+        }
+        if (type == MTMD_INPUT_CHUNK_TYPE_AUDIO && !tokens_audio) {
+            throw std::runtime_error("type is AUDIO but tokens_audio is missing");
+        }
+    }
 };
 
 struct mtmd_input_chunks {
@@ -2043,6 +2255,42 @@ void mtmd_input_chunk_free(mtmd_input_chunk * chunk) {
     }
 }
 
+int32_t mtmd_input_chunk_save(const mtmd_input_chunk * chunk, char * out_buf, size_t out_len, size_t * expected_out_len) {
+    try {
+        mtmd_serialization ser(MTMD_SERIALIZATION_VERSION);
+        chunk->serialize(ser);
+
+        if (expected_out_len) {
+            *expected_out_len = ser.data.size();
+        }
+        if (!out_buf) {
+            // caller is only querying the required size
+            return 0;
+        }
+        if (out_len < ser.data.size()) {
+            LOG_ERR("%s: out_buf is too small, need %zu bytes, got %zu\n", __func__, ser.data.size(), out_len);
+            return -1;
+        }
+        std::memcpy(out_buf, ser.data.data(), ser.data.size());
+        return 0;
+    } catch (const std::exception & e) {
+        LOG_ERR("%s: %s\n", __func__, e.what());
+        return -1;
+    }
+}
+
+mtmd_input_chunk * mtmd_input_chunk_load(const char * buf, size_t len) {
+    try {
+        mtmd_serialization ser(MTMD_SERIALIZATION_VERSION, buf, len);
+        mtmd::input_chunk_ptr chunk(new mtmd_input_chunk());
+        chunk->deserialize(ser);
+        return chunk.release();
+    } catch (const std::exception & e) {
+        LOG_ERR("%s: %s\n", __func__, e.what());
+        return nullptr;
+    }
+}
+
 // mtmd_image_tokens
 
 size_t mtmd_image_tokens_get_n_tokens(const mtmd_image_tokens * image_tokens) {

tools/mtmd/mtmd.h

@@ -55,6 +55,7 @@ enum mtmd_input_chunk_type {
     MTMD_INPUT_CHUNK_TYPE_TEXT,
     MTMD_INPUT_CHUNK_TYPE_IMAGE,
     MTMD_INPUT_CHUNK_TYPE_AUDIO,
+    MTMD_INPUT_CHUNK_TYPE_COUNT, // for validation
 };
 
 // opaque types
@@ -232,6 +233,15 @@ MTMD_API llama_pos                  mtmd_input_chunk_get_n_pos       (const mtmd
 MTMD_API mtmd_input_chunk * mtmd_input_chunk_copy(const mtmd_input_chunk * chunk);
 MTMD_API void               mtmd_input_chunk_free(mtmd_input_chunk * chunk);
 
+// save/load an input chunk to/from a buffer (useful for KV save/load)
+// important: only chunk's metadata will be saved, the actual image/audio data will not be saved
+// the loaded chunk will always be a placeholder, cannot be used for mtmd_encode() or mtmd_batch_encode()
+// out_buf can be nullptr (to query expected_out_len)
+// returns 0 on success, non-zero on failure
+MTMD_API int32_t            mtmd_input_chunk_save(const mtmd_input_chunk * chunk, char * out_buf, size_t out_len, size_t * expected_out_len);
+// returns nullptr on failure
+MTMD_API mtmd_input_chunk * mtmd_input_chunk_load(const char * buf, size_t len);
+
 
 // mtmd_image_tokens
 //

tools/server/server-models.cpp

@@ -2079,9 +2079,8 @@ server_http_proxy::server_http_proxy(
         return has_next; // false if EOF or pipe broken
     };
 
-    // wire up the HTTP client
-    // note: do NOT capture `this` pointer, as it may be destroyed before the thread ends
-    httplib::ResponseHandler response_handler = [pipe, cli](const httplib::Response & response) {
+    // build the header message forwarded to the reader thread, stripping internal proxy headers
+    auto make_header_msg = [](const httplib::Response & response) {
         msg_t msg;
         msg.status = response.status;
         for (const auto & [key, value] : response.headers) {
@@ -2095,7 +2094,17 @@ server_http_proxy::server_http_proxy(
             }
             msg.headers[key] = value;
         }
-        return pipe->write(std::move(msg)); // send headers first
+        return msg;
+    };
+
+    // true once response_handler has already forwarded the headers
+    auto headers_sent = std::make_shared<std::atomic<bool>>(false);
+
+    // wire up the HTTP client
+    // note: do NOT capture `this` pointer, as it may be destroyed before the thread ends
+    httplib::ResponseHandler response_handler = [pipe, headers_sent, make_header_msg](const httplib::Response & response) {
+        headers_sent->store(true);
+        return pipe->write(make_header_msg(response)); // send headers first
     };
     httplib::ContentReceiverWithProgress content_receiver = [pipe](const char * data, size_t data_length, size_t, size_t) {
         // send data chunks
@@ -2169,13 +2178,16 @@ server_http_proxy::server_http_proxy(
 
     // start the proxy thread
     SRV_DBG("start proxy thread %s %s\n", req.method.c_str(), req.path.c_str());
-    this->thread = std::thread([cli, pipe, req]() {
+    this->thread = std::thread([cli, pipe, req, headers_sent, make_header_msg]() {
         auto result = cli->send(std::move(req));
         if (result.error() != httplib::Error::Success) {
             auto err_str = httplib::to_string(result.error());
             SRV_ERR("http client error: %s\n", err_str.c_str());
             pipe->write({{}, 500, "", ""}); // header
             pipe->write({{}, 0, "proxy error: " + err_str, ""}); // body
+        } else if (!headers_sent->load()) {
+            // httplib skips response_handler for bodyless statuses like 204, send headers here instead
+            pipe->write(make_header_msg(*result));
         }
         pipe->close_write(); // signal EOF to reader
         SRV_DBG("%s", "client request thread ended\n");

tools/server/server-tools.cpp

@@ -30,20 +30,14 @@ namespace fs = std::filesystem;
 // internal helpers
 //
 
-#if defined(_WIN32)
-// A chunk can end in the middle of a multi-byte sequence, so the incomplete
-// tail is dropped before validating what precedes it.
-static bool is_utf8_text(const std::string & text) {
-    return is_valid_utf8(text.substr(0, validate_utf8(text)));
-}
-
-// A child process writes its output in the OEM code page, which is not UTF-8
-// on a western Windows install, so accented text reaches the JSON layer as
-// invalid bytes and is replaced there. Text that already decodes as UTF-8 is
-// returned untouched, so a child that emits UTF-8 is never decoded twice.
-// run() spawns without a console, so the console code page does not apply.
+// a child process writes in the OEM code page, so accented output would reach
+// the JSON layer as invalid bytes. run() spawns without a console, so the
+// console code page never applies
 static std::string console_output_to_utf8(const std::string & text) {
-    if (text.empty() || is_utf8_text(text)) {
+#if defined(_WIN32)
+    // a chunk can end mid sequence, so the incomplete tail is dropped first
+    if (text.empty() || is_valid_utf8(text.substr(0, validate_utf8(text)))) {
+        // never decode twice a child that already emits UTF-8
         return text;
     }
 
@@ -64,12 +58,10 @@ static std::string console_output_to_utf8(const std::string & text) {
     std::string utf8(utf8_len, '\0');
     WideCharToMultiByte(CP_UTF8, 0, wide.data(), wide_len, utf8.data(), utf8_len, nullptr, nullptr);
     return utf8;
-}
 #else
-static std::string console_output_to_utf8(const std::string & text) {
     return text;
-}
 #endif
+}
 
 json server_tool::to_json() const {
     return {
@@ -94,14 +86,30 @@ enum class list_kind {
     all,   // both
 };
 
+// a narrow path uses the active code page on Windows, so every crossing between
+// a std::string (always UTF-8 here) and fs::path is converted explicitly
+static fs::path path_from_utf8(const std::string & s) {
+    return fs::u8path(s);
+}
+
+// '/' separators on every platform: Windows accepts them, the web UI needs them
+static std::string path_to_utf8(const fs::path & p) {
+    const auto s = p.generic_u8string();
+    return std::string(s.begin(), s.end());
+}
+
 // home directory, read once at first use (getenv is not thread safe against setenv)
 static const std::string & home_dir() {
     static const std::string home = [] {
-        const char * h = getenv("HOME");
 #ifdef _WIN32
-        if (h == nullptr) h = getenv("USERPROFILE");
-#endif
+        // the narrow getenv would return the profile path in the active code page
+        const wchar_t * w = _wgetenv(L"HOME");
+        if (w == nullptr) w = _wgetenv(L"USERPROFILE");
+        return w ? path_to_utf8(fs::path(w)) : std::string();
+#else
+        const char * h = getenv("HOME");
         return h ? std::string(h) : std::string();
+#endif
     }();
     return home;
 }
@@ -140,11 +148,14 @@ public:
         std::string rel; // '/'-separated, relative to `base`
         bool is_dir = false;
     };
-    // entries relative to `base`; sets `err` if `base` isn't a directory
+    struct list_result {
+        std::vector<list_entry> entries;
+        std::string err;        // set when `base` is not a directory
+        bool truncated = false; // set when the walk could not see everything
+    };
+    // entries relative to `base`, which must already be resolved (absolute)
     // max_depth == 0 means unlimited, 1 means direct children of `base` only
-    // `base` must already be resolved (absolute); `caller_path` is the path the
-    // caller passed, used only for error messages
-    virtual std::vector<list_entry> list_entries(const std::string & base, const std::string & caller_path, int max_depth, list_kind kind, std::string & err, bool & truncated) const = 0;
+    virtual list_result list_entries(const std::string & base, int max_depth, list_kind kind) const = 0;
     // on_chunk, if set, is called with each chunk of output as it is read (before truncation cuts in);
     // returning false terminates the process early (e.g. the client disconnected)
     virtual exec_result run(
@@ -162,37 +173,47 @@ public:
     // expands a leading `~`, then resolves `path` against `cwd` (or the server
     // working directory when `cwd` is unset); the result is always absolute
     std::string resolve(const std::string & path) const override {
-        std::string p = expand_home(path);
-        if (fs::path(p).is_absolute()) {
-            return p;
+        const std::string p = expand_home(path);
+
+        fs::path full = path_from_utf8(p);
+        if (!full.is_absolute()) {
+            if (cwd.empty()) {
+                std::error_code ec;
+                const fs::path cur = fs::current_path(ec);
+                if (ec) return p;
+                full = cur / full;
+            } else {
+                full = path_from_utf8(cwd) / full;
+            }
         }
-        if (cwd.empty()) {
-            std::error_code ec;
-            fs::path cur = fs::current_path(ec);
-            if (ec) return p;
-            return (cur / p).string();
+
+        // drop "." and ".." so they never reach git or the client
+        full = full.lexically_normal();
+        // a trailing ".." normalizes to a path that ends with a separator
+        if (!full.has_filename() && full != full.root_path()) {
+            full = full.parent_path();
         }
-        return (fs::path(cwd) / p).string();
+        return path_to_utf8(full);
     }
 
     bool is_directory(const std::string & path) const override {
         std::error_code ec;
-        return fs::is_directory(resolve(path), ec) && !ec;
+        return fs::is_directory(path_from_utf8(resolve(path)), ec) && !ec;
     }
 
     bool is_regular_file(const std::string & path) const override {
         std::error_code ec;
-        return fs::is_regular_file(resolve(path), ec) && !ec;
+        return fs::is_regular_file(path_from_utf8(resolve(path)), ec) && !ec;
     }
 
     bool file_size(const std::string & path, uintmax_t & out_size) const override {
         std::error_code ec;
-        out_size = fs::file_size(resolve(path), ec);
+        out_size = fs::file_size(path_from_utf8(resolve(path)), ec);
         return !ec;
     }
 
     bool read_file(const std::string & path, std::string & out) const override {
-        std::ifstream f(resolve(path), std::ios::binary);
+        std::ifstream f(path_from_utf8(resolve(path)), std::ios::binary);
         if (!f) return false;
         std::ostringstream ss;
         ss << f.rdbuf();
@@ -202,7 +223,7 @@ public:
 
     bool write_file(const std::string & path, const std::string & content) const override {
         std::error_code ec;
-        fs::path fpath(resolve(path));
+        fs::path fpath = path_from_utf8(resolve(path));
         if (fpath.has_parent_path()) {
             fs::create_directories(fpath.parent_path(), ec);
             if (ec) return false;
@@ -213,13 +234,13 @@ public:
         return (bool) f;
     }
 
-    std::vector<list_entry> list_entries(const std::string & base, const std::string & caller_path, int max_depth, list_kind kind, std::string & err, bool & truncated) const override {
-        err.clear();
-        truncated = false;
+    list_result list_entries(const std::string & base, int max_depth, list_kind kind) const override {
+        list_result out;
+
         std::error_code ec;
         if (!fs::is_directory(base, ec) || ec) {
-            err = "path does not exist or is not a directory: " + caller_path;
-            return {};
+            out.err = "path does not exist or is not a directory";
+            return out;
         }
 
         const auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(SERVER_TOOL_LIST_ENTRIES_TIMEOUT);
@@ -231,7 +252,6 @@ public:
                 SERVER_TOOL_GIT_LS_FILES_MAX_OUTPUT, SERVER_TOOL_LIST_ENTRIES_TIMEOUT);
 
             if (res.exit_code == 0 && !res.timed_out) {
-                std::vector<list_entry> result;
                 std::istringstream iss(res.output);
                 std::string line;
                 while (std::getline(iss, line)) {
@@ -239,15 +259,16 @@ public:
                     if (line.empty()) continue;
                     std::replace(line.begin(), line.end(), '\\', '/');
                     if (max_depth > 0 && entry_depth(line) > max_depth) continue;
-                    if (is_regular_file((fs::path(base) / line).string())) {
-                        result.push_back({line, false});
+                    if (is_regular_file(path_to_utf8(path_from_utf8(base) / path_from_utf8(line)))) {
+                        out.entries.push_back({line, false});
                     }
                 }
-                return result;
+                return out;
             }
         }
 
-        return list_entries_fallback(base, max_depth, kind, deadline, truncated);
+        out.entries = list_entries_fallback(base, max_depth, kind, deadline, out.truncated);
+        return out;
     }
 
     exec_result run(
@@ -326,6 +347,42 @@ public:
 private:
     std::string cwd;
 
+    // a link can point back to an ancestor and loop forever, so it is never walked
+    static bool is_link(const fs::directory_entry & entry) {
+        std::error_code ec;
+        if (entry.is_symlink(ec) || ec) {
+            return true;
+        }
+#if defined(_WIN32)
+        // a junction looks like a plain directory to std::filesystem, so read the reparse tag
+        WIN32_FIND_DATAW data;
+        const HANDLE h = FindFirstFileW(entry.path().c_str(), &data);
+        if (h == INVALID_HANDLE_VALUE) {
+            return false;
+        }
+        FindClose(h);
+        if ((data.dwFileAttributes & FILE_ATTRIBUTE_REPARSE_POINT) == 0) {
+            return false;
+        }
+        // other reparse points (cloud placeholder, dedup stub) are real directories
+        return data.dwReserved0 == IO_REPARSE_TAG_SYMLINK || data.dwReserved0 == IO_REPARSE_TAG_MOUNT_POINT;
+#else
+        return false;
+#endif
+    }
+
+    // NTFS is case insensitive, so Build and build are the same directory
+    static std::string get_effective_name(const std::string & fname) {
+#if defined(_WIN32)
+        std::string lowered = fname;
+        std::transform(lowered.begin(), lowered.end(), lowered.begin(),
+                       [](unsigned char c) { return (char) std::tolower(c); });
+        return lowered;
+#else
+        return fname;
+#endif
+    }
+
     static const std::unordered_set<std::string> & junk_dir_names() {
         static const std::unordered_set<std::string> names = {
             ".git", ".svn", ".hg", "node_modules", "__pycache__",
@@ -337,46 +394,53 @@ private:
     std::vector<list_entry> list_entries_fallback(const std::string & base, int max_depth, list_kind kind,
                                                   std::chrono::steady_clock::time_point deadline, bool & truncated) const {
         std::vector<list_entry> result;
-        std::error_code ec;
 
         std::vector<std::tuple<fs::path, fs::path, int>> stack;
-        stack.emplace_back(fs::path(base), fs::path(), 0);
+        stack.emplace_back(path_from_utf8(base), fs::path(), 0);
 
         while (!stack.empty()) {
-            auto [dir, rel_dir, depth] = stack.back();
+            if (std::chrono::steady_clock::now() >= deadline) {
+                truncated = true;
+                return result;
+            }
+
+            auto [dir, rel_dir, depth] = std::move(stack.back());
             stack.pop_back();
 
-            // the throwing increment would escape the tool on a directory that
-            // goes away mid walk, so step the iterator explicitly
+            std::error_code ec;
+            // step the iterator by hand: the throwing increment escapes on a directory that goes away
             fs::directory_iterator it(dir, fs::directory_options::skip_permission_denied, ec);
+            // permission errors are skipped above, so this is a subtree the caller never sees
+            if (ec) {
+                truncated = true;
+                continue;
+            }
             for (const fs::directory_iterator end; it != end; it.increment(ec)) {
-                if (ec) break;
+                if (ec) {
+                    truncated = true;
+                    break;
+                }
                 if (std::chrono::steady_clock::now() >= deadline) {
                     truncated = true;
                     return result;
                 }
                 const fs::directory_entry & entry = *it;
-                std::string fname = entry.path().filename().string();
+                const fs::path fname = entry.path().filename();
                 std::error_code tec;
-                if (entry.is_directory(tec)) {
-                    std::string rel = (rel_dir / fname).string();
-                    std::replace(rel.begin(), rel.end(), '\\', '/');
+                const bool is_dir = entry.is_directory(tec);
+                if (tec) continue;
+                if (is_dir) {
                     if (kind == list_kind::dirs || kind == list_kind::all) {
-                        result.push_back({rel, true});
+                        result.push_back({path_to_utf8(rel_dir / fname), true});
                     }
-                    // junk directories stay selectable but are never walked: they
-                    // hold nothing worth searching and can be enormous
-                    if (junk_dir_names().count(fname) > 0) continue;
-                    // do not descend into symlinks: a link can point back to an
-                    // ancestor and loop forever
-                    if (!entry.is_symlink(tec) && (max_depth == 0 || depth + 1 < max_depth)) {
+                    // junk directories stay selectable but are never walked: they can be enormous
+                    if (junk_dir_names().count(get_effective_name(path_to_utf8(fname))) > 0) continue;
+                    if (!is_link(entry) && (max_depth == 0 || depth + 1 < max_depth)) {
                         stack.emplace_back(entry.path(), rel_dir / fname, depth + 1);
                     }
                 } else if (entry.is_regular_file(tec)) {
-                    std::string rel = (rel_dir / fname).string();
-                    std::replace(rel.begin(), rel.end(), '\\', '/');
                     if (kind == list_kind::files || kind == list_kind::all) {
-                        result.push_back({rel, false});
+                        result.push_back({path_to_utf8(rel_dir / fname), false});
                     }
                 }
             }
@@ -394,7 +458,7 @@ static std::unique_ptr<tools_io> make_tools_io(const json & params) {
 // no '/' in pattern -> match basename at any depth; else match full relative path
 static bool path_glob_match(const std::string & pattern, const std::string & rel_path) {
     if (pattern.find('/') == std::string::npos) {
-        return glob_match(pattern, fs::path(rel_path).filename().string());
+        return glob_match(pattern, path_to_utf8(path_from_utf8(rel_path).filename()));
     }
     if (pattern == "**" || pattern.rfind("**/", 0) == 0 || pattern.rfind('/', 0) == 0) {
         return glob_match(pattern, rel_path);
@@ -491,7 +555,7 @@ struct server_tool_read_file : server_tool {
 // file_glob_search: find files matching a glob pattern under a base directory
 //
 
-static constexpr size_t SERVER_TOOL_FILE_SEARCH_MAX_RESULTS = 100;
+static constexpr int SERVER_TOOL_FILE_SEARCH_MAX_RESULTS = 100;
 static constexpr const char * SERVER_TOOL_FILE_SEARCH_TYPE_FILE = "file";
 static constexpr const char * SERVER_TOOL_FILE_SEARCH_TYPE_DIR  = "dir";
 static constexpr const char * SERVER_TOOL_FILE_SEARCH_TYPE_ALL  = "all";
@@ -525,7 +589,7 @@ struct server_tool_file_glob_search : server_tool {
                         {"exclude",   {{"type", "string"},  {"description", "Glob pattern for files to exclude"}}},
                         {"type",      {{"type", "string"},  {"description", "Entry type to return: \"file\" (default), \"dir\" or \"all\""}}},
                         {"max_depth", {{"type", "integer"}, {"description", "Maximum depth to descend into subdirectories (default: 0 = unlimited; 1 = direct children only)"}}},
-                        {"limit",     {{"type", "integer"}, {"description", string_format("Maximum number of results to return (default %zu; values below 1 fall back to the default)", SERVER_TOOL_FILE_SEARCH_MAX_RESULTS)}}},
+                        {"limit",     {{"type", "integer"}, {"description", string_format("Maximum number of results to return, capped at %d (default %d)", SERVER_TOOL_FILE_SEARCH_MAX_RESULTS, SERVER_TOOL_FILE_SEARCH_MAX_RESULTS)}}},
                     }},
                     {"required", json::array({"path"})},
                 }},
@@ -536,17 +600,18 @@ struct server_tool_file_glob_search : server_tool {
     json invoke(json params, server_tool::stream *) const override {
         auto io = make_tools_io(params);
 
-        std::string base      = io->resolve(params.at("path").get<std::string>());
-        // normalize to forward slashes so the web UI (which assumes '/') can
-        // join the relative entries into absolute paths on Windows too
-        std::replace(base.begin(), base.end(), '\\', '/');
+        const std::string path = params.at("path").get<std::string>();
+
+        std::string base      = io->resolve(path);
         std::string include   = json_value(params, "include", std::string("**"));
         std::string exclude   = json_value(params, "exclude", std::string(""));
         std::string type      = json_value(params, "type",    std::string("file"));
         int         max_depth = std::max(0, json_value(params, "max_depth", 0));
-        int         limit     = json_value(params, "limit", (int) SERVER_TOOL_FILE_SEARCH_MAX_RESULTS);
-        if (limit < 1) limit = SERVER_TOOL_FILE_SEARCH_MAX_RESULTS;
-        limit = std::min(limit, (int) SERVER_TOOL_FILE_SEARCH_MAX_RESULTS);
+        const int   limit_req = json_value(params, "limit", SERVER_TOOL_FILE_SEARCH_MAX_RESULTS);
+        if (limit_req < 1) {
+            return {{"error", "invalid limit: " + std::to_string(limit_req) + " (expected 1 or more)"}};
+        }
+        const int   limit     = std::min(limit_req, SERVER_TOOL_FILE_SEARCH_MAX_RESULTS);
 
         list_kind kind;
         if (type == SERVER_TOOL_FILE_SEARCH_TYPE_FILE) {
@@ -559,15 +624,13 @@ struct server_tool_file_glob_search : server_tool {
             return {{"error", "invalid type: " + type + " (expected \"file\", \"dir\" or \"all\")"}};
         }
 
-        std::string err;
-        bool truncated = false;
-        auto entries = io->list_entries(base, params.at("path").get<std::string>(), max_depth, kind, err, truncated);
-        if (!err.empty()) {
-            return {{"error", err}};
+        const auto listing = io->list_entries(base, max_depth, kind);
+        if (!listing.err.empty()) {
+            return {{"error", listing.err + ": " + path}};
         }
 
         std::vector<tools_io::list_entry> matches;
-        for (const auto & entry : entries) {
+        for (const auto & entry : listing.entries) {
             if (!path_glob_match(include, entry.rel)) continue;
             if (!exclude.empty() && path_glob_match(exclude, entry.rel)) continue;
             matches.push_back(entry);
@@ -592,8 +655,8 @@ struct server_tool_file_glob_search : server_tool {
                 "[%zu results limit reached (%zu total matches). Refine the glob pattern to narrow the search.]\n",
                 shown, total);
         }
-        if (truncated) {
-            output_text << "[search timed out, results truncated]\n";
+        if (listing.truncated) {
+            output_text << "[results truncated: time budget or unreadable directory]\n";
         }
 
         // `base` is always absolute (resolve falls back to the server cwd), so
@@ -688,16 +751,14 @@ struct server_tool_grep_search : server_tool {
         if (io->is_regular_file(abs_path)) {
             files.emplace_back(abs_path, path);
         } else if (io->is_directory(abs_path)) {
-            std::string err;
-            bool truncated = false;
-            auto candidates = io->list_entries(abs_path, path, 0, list_kind::files, err, truncated);
-            if (!err.empty()) {
-                return {{"error", err}};
+            const auto listing = io->list_entries(abs_path, 0, list_kind::files);
+            if (!listing.err.empty()) {
+                return {{"error", listing.err + ": " + path}};
             }
-            for (const auto & entry : candidates) {
+            for (const auto & entry : listing.entries) {
                 if (!path_glob_match(include, entry.rel)) continue;
                 if (!exclude.empty() && path_glob_match(exclude, entry.rel)) continue;
-                files.emplace_back((fs::path(abs_path) / entry.rel).string(), entry.rel);
+                files.emplace_back(path_to_utf8(path_from_utf8(abs_path) / path_from_utf8(entry.rel)), entry.rel);
             }
         } else {
             return {{"error", "path does not exist: " + path}};
@@ -1306,7 +1367,7 @@ struct server_tool_get_info : server_tool {
         std::string cwd = json_value(params, "cwd", std::string());
         if (cwd.empty()) {
             std::error_code ec;
-            cwd = fs::current_path(ec).string();
+            cwd = path_to_utf8(fs::current_path(ec));
         }
 
         return {

tools/server/tests/unit/test_proxy.py

@@ -1,5 +1,7 @@
 import pytest
 from utils import *
+import threading
+from http.server import BaseHTTPRequestHandler, ThreadingHTTPServer
 
 server = ServerPreset.tinyllama2()
 
@@ -39,3 +41,31 @@ def test_mcp_proxy_custom_port():
     res = server.make_request("GET", f"/cors-proxy?url=http://{server.server_host}:{server.server_port}/models")
     assert res.status_code == 200
     assert "data" in res.body
+
+
+def test_mcp_proxy_no_content():
+    # note: see issue #26598
+    class NoContentHandler(BaseHTTPRequestHandler):
+        def do_POST(self):
+            self.send_response(204)
+            self.end_headers()
+
+        def log_message(self, format, *args):
+            pass
+
+    target = ThreadingHTTPServer(("127.0.0.1", 0), NoContentHandler)
+    target_thread = threading.Thread(target=target.serve_forever, daemon=True)
+    target_thread.start()
+
+    try:
+        global server
+        server.ui_mcp_proxy = True
+        server.start()
+
+        res = server.make_request("POST", f"/cors-proxy?url=http://127.0.0.1:{target.server_port}/", data={})
+        assert res.status_code == 204
+        assert res.body in (None, b"", "")
+    finally:
+        target.shutdown()
+        target.server_close()
+

tools/server/tests/unit/test_tools_builtin.py

@@ -214,6 +214,27 @@ def test_tools_builtin_file_glob_search_max_depth_and_limit(tmp_path):
     assert "Total matches: 3" in res["plain_text_response"]
 
 
+def test_tools_builtin_file_glob_search_junk_dirs(tmp_path):
+    global server
+    server.start()
+
+    (tmp_path / "build" / "nested").mkdir(parents=True)
+    (tmp_path / "build" / "artifact.txt").write_text("built")
+    (tmp_path / "src").mkdir()
+    (tmp_path / "src" / "main.cpp").write_text("int main() {}")
+
+    # a junk directory stays selectable as a working directory
+    res = call_tool("file_glob_search", {"path": str(tmp_path), "type": "dir", "max_depth": 1})
+    assert "build" in [e["path"] for e in res["entries"]]
+
+    # but it is never walked, so nothing inside it shows up
+    res = call_tool("file_glob_search", {"path": str(tmp_path), "type": "all"})
+    paths = [e["path"] for e in res["entries"]]
+    assert "src/main.cpp" in paths
+    assert "build/artifact.txt" not in paths
+    assert "build/nested" not in paths
+
+
 def test_tools_builtin_file_glob_search_rejects_invalid_type(tmp_path):
     global server
     server.start()

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormWorkingDirectory.svelte

@@ -87,8 +87,9 @@
 	let searchSeq = 0;
 
 	// Cache of the last file_glob_search result per (parent, include, max_depth),
-	// so repeated queries in the same directory don't re-walk the tree. Entries
-	// expire after a short TTL.
+	// so repeated queries in the same directory don't re-walk the tree. Entering
+	// a directory hits it every time: the children listed for an exactly typed
+	// segment are what the next keystroke, the trailing slash, asks for again.
 	const SEARCH_CACHE_TTL_MS = 2000;
 	const searchCache = new SvelteMap<string, { results: GlobEntry[]; base: string; at: number }>();
 
@@ -161,7 +162,11 @@
 		if (typeof res.error === 'string') return { base: '', entries: [], error: res.error };
 		const base = typeof res.base === 'string' ? res.base : '';
 		const entries = Array.isArray(res.entries) ? (res.entries as GlobEntry[]) : [];
-		searchCache.set(key, { results: entries, base, at: Date.now() });
+		const now = Date.now();
+		for (const [k, v] of searchCache) {
+			if (now - v.at >= SEARCH_CACHE_TTL_MS) searchCache.delete(k);
+		}
+		searchCache.set(key, { results: entries, base, at: now });
 		return { base, entries };
 	}
 

tools/ui/src/lib/stores/models.svelte.ts

@@ -1,7 +1,12 @@
 import { base } from '$app/paths';
 import { SvelteMap, SvelteSet } from 'svelte/reactivity';
 import { toast } from 'svelte-sonner';
-import { ServerModelStatus, ServerModelsSseEventType, ModelModality } from '$lib/enums';
+import {
+	ServerModelStatus,
+	ServerModelsSseEventType,
+	ModelModality,
+	FileTypeCategory
+} from '$lib/enums';
 import { ModelsService } from '$lib/services/models.service';
 import { PropsService } from '$lib/services/props.service';
 import { serverStore, isRouterMode } from '$lib/stores/server.svelte';
@@ -394,6 +399,7 @@ class ModelsStore {
 				model: modelId,
 				description: details?.description,
 				capabilities: rawCapabilities.filter((value: unknown): value is string => Boolean(value)),
+				modalities: this.buildArchitectureModalities(item.architecture),
 				details: details?.details,
 				meta: item.meta ?? null,
 				parsedId: ModelsService.parseModelId(modelId),
@@ -1000,6 +1006,21 @@ class ModelsStore {
 		};
 	}
 
+	/** Map the router modalities, the only source available while a model is not loaded. */
+	private buildArchitectureModalities(
+		architecture: ApiModelDataEntry['architecture']
+	): ModelModalities | undefined {
+		if (!architecture) return undefined;
+
+		const inputs = architecture.input_modalities;
+
+		return {
+			vision: inputs.includes(FileTypeCategory.IMAGE),
+			audio: inputs.includes(FileTypeCategory.AUDIO),
+			video: inputs.includes(FileTypeCategory.VIDEO)
+		};
+	}
+
 	clear(): void {
 		this.unsubscribeStatus();
 		this.statusWaiters.forEach((waiter) => waiter.reject(new Error('Models store cleared')));

tools/ui/src/lib/types/api.d.ts

@@ -98,10 +98,21 @@ export interface ApiModelDataEntry {
 	aliases?: string[];
 	/** Informational tags for this model */
 	tags?: string[];
+	/** Modality capabilities, reported by the router for every model regardless of load state */
+	architecture?: ApiModelArchitecture;
 	/** Legacy meta field (may be present in older responses) */
 	meta?: Record<string, unknown> | null;
 }
 
+/**
+ * Modality capabilities of a model, as advertised by the ROUTER /models endpoint.
+ * Read from the model manifest, so it is available before the model is loaded.
+ */
+export interface ApiModelArchitecture {
+	/** Accepted input modalities, always contains "text" */
+	input_modalities: string[];
+}
+
 /**
  * Load stage reported by the /models/sse feed, in load order.
  */