tests : make MoE router selection independent of the backend noise (#30818)

* tests : make MoE router selection independent of the backend noise

The router scores are in the range [0, 1), so a random draw can leave the
last expert (or expert group) still selected and the first one dropped tied
within the noise between the backends. They then route a token to different
experts and the logits differ by 4e-3 instead of the usual 5e-7.

Beam is the only arch here with a real top-k (8 experts, 2 used) and a group
top-k (4 groups, 2 used), so it is the only one that hits this: 2 of 20 seeds
failed on an M2 Ultra, the other MoE archs select 2 of 2 and cannot tie.

Fill the selection bias with a ramp of step 2 to give every expert a score
range of its own, which turns the top-k into a function of the expert index
alone. The bias affects only the selection, the weights still come from the
unbiased probs, so the MoE output keeps its magnitude. Beam now matches the
CPU reference within 4e-7 for all of 160 seeds.

Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)

* tests : tell the CPU and the GPU of an Apple SoC apart in the device table

Both devices report the same description ("Apple M2 Ultra" here), so the
device column of the test table had two indistinguishable rows and a failing
row did not say whether the Metal device or the host CPU produced it.

Add the device name, which is also what -b takes, unless the description
already is the name.

Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)

.flake8

@@ -5,8 +5,8 @@ exclude =
     # Do not traverse examples and tools
     examples,
     tools,
-    # Reference scripts for skills
-    skills,
+    # Do not traverse .venv directories
+    .venv*,
     # Do not include package initializers
     __init__.py,
     # No need to traverse our git directory

.github/workflows/build-and-test-snapdragon.yml

@@ -35,7 +35,7 @@ jobs:
   android-ndk-snapdragon:
     runs-on: ubuntu-24.04  # previously ubuntu-latest
     container:
-      image: 'ghcr.io/snapdragon-toolchain/arm64-android:v0.7'
+      image: 'ghcr.io/snapdragon-toolchain/arm64-android:v0.8'
     defaults:
       run:
         shell: bash
@@ -69,7 +69,7 @@ jobs:
   linux-iot-snapdragon:
     runs-on: ubuntu-24.04  # previously ubuntu-latest
     container:
-      image: 'ghcr.io/snapdragon-toolchain/arm64-linux:v0.7'
+      image: 'ghcr.io/snapdragon-toolchain/arm64-linux:v0.8'
     defaults:
       run:
         shell: bash

.github/workflows/build-wasi.yml

@@ -0,0 +1,142 @@
+name: CI (wasi)
+
+on:
+  workflow_dispatch: # allows manual triggering
+  push:
+    branches:
+      - master
+    paths: [
+      '.github/workflows/build-wasi.yml',
+      '**/CMakeLists.txt',
+      '**/.cmake',
+      '**/*.h',
+      '**/*.hpp',
+      '**/*.c',
+      '**/*.cpp',
+      '**/*.wgsl',
+      '**/*.tmpl',
+      'ggml/src/ggml-webgpu/wgsl-shaders/embed_wgsl.py'
+    ]
+
+  pull_request:
+    types: [opened, synchronize, reopened]
+    paths: [
+      '.github/workflows/build-wasi.yml',
+      '**/CMakeLists.txt',
+      '**/.cmake',
+      '**/*.h',
+      '**/*.hpp',
+      '**/*.c',
+      '**/*.cpp',
+      '**/*.wgsl',
+      '**/*.tmpl',
+      'ggml/src/ggml-webgpu/wgsl-shaders/embed_wgsl.py'
+    ]
+
+cache-mode: none
+permissions:
+  contents: read
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
+  cancel-in-progress: true
+
+env:
+  WASI_SDK_VERSION: 34
+  DAWN_VERSION: v20260908.214631
+  DAWN_COMMIT: 94c3c9cc0d5fb2e85aebb370fa8d37b71aa34655
+  WASI_WEBGPU_HEADERS_COMMIT: 917765edd7cb8f8cf95bda1217514c78f71b582b
+
+jobs:
+  ubuntu-webgpu:
+    runs-on: ubuntu-24.04-arm
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+
+      - name: ccache
+        uses: ggml-org/ccache-action@v1.2.24
+        with:
+          restore: false
+          save: false
+
+      - name: Install wasi-sdk
+        uses: bytecodealliance/setup-wasi-sdk-action@77d09f3df1e8da344d8403cf8631af8977938b4c # v1.0.0
+        with:
+          version: ${{ env.WASI_SDK_VERSION }}
+          add-to-path: false
+
+      - name: wasi:webgpu Dependency
+        run: |
+          mkdir -p wasi-webgpu/dawn-headers
+          curl -fL "https://github.com/google/dawn/releases/download/${DAWN_VERSION}/dawn-headers-${DAWN_COMMIT}.tar.gz" |
+            tar -xz -C wasi-webgpu/dawn-headers --strip-components=1
+
+          git init -q wasi-webgpu/wasi-webgpu-headers
+          cd wasi-webgpu/wasi-webgpu-headers
+          git fetch -q --depth 1 https://github.com/wasi-gfx/wasi-webgpu-headers.git ${WASI_WEBGPU_HEADERS_COMMIT}
+          git checkout -q FETCH_HEAD
+          git submodule update -q --init --depth 1
+          cd ../..
+
+          # ggml-webgpu does find_package(Dawn) and links dawn::webgpu_dawn. There is no Dawn on WASI,
+          # so this target instead compiles wasi-webgpu-headers (webgpu.h over wasi:webgpu imports)
+          # and exposes Dawn's webgpu_cpp.h on top of it.
+          cat > wasi-webgpu/DawnConfig.cmake << 'EOF'
+          set(_wasi_webgpu "${CMAKE_CURRENT_LIST_DIR}/wasi-webgpu-headers")
+          add_library(dawn::webgpu_dawn INTERFACE IMPORTED)
+          set_target_properties(dawn::webgpu_dawn PROPERTIES
+              INTERFACE_SOURCES "${_wasi_webgpu}/webgpu.c;${_wasi_webgpu}/async_futures.c;${_wasi_webgpu}/imports.c;${_wasi_webgpu}/imports_component_type.o"
+              INTERFACE_INCLUDE_DIRECTORIES "${CMAKE_CURRENT_LIST_DIR}/dawn-headers/include/webgpu_upstream;${_wasi_webgpu}")
+          EOF
+
+      - name: ccache-buckets-restore
+        uses: ./.github/actions/ccache-buckets
+        env:
+          HF_TOKEN: ${{ secrets.HF_TOKEN_CI }}
+        with:
+          key: webgpu-ubuntu-24.04-arm-wasi
+          folder: llama.cpp
+          hf_bucket: ggml-org/cache
+
+      # llama-common does not build for WASI, so the executables are linked by hand
+      - name: Build
+        id: cmake_build
+        run: |
+          WASI_FLAGS="-msimd128 -D_WASI_EMULATED_SIGNAL -fwasm-exceptions -mllvm -wasm-use-legacy-eh=false"
+          cmake -B build-wasi \
+            -G "Ninja" \
+            -DCMAKE_TOOLCHAIN_FILE=${WASI_SDK_PATH}/share/cmake/wasi-sdk-p3.cmake \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DCMAKE_C_FLAGS="${WASI_FLAGS}" \
+            -DCMAKE_CXX_FLAGS="${WASI_FLAGS}" \
+            -DCMAKE_HAVE_LIBC_PTHREAD=1 \
+            -DBUILD_SHARED_LIBS=OFF \
+            -DGGML_NATIVE=OFF \
+            -DGGML_OPENMP=OFF \
+            -DGGML_WEBGPU=ON \
+            -DDawn_DIR=$PWD/wasi-webgpu \
+            -DLLAMA_BUILD_COMMON=OFF
+
+          time cmake --build build-wasi --config Release --target llama -j $(nproc)
+
+          LIBS="build-wasi/ggml/src/libggml.a build-wasi/ggml/src/ggml-webgpu/libggml-webgpu.a build-wasi/ggml/src/libggml-cpu.a build-wasi/ggml/src/libggml-base.a"
+          LDFLAGS="-lunwind -ldl -Wl,-z,stack-size=8388608"
+          CXX="${WASI_SDK_PATH}/bin/clang++ --target=wasm32-wasip3 -O3 -std=c++17 ${WASI_FLAGS} -I include -I ggml/include"
+
+          $CXX examples/simple/simple.cpp build-wasi/src/libllama.a $LIBS $LDFLAGS -o build-wasi/llama-simple.wasm
+          $CXX tests/test-backend-ops.cpp $LIBS $LDFLAGS -o build-wasi/test-backend-ops.wasm
+
+      - name: ccache-buckets-save
+        if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
+        uses: ./.github/actions/ccache-buckets
+        env:
+          HF_TOKEN: ${{ secrets.HF_TOKEN_CACHE_OUTPUT }}
+        with:
+          key: webgpu-ubuntu-24.04-arm-wasi
+          folder: llama.cpp
+          evict-old-files: 1d
+          hf_bucket: ggml-org/cache
+          save: true

.github/workflows/release.yml

@@ -35,7 +35,7 @@ permissions:
 env:
   GH_TOKEN: ${{ github.token }}
   BRANCH_NAME: ${{ github.head_ref || github.ref_name }}
-  CMAKE_ARGS: "-DLLAMA_BUILD_EXAMPLES=OFF -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_TOOLS=ON -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON"
+  CMAKE_ARGS: "-DLLAMA_BUILD_EXAMPLES=OFF -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_TOOLS=ON -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DLLAMA_CONNECT=ON"
 
 # note: run this workflow one at a time for better cache reuse
 concurrency:
@@ -547,7 +547,7 @@ jobs:
     if: ${{ needs.check-release.outputs.should_release == 'true' }}
 
     runs-on: ubuntu-latest
-    container: 'ghcr.io/snapdragon-toolchain/arm64-android:v0.7'
+    container: 'ghcr.io/snapdragon-toolchain/arm64-android:v0.8'
 
     defaults:
       run:
@@ -600,7 +600,7 @@ jobs:
     if: ${{ needs.check-release.outputs.should_release == 'true' }}
 
     runs-on: ubuntu-latest
-    container: 'ghcr.io/snapdragon-toolchain/arm64-linux:v0.7'
+    container: 'ghcr.io/snapdragon-toolchain/arm64-linux:v0.8'
 
     defaults:
       run:
@@ -1478,6 +1478,7 @@ jobs:
             -DCMAKE_INSTALL_RPATH='$ORIGIN' \
             -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
             -DGGML_NATIVE=OFF \
+            -DLLAMA_CONNECT=ON \
             -DGGML_SYCL_F16=${{ matrix.fp16 }}
           time cmake --build build --config Release -j $(nproc)
 

CMakeLists.txt

@@ -136,6 +136,7 @@ option(LLAMA_BUILD_SERVER    "llama: build server example"
 option(LLAMA_BUILD_APP       "llama: build the unified binary"                                                   ${LLAMA_STANDALONE})
 option(LLAMA_BUILD_UI        "llama: build the embedded Web UI for server"                                       OFF)
 option(LLAMA_USE_PREBUILT_UI "llama: use prebuilt UI from HF Bucket when available" ON)
+option(LLAMA_CONNECT         "llama: include the llama-connect tool (downloads a prebuilt binary)"               OFF)
 
 option(LLAMA_TOOLS_INSTALL "llama: install tools" ${LLAMA_TOOLS_INSTALL_DEFAULT})
 option(LLAMA_TESTS_INSTALL "llama: install tests" ON)

cmake/common.cmake

@@ -2,7 +2,7 @@ include("ggml/cmake/common.cmake")
 
 function(llama_add_compile_flags)
     if (LLAMA_FATAL_WARNINGS)
-        if (CMAKE_CXX_COMPILER_ID MATCHES "GNU" OR CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+        if (CMAKE_CXX_COMPILER_ID MATCHES "GNU" OR CMAKE_CXX_COMPILER_ID MATCHES "Clang" OR CMAKE_CXX_COMPILER_ID MATCHES "IntelLLVM")
             list(APPEND C_FLAGS   -Werror)
             list(APPEND CXX_FLAGS -Werror)
         elseif (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")

common/arg.cpp

@@ -3494,6 +3494,26 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.ui = value;
         }
     ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_UI"));
+    add_opt(common_arg(
+        {"--connect"},
+        string_format("(experimental) open a peer-to-peer tunnel for server using llama-connect (default: %s)", params.server_connect ? "enabled" : "disabled"),
+        [](common_params & params) {
+            params.server_connect = true;
+        }
+    ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CONNECT"));
+    add_opt(common_arg(
+        {"--connect-code"}, "CODE",
+        "manually specify a 40-character code for --connect (default: generate a new one for each run)",
+        [](common_params & params, const std::string & value) {
+            std::string val = value;
+            string_replace_all(val, "-", "");
+            string_replace_all(val, " ", "");
+            if (val.size() != 40) {
+                throw std::invalid_argument(string_format("error: invalid connect code '%s', must be 40 characters\n", value.c_str()));
+            }
+            params.server_connect_code = val;
+        }
+    ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CONNECT_CODE"));
     add_opt(common_arg(
         {"--embedding", "--embeddings"},
         string_format("restrict to only support embedding use case; use only with dedicated embedding models (default: %s)", params.embedding ? "enabled" : "disabled"),

common/common.h

@@ -642,6 +642,10 @@ struct common_params {
     int32_t cache_ram_mib       = 8192;  // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.
     bool    cache_ram_explicit  = false; // set by --cache-ram, decision models disable the cache otherwise
 
+    // llama-connect params
+    bool        server_connect      = false;
+    std::string server_connect_code = "";
+
     std::string public_path   = "";                                                                         // NOLINT
     std::string api_prefix    = "";                                                                         // NOLINT
     std::string chat_template = "";                                                                         // NOLINT

common/jinja/value.cpp

@@ -1028,10 +1028,6 @@ const func_builtins & value_bool_t::get_builtins() const {
     return builtins;
 }
 
-[[noreturn]] static value array_unique_not_implemented(const func_args &) {
-    throw not_implemented_exception("Array unique builtin not implemented");
-}
-
 const func_builtins & value_array_t::get_builtins() const {
     static const func_builtins builtins = {
         {"default", default_value},
@@ -1257,7 +1253,36 @@ const func_builtins & value_array_t::get_builtins() const {
             }
             return result;
         }},
-        {"unique", array_unique_not_implemented},
+        {"unique", [](const func_args & args) -> value {
+            args.ensure_count(1, 3);
+            args.ensure_vals<value_array>();
+            value val       = args.get_pos(0);
+            value val_case  = args.get_kwarg_or_pos("case_sensitive", 1);
+            value attribute = args.get_kwarg_or_pos("attribute",      2);
+            // FIXME: unique is currently always case sensitive
+            //const bool case_sensitive = val_case->as_bool(); // undefined == false
+            value undef = mk_val<value_undefined>();
+            std::vector<value> result;
+            std::vector<value> seen;
+            for (const auto & item : val->as_array()) {
+                value key = item;
+                if (!attribute->is_undefined()) {
+                    key = get_attribute(item, attribute, undef);
+                }
+                bool dup = false;
+                for (const auto & s : seen) {
+                    if (value_compare(s, key, value_compare_op::eq)) {
+                        dup = true;
+                        break;
+                    }
+                }
+                if (!dup) {
+                    seen.push_back(key);
+                    result.push_back(item);
+                }
+            }
+            return mk_val<value_array>(std::move(result));
+        }},
     };
     return builtins;
 }

conversion/base.py

@@ -10,6 +10,7 @@ import json
 import os
 import re
 import sys
+import tempfile
 from enum import IntEnum
 from pathlib import Path
 from hashlib import sha256
@@ -2169,20 +2170,34 @@ class TextModel(ModelBase):
         special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
         special_vocab.add_to_gguf(self.gguf_writer)
 
+    def _is_hybriddna(self) -> bool:
+        tokenizer_config_file = self.dir_model / "tokenizer_config.json"
+        if not tokenizer_config_file.is_file():
+            return False
+        with open(tokenizer_config_file, "r", encoding="utf-8") as f:
+            return json.load(f).get("tokenizer_class") == "HybridDNATokenizer"
+
     def _set_vocab_hybriddna(self):
         from transformers import AutoTokenizer
         tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
-        vocab_size = self.hparams.get("vocab_size", len(tokenizer.vocab))  # ty: ignore[unresolved-attribute]
-        assert max(tokenizer.vocab.values()) < vocab_size  # ty: ignore[unresolved-attribute]
+        vocab = tokenizer.get_vocab()
+        vocab_size = self.hparams.get("vocab_size", len(vocab))
+        assert max(vocab.values()) < vocab_size
 
-        reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}  # ty: ignore[unresolved-attribute]
-        # k-mers can share text with a base-vocab BPE token (e.g. CCCCCC) and get
-        # dropped by get_vocab(); a reserved marker suffix (U+E000) keeps each
-        # k-mer's own id (llama.cpp strips it on detokenization)
-        for kmer in tokenizer.kmers:  # ty: ignore[unresolved-attribute]
+        reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in vocab.items()}
+        # k-mers can share text with a base-vocab BPE token (e.g. CCCCCC) and get dropped by get_vocab()
+        # a reserved marker suffix (U+E000) keeps each k-mer's own id (llama.cpp strips it on detokenization)
+        # tail="pad_n" (Carbon-14B) adds N-padded tail tokens ("CT" -> "CTNNNN") that get the same marker
+        # llama.cpp detects them to pad with N and strips the padding on detokenization, tail="char" is unsupported
+        tail = getattr(tokenizer, "tail", "pad_a")
+        if tokenizer.k != 6 or tail not in ("pad_a", "pad_n"):
+            raise NotImplementedError(f"HybridDNATokenizer with k={tokenizer.k}, tail={tail!r}")
+        for kmer in [*tokenizer.kmers, *getattr(tokenizer, "npad_tokens", [])]:
             reverse_vocab[tokenizer.dna_token_to_id[kmer]] = kmer + "\ue000"  # ty: ignore[unresolved-attribute]
-        added_vocab = tokenizer.get_added_vocab()  # ty: ignore[unresolved-attribute]
-        added_tokens_decoder = tokenizer.added_tokens_decoder  # ty: ignore[unresolved-attribute]
+        # older HybridDNATokenizer versions (Carbon-500M/3B/8B) do not register the base tokenizer's added tokens (<|im_start|>, <think>, ...) as their own
+        # but they are still encoded as special tokens: take their types from the base tokenizer
+        base_tokenizer = getattr(tokenizer, "_base_tokenizer", tokenizer)
+        added_tokens_decoder = {**base_tokenizer.added_tokens_decoder, **tokenizer.added_tokens_decoder}  # ty: ignore[unresolved-attribute]
 
         tokens: list[str] = []
         toktypes: list[int] = []
@@ -2192,7 +2207,7 @@ class TextModel(ModelBase):
                 toktypes.append(gguf.TokenType.UNUSED)
             else:
                 token: str = reverse_vocab[i]
-                if token in added_vocab:
+                if i in added_tokens_decoder:
                     if added_tokens_decoder[i].special or self.does_token_look_special(token):
                         toktypes.append(gguf.TokenType.CONTROL)
                     else:
@@ -2204,10 +2219,19 @@ class TextModel(ModelBase):
         tokpre = self.get_vocab_base_pre(tokenizer)
         self.gguf_writer.add_tokenizer_model("hybriddna")
         self.gguf_writer.add_tokenizer_pre(tokpre)
+        self.gguf_writer.add_dna_auto_tags(getattr(tokenizer, "auto_dna_tags", False))
         self.gguf_writer.add_token_list(tokens)
         self.gguf_writer.add_token_types(toktypes)
 
         special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
+        if not special_vocab.merges:
+            # Carbon-500M/3B/8B ship only tokenizer.py, no tokenizer.json
+            # take the merges and special token ids from the base BPE tokenizer, keeping the model's own chat template (if any)
+            with tempfile.TemporaryDirectory() as tmp:
+                base_tokenizer.save_pretrained(tmp)
+                base_special_vocab = gguf.SpecialVocab(Path(tmp), load_merges=True)
+            base_special_vocab.chat_template = special_vocab.chat_template
+            special_vocab = base_special_vocab
         special_vocab.add_to_gguf(self.gguf_writer)
 
     def _set_vocab_qwen(self):

conversion/llama.py

@@ -129,8 +129,9 @@ class LlamaModel(TextModel):
                 tokenizer_config_json = json.load(f)
                 if (add_prefix_space := tokenizer_config_json.get("add_prefix_space")) is not None:
                     self.gguf_writer.add_add_space_prefix(add_prefix_space)
-                if tokenizer_config_json.get("tokenizer_class") == "HybridDNATokenizer":
-                    return self._set_vocab_hybriddna()
+
+        if self._is_hybriddna():
+            return self._set_vocab_hybriddna()
 
         try:
             self._set_vocab_sentencepiece()

conversion/qwen.py

@@ -59,6 +59,10 @@ class Qwen2Model(TextModel):
     model_arch = gguf.MODEL_ARCH.QWEN2
 
     def set_vocab(self):
+        # HuggingFaceBio/Carbon-14B: Qwen3 with the HybridDNATokenizer (Qwen3 BPE plus DNA 6-mers)
+        if self._is_hybriddna():
+            self._set_vocab_hybriddna()
+            return
         try:
             self._set_vocab_sentencepiece()
         except FileNotFoundError:

docs/backend/snapdragon/README.md

@@ -6,8 +6,8 @@ The cross-compilation toolchain images are provided by the
 [Qualcomm Snapdragon Toolchain registry](https://github.com/snapdragon-toolchain).
 These Docker images include the Android NDK, OpenCL SDK, Hexagon SDK, CMake, and the necessary cross-compilers:
 
-* **Android toolchain**: `ghcr.io/snapdragon-toolchain/arm64-android:v0.7`
-* **Linux toolchain**: `ghcr.io/snapdragon-toolchain/arm64-linux:v0.7`
+* **Android toolchain**: `ghcr.io/snapdragon-toolchain/arm64-android:v0.8`
+* **Linux toolchain**: `ghcr.io/snapdragon-toolchain/arm64-linux:v0.8`
 
 The unified build utility (`scripts/snapdragon/build.py`) automatically pulls
 and orchestrates these containers to perform target compilation.
@@ -40,7 +40,7 @@ Alternatively, you can build llama.cpp manually by entering the cross-compilatio
 
 ```bash
 # Start the cross-compilation container manually:
-~/src/llama.cpp$ docker run -it --rm -u $(id -u):$(id -g) --volume $(pwd):/workspace --platform linux/amd64 ghcr.io/snapdragon-toolchain/arm64-android:v0.7
+~/src/llama.cpp$ docker run -it --rm -u $(id -u):$(id -g) --volume $(pwd):/workspace --platform linux/amd64 ghcr.io/snapdragon-toolchain/arm64-android:v0.8
 
 # Inside the container, build the project using presets:
 [d]/workspace> cp docs/backend/snapdragon/CMakeUserPresets.json .
@@ -53,7 +53,7 @@ Preset CMake variables:
   GGML_HEXAGON="ON"
   GGML_OPENCL="ON"
   GGML_OPENMP="OFF"
-  HEXAGON_SDK_ROOT="/opt/hexagon/6.6.0.0"
+  HEXAGON_SDK_ROOT="/opt/hexagon/6.7.0.0"
 ...
 -- Including OpenCL backend
 -- Including Hexagon backend

docs/backend/snapdragon/linux.md

@@ -3,7 +3,7 @@
 The cross-compilation is performed using the Snapdragon Linux Docker toolchain image (see
 [github.com/snapdragon-toolchain](https://github.com/snapdragon-toolchain)):
 
-* **Linux toolchain**: `ghcr.io/snapdragon-toolchain/arm64-linux:v0.7`
+* **Linux toolchain**: `ghcr.io/snapdragon-toolchain/arm64-linux:v0.8`
 
 The unified build utility (`scripts/snapdragon/build.py`) automatically pulls
 and orchestrates this container to perform target compilation. You only need to
@@ -29,7 +29,7 @@ Alternatively, you can build llama.cpp manually by entering the cross-compilatio
 
 ```bash
 # Start the cross-compilation container manually:
-~/src/llama.cpp$ docker run -it --rm -u $(id -u):$(id -g) --volume $(pwd):/workspace --platform linux/amd64 ghcr.io/snapdragon-toolchain/arm64-linux:v0.7
+~/src/llama.cpp$ docker run -it --rm -u $(id -u):$(id -g) --volume $(pwd):/workspace --platform linux/amd64 ghcr.io/snapdragon-toolchain/arm64-linux:v0.8
 
 # Inside the container, build the project using presets:
 [d]/workspace> cp docs/backend/snapdragon/CMakeUserPresets.json .

docs/backend/snapdragon/windows.md

@@ -8,7 +8,7 @@ Native Windows 11 arm64 builds have the following tool dependencies:
   - UCRT and Driver Kit
 - LLVM core libraries and Clang compiler (winget)
 - CMake, Git, Python (winget)
-- Hexagon SDK Community Edition 6.6 or later (see below)
+- Hexagon SDK Community Edition 6.7 or later (see below)
 - OpenCL SDK 2.3 or later (see below)
 
 Note: The rest of the **Windows** build process assumes that you're running natively in Powershell.
@@ -54,15 +54,15 @@ c:\Qualcomm\OpenCL_SDK\2.3.2
 
 Either use the trimmed down version (optimized for CI) from
 
-    https://github.com/snapdragon-toolchain/hexagon-sdk/releases/download/v6.6.0.0/hexagon-sdk-v6.6.0.0-arm64-wos.tar.xz
+    https://github.com/snapdragon-toolchain/hexagon-sdk/releases/download/v6.7.0.0/hexagon-sdk-v6.7.0.0-arm64-wos.tar.xz
 
 Or download the complete official version from
 
-    https://softwarecenter.qualcomm.com/catalog/item/Hexagon_SDK?version=6.6.0.0
+    https://softwarecenter.qualcomm.com/catalog/item/Hexagon_SDK?version=6.7.0.0
 
 Unzip/untar the archive into
 ```
-c:\Qualcomm\Hexagon_SDK\6.6.0.0
+c:\Qualcomm\Hexagon_SDK\6.7.0.0
 ```
 
 ## Install the latest Adreno GPU driver
@@ -150,8 +150,8 @@ The overall Hexagon backend build procedure for Windows on Snapdragon is the sam
 However, additional settings are required for generating and signing HTP Ops libraries.
 ```
 > $env:OPENCL_SDK_ROOT="C:\Qualcomm\OpenCL_SDK\2.3.2"
-> $env:HEXAGON_SDK_ROOT="C:\Qualcomm\Hexagon_SDK\6.6.0.0"
-> $env:HEXAGON_TOOLS_ROOT="C:\Qualcomm\Hexagon_SDK\6.6.0.0\tools\HEXAGON_Tools\19.0.07"
+> $env:HEXAGON_SDK_ROOT="C:\Qualcomm\Hexagon_SDK\6.7.0.0"
+> $env:HEXAGON_TOOLS_ROOT="C:\Qualcomm\Hexagon_SDK\6.7.0.0\tools\HEXAGON_Tools\21.0.04"
 > $env:HEXAGON_HTP_CERT="c:\Users\MyUsers\Certs\ggml-htp-v1.pfx"
 > $env:WINDOWS_SDK_BIN="C:\Program Files (x86)\Windows Kits\10\bin\10.0.26100.0"
 

docs/build.md

@@ -852,6 +852,36 @@ WebGPU allows cross-platform access to the GPU from supported browsers. We utili
 
 Follow the instructions [here](https://dawn.googlesource.com/dawn/+/refs/heads/main/src/emdawnwebgpu/) to download or build the emdawnwebgpu package (Note that it might be safer to build the emdawnwebgpu package locally, so that it stays in sync with the version of Dawn you have installed above). When building using CMake, the path to the emdawnwebgpu port file needs to be set with the flag `EMDAWNWEBGPU_DIR`.
 
+## WASI
+
+Requires [wasi-sdk](https://github.com/WebAssembly/wasi-sdk) 34+ with `WASI_SDK_PATH` set. Only the libraries build (`llama-common` does not support WASI yet), so link your program against them:
+
+```bash
+FLAGS="-msimd128 -D_WASI_EMULATED_SIGNAL -fwasm-exceptions -mllvm -wasm-use-legacy-eh=false"
+cmake -B build-wasi -DCMAKE_TOOLCHAIN_FILE=$WASI_SDK_PATH/share/cmake/wasi-sdk-p2.cmake \
+    -DCMAKE_C_FLAGS="$FLAGS" -DCMAKE_CXX_FLAGS="$FLAGS" -DCMAKE_HAVE_LIBC_PTHREAD=1 \
+    -DBUILD_SHARED_LIBS=OFF -DGGML_NATIVE=OFF -DGGML_OPENMP=OFF -DLLAMA_BUILD_COMMON=OFF
+cmake --build build-wasi --target llama
+
+$WASI_SDK_PATH/bin/clang++ --target=wasm32-wasip3 -O3 $FLAGS -I include -I ggml/include examples/simple/simple.cpp \
+    build-wasi/src/libllama.a build-wasi/ggml/src/libggml*.a -lunwind -ldl -Wl,-z,stack-size=8388608 -o llama-simple.wasm
+wasmtime run -W exceptions=y --dir models llama-simple.wasm -m models/model.gguf "Hello"
+```
+
+### wasi:webgpu
+
+The WebGPU backend can target [wasi:webgpu](https://github.com/WebAssembly/wasi-webgpu) via [wasi-webgpu-headers](https://github.com/wasi-gfx/wasi-webgpu-headers) (with its submodule) and Dawn's `dawn-headers` release. Point `Dawn_DIR` at a directory containing this `DawnConfig.cmake` next to both:
+
+```cmake
+set(_wasi_webgpu "${CMAKE_CURRENT_LIST_DIR}/wasi-webgpu-headers")
+add_library(dawn::webgpu_dawn INTERFACE IMPORTED)
+set_target_properties(dawn::webgpu_dawn PROPERTIES
+    INTERFACE_SOURCES "${_wasi_webgpu}/webgpu.c;${_wasi_webgpu}/async_futures.c;${_wasi_webgpu}/imports.c;${_wasi_webgpu}/imports_component_type.o"
+    INTERFACE_INCLUDE_DIRECTORIES "${CMAKE_CURRENT_LIST_DIR}/dawn-headers/include/webgpu_upstream;${_wasi_webgpu}")
+```
+
+Then build as above, but with `wasi-sdk-p3.cmake` / `--target=wasm32-wasip3`, `-DGGML_WEBGPU=ON -DDawn_DIR=...` at configure, and `build-wasi/ggml/src/ggml-webgpu/libggml-webgpu.a` added at link. See `.github/workflows/build-wasi.yml` for a more complete example.
+
 ## IBM Z & LinuxONE
 
 To read documentation for how to build on IBM Z & LinuxONE, [click here](./build-s390x.md)

ggml/cmake/common.cmake

@@ -2,10 +2,16 @@ function(ggml_get_flags CCID CCVER)
     set(C_FLAGS "")
     set(CXX_FLAGS "")
 
-    if (CCID MATCHES "Clang")
+    if (CCID MATCHES "Clang" OR CCID MATCHES "IntelLLVM")
         set(C_FLAGS   -Wunreachable-code-break -Wunreachable-code-return)
         set(CXX_FLAGS -Wunreachable-code-break -Wunreachable-code-return -Wmissing-prototypes -Wextra-semi)
 
+        if (CCID STREQUAL "IntelLLVM")
+            # for Intel DPC++ and LLVM SYCL, using debug does not disable optimization
+            list(APPEND C_FLAGS    -Rno-debug-disables-optimization)
+            list(APPEND CXX_FLAGS  -Rno-debug-disables-optimization)
+        endif()
+
         if (
             (CCID STREQUAL "Clang"      AND CCVER VERSION_GREATER_EQUAL 3.8.0) OR
             (CCID STREQUAL "AppleClang" AND CCVER VERSION_GREATER_EQUAL 7.3.0)

ggml/src/ggml-hexagon/CMakeLists.txt

@@ -21,6 +21,7 @@ message(STATUS "hexagon: using ${HEXAGON_SDK_ROOT} and ${HEXAGON_TOOLS_ROOT} for
 
 include(${HEXAGON_SDK_ROOT}/build/cmake/hexagon_fun.cmake)
 include(ExternalProject)
+find_package(Python3 REQUIRED COMPONENTS Interpreter)
 
 option(GGML_HEXAGON_HTP_DEBUG  "ggml-hexagon: enable HTP debug output" OFF)
 set(GGML_HEXAGON_HTP_BUILD_TYPE "Release" CACHE STRING "ggml-hexagon: HTP skel build type (Release, RelWithDebInfo, Debug)")
@@ -74,7 +75,8 @@ function(build_htp_skel V)
             -DHEXAGON_TOOLS_ROOT=${HEXAGON_TOOLS_ROOT}
             -DHEXAGON_HTP_DEBUG=${GGML_HEXAGON_HTP_DEBUG}
             -DDSP_VERSION=${V}
-            -DPREBUILT_LIB_DIR="toolv19_${V}")
+            -DPython3_EXECUTABLE=${Python3_EXECUTABLE}
+            -DPREBUILT_LIB_DIR="toolv21_${V}")
     list(APPEND HTP_SKELS ${CMAKE_CURRENT_BINARY_DIR}/libggml-htp-${V}.so)
     list(APPEND HTP_PROJECTS htp-${V})
     set(HTP_SKELS ${HTP_SKELS} PARENT_SCOPE)
@@ -85,6 +87,7 @@ build_htp_skel(v73)
 build_htp_skel(v75)
 build_htp_skel(v79)
 build_htp_skel(v81)
+build_htp_skel(v85)
 
 # Install Hexagon skels required at runtime
 install(FILES ${HTP_SKELS} TYPE LIB)

ggml/src/ggml-hexagon/ggml-hexagon.cpp

@@ -8399,9 +8399,9 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
             if (opt_arch < 73) {
                 GGML_LOG_WARN("ggml-hex: Hexagon arch v%d is under supported range, capping at v73\n", opt_arch);
                 opt_arch = 73;
-            } else if (opt_arch > 81) {
-                GGML_LOG_WARN("ggml-hex: Hexagon arch v%d is over supported range, capping at v81\n", opt_arch);
-                opt_arch = 81;
+            } else if (opt_arch > 85) {
+                GGML_LOG_WARN("ggml-hex: Hexagon arch v%d is over supported range, capping at v85\n", opt_arch);
+                opt_arch = 85;
             }
         }
     } else {

ggml/src/ggml-hexagon/htp/CMakeLists.txt

@@ -53,9 +53,17 @@ target_compile_definitions(${HTP_LIB} PRIVATE
     $<IF:$<BOOL:${HEXAGON_HTP_DEBUG}>,HTP_DEBUG=1,NDEBUG=1>
     $<IF:$<BOOL:${HEXAGON_HTP_DEBUG}>,FARF_HIGH=1,>)
 
-
 build_idl(htp_iface.idl ${HTP_LIB})
 
 set_target_properties(${HTP_LIB} PROPERTIES EXPORT_COMPILE_COMMANDS ON)
 
+if (NOT Python3_EXECUTABLE)
+    find_package(Python3 REQUIRED COMPONENTS Interpreter)
+endif()
+
+add_custom_command(TARGET ${HTP_LIB} POST_BUILD
+    COMMAND ${CMAKE_COMMAND} -E echo "Adding SWIV section $<TARGET_FILE_NAME:${HTP_LIB}>"
+    COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/htp-swiv.py -i $<TARGET_FILE:${HTP_LIB}> -o $<TARGET_FILE:${HTP_LIB}>
+    VERBATIM)
+
 install(TARGETS ${HTP_LIB})

ggml/src/ggml-hexagon/htp/htp-swiv.py

@@ -0,0 +1,512 @@
+#!/usr/bin/env python3
+
+import argparse
+import logging
+import os
+import shutil
+import struct
+import sys
+import tempfile
+import zlib
+
+ELF_MAGIC = b'\x7fELF'
+ELF_CLASS_INDEX = 4
+ELFCLASS32 = 1
+ELFCLASS64 = 2
+ELF_VERSION_INDEX = 6
+EV_CURRENT = 1
+MAX_PHDR_COUNT = 100
+
+LOAD_TYPE = 0x1
+
+SHT_NOBITS = 8
+
+SWIV_MAGIC = b'SWIV'
+SWIV_SEG_SIZE_MIN = 16
+# SHT_LOUSER plus ASCII "SWIV".
+SWIV_SECTION_TYPE = 0xD3574956
+
+logger = logging.getLogger("htp-swiv")
+
+
+class Elf_Ehdr_common:
+    s = struct.Struct('16sHHI')
+
+    def __init__(self, data):
+        unpacked_data = (Elf_Ehdr_common.s).unpack(data)
+        self.e_ident = unpacked_data[0]
+        self.e_type = unpacked_data[1]
+        self.e_machine = unpacked_data[2]
+        self.e_version = unpacked_data[3]
+
+
+class Elf32_Ehdr:
+    s = struct.Struct('16sHHIIIIIHHHHHH')
+
+    def __init__(self, data):
+        unpacked_data = (Elf32_Ehdr.s).unpack(data)
+        self.e_ident = unpacked_data[0]
+        self.e_type = unpacked_data[1]
+        self.e_machine = unpacked_data[2]
+        self.e_version = unpacked_data[3]
+        self.e_entry = unpacked_data[4]
+        self.e_phoff = unpacked_data[5]
+        self.e_shoff = unpacked_data[6]
+        self.e_flags = unpacked_data[7]
+        self.e_ehsize = unpacked_data[8]
+        self.e_phentsize = unpacked_data[9]
+        self.e_phnum = unpacked_data[10]
+        self.e_shentsize = unpacked_data[11]
+        self.e_shnum = unpacked_data[12]
+        self.e_shstrndx = unpacked_data[13]
+
+    def get_packed_data(self):
+        values = [self.e_ident,
+                  self.e_type,
+                  self.e_machine,
+                  self.e_version,
+                  self.e_entry,
+                  self.e_phoff,
+                  self.e_shoff,
+                  self.e_flags,
+                  self.e_ehsize,
+                  self.e_phentsize,
+                  self.e_phnum,
+                  self.e_shentsize,
+                  self.e_shnum,
+                  self.e_shstrndx
+                  ]
+
+        return (Elf32_Ehdr.s).pack(*values)
+
+
+class Elf32_Phdr:
+    s = struct.Struct('I' * 8)
+
+    def __init__(self, data):
+        unpacked_data = (Elf32_Phdr.s).unpack(data)
+        self.p_type = unpacked_data[0]
+        self.p_offset = unpacked_data[1]
+        self.p_vaddr = unpacked_data[2]
+        self.p_paddr = unpacked_data[3]
+        self.p_filesz = unpacked_data[4]
+        self.p_memsz = unpacked_data[5]
+        self.p_flags = unpacked_data[6]
+        self.p_align = unpacked_data[7]
+
+    def get_packed_data(self):
+        values = [self.p_type,
+                  self.p_offset,
+                  self.p_vaddr,
+                  self.p_paddr,
+                  self.p_filesz,
+                  self.p_memsz,
+                  self.p_flags,
+                  self.p_align
+                  ]
+
+        return (Elf32_Phdr.s).pack(*values)
+
+
+class Elf32_Shdr:
+    s = struct.Struct('I' * 10)
+
+    def __init__(self, data):
+        unpacked_data = (Elf32_Shdr.s).unpack(data)
+        self.sh_name = unpacked_data[0]
+        self.sh_type = unpacked_data[1]
+        self.sh_flags = unpacked_data[2]
+        self.sh_addr = unpacked_data[3]
+        self.sh_offset = unpacked_data[4]
+        self.sh_size = unpacked_data[5]
+        self.sh_link = unpacked_data[6]
+        self.sh_info = unpacked_data[7]
+        self.sh_addralign = unpacked_data[8]
+        self.sh_entsize = unpacked_data[9]
+        self.sh_name_str = ""
+
+    def get_packed_data(self):
+        values = [self.sh_name,
+                  self.sh_type,
+                  self.sh_flags,
+                  self.sh_addr,
+                  self.sh_offset,
+                  self.sh_size,
+                  self.sh_link,
+                  self.sh_info,
+                  self.sh_addralign,
+                  self.sh_entsize]
+
+        return (Elf32_Shdr.s).pack(*values)
+
+
+class Elf64_Ehdr:
+    s = struct.Struct('16sHHIQQQIHHHHHH')
+
+    def __init__(self, data):
+        unpacked_data = (Elf64_Ehdr.s).unpack(data)
+        self.e_ident = unpacked_data[0]
+        self.e_type = unpacked_data[1]
+        self.e_machine = unpacked_data[2]
+        self.e_version = unpacked_data[3]
+        self.e_entry = unpacked_data[4]
+        self.e_phoff = unpacked_data[5]
+        self.e_shoff = unpacked_data[6]
+        self.e_flags = unpacked_data[7]
+        self.e_ehsize = unpacked_data[8]
+        self.e_phentsize = unpacked_data[9]
+        self.e_phnum = unpacked_data[10]
+        self.e_shentsize = unpacked_data[11]
+        self.e_shnum = unpacked_data[12]
+        self.e_shstrndx = unpacked_data[13]
+
+    def get_packed_data(self):
+        values = [self.e_ident,
+                  self.e_type,
+                  self.e_machine,
+                  self.e_version,
+                  self.e_entry,
+                  self.e_phoff,
+                  self.e_shoff,
+                  self.e_flags,
+                  self.e_ehsize,
+                  self.e_phentsize,
+                  self.e_phnum,
+                  self.e_shentsize,
+                  self.e_shnum,
+                  self.e_shstrndx
+                  ]
+
+        return (Elf64_Ehdr.s).pack(*values)
+
+
+class Elf64_Phdr:
+    s = struct.Struct('IIQQQQQQ')
+
+    def __init__(self, data):
+        unpacked_data = (Elf64_Phdr.s).unpack(data)
+        self.p_type = unpacked_data[0]
+        self.p_flags = unpacked_data[1]
+        self.p_offset = unpacked_data[2]
+        self.p_vaddr = unpacked_data[3]
+        self.p_paddr = unpacked_data[4]
+        self.p_filesz = unpacked_data[5]
+        self.p_memsz = unpacked_data[6]
+        self.p_align = unpacked_data[7]
+
+    def get_packed_data(self):
+        values = [self.p_type,
+                  self.p_flags,
+                  self.p_offset,
+                  self.p_vaddr,
+                  self.p_paddr,
+                  self.p_filesz,
+                  self.p_memsz,
+                  self.p_align
+                  ]
+
+        return (Elf64_Phdr.s).pack(*values)
+
+
+class Elf64_Shdr:
+    s = struct.Struct('IIQQQQIIQQ')
+
+    def __init__(self, data):
+        unpacked_data = (Elf64_Shdr.s).unpack(data)
+        self.sh_name = unpacked_data[0]
+        self.sh_type = unpacked_data[1]
+        self.sh_flags = unpacked_data[2]
+        self.sh_addr = unpacked_data[3]
+        self.sh_offset = unpacked_data[4]
+        self.sh_size = unpacked_data[5]
+        self.sh_link = unpacked_data[6]
+        self.sh_info = unpacked_data[7]
+        self.sh_addralign = unpacked_data[8]
+        self.sh_entsize = unpacked_data[9]
+        self.sh_name_str = ""
+
+    def get_packed_data(self):
+        values = [self.sh_name,
+                  self.sh_type,
+                  self.sh_flags,
+                  self.sh_addr,
+                  self.sh_offset,
+                  self.sh_size,
+                  self.sh_link,
+                  self.sh_info,
+                  self.sh_addralign,
+                  self.sh_entsize]
+
+        return (Elf64_Shdr.s).pack(*values)
+
+
+class SWIV_Seg:
+    s = struct.Struct('4sIQ')
+
+    def __init__(self):
+        self.magic = SWIV_MAGIC
+        self.crc = 0
+        self.padding = 0
+
+    def get_packed_data(self):
+        values = [self.magic,
+                  self.crc,
+                  self.padding
+                  ]
+
+        return (SWIV_Seg.s).pack(*values)
+
+
+def is_elf64(elf_header):
+    return elf_header.e_ident[ELF_CLASS_INDEX] == ELFCLASS64
+
+
+def verify_elf_header(elf_header):
+    ident = elf_header.e_ident
+    return (ident.startswith(ELF_MAGIC)
+            and ident[ELF_CLASS_INDEX] in (ELFCLASS32, ELFCLASS64)
+            and ident[ELF_VERSION_INDEX] == EV_CURRENT)
+
+
+def file_copy_offset(in_fp, in_off, out_fp, out_off, num_bytes):
+    in_fp.seek(in_off)
+    read_in = in_fp.read(num_bytes)
+    out_fp.seek(out_off)
+    out_fp.write(read_in)
+
+    return num_bytes
+
+
+def preprocess_elf_file(elf_file_name):
+    with open(elf_file_name, 'rb') as elf_fp:
+        elf_header = Elf_Ehdr_common(elf_fp.read(Elf_Ehdr_common.s.size))
+
+        if verify_elf_header(elf_header) is False:
+            raise RuntimeError("Input ELF verification failed : " + elf_file_name)
+
+        elf_fp.seek(0)
+
+        if is_elf64(elf_header):
+            elf_header = Elf64_Ehdr(elf_fp.read(Elf64_Ehdr.s.size))
+        else:
+            elf_header = Elf32_Ehdr(elf_fp.read(Elf32_Ehdr.s.size))
+
+        phdr_table = []
+
+        phdr_size = elf_header.e_phentsize
+
+        file_offset = elf_header.e_phoff
+        elf_fp.seek(file_offset)
+
+        for _ in range(elf_header.e_phnum):
+            if is_elf64(elf_header):
+                phdr_table.append(Elf64_Phdr(elf_fp.read(phdr_size)))
+            else:
+                phdr_table.append(Elf32_Phdr(elf_fp.read(phdr_size)))
+
+    return [elf_header, phdr_table]
+
+
+def process_section_header(elf_file_name):
+    with open(elf_file_name, 'rb') as elf_fp:
+        elf_header = Elf_Ehdr_common(elf_fp.read(Elf_Ehdr_common.s.size))
+
+        if verify_elf_header(elf_header) is False:
+            raise RuntimeError("Input ELF verification failed : " + elf_file_name)
+
+        elf_fp.seek(0)
+
+        if is_elf64(elf_header):
+            elf_header = Elf64_Ehdr(elf_fp.read(Elf64_Ehdr.s.size))
+        else:
+            elf_header = Elf32_Ehdr(elf_fp.read(Elf32_Ehdr.s.size))
+
+        shdr_table = []
+
+        shdr_size = elf_header.e_shentsize
+        elf_fp.seek(elf_header.e_shoff)
+
+        for _ in range(elf_header.e_shnum):
+            if is_elf64(elf_header):
+                shdr_table.append(Elf64_Shdr(elf_fp.read(shdr_size)))
+            else:
+                shdr_table.append(Elf32_Shdr(elf_fp.read(shdr_size)))
+
+    return shdr_table
+
+
+def crc32_update_file(elf_fp, offset, size, crc):
+    elf_fp.seek(offset)
+    remaining = size
+    while remaining > 0:
+        chunk = elf_fp.read(min(1024 * 1024, remaining))
+        if not chunk:
+            raise RuntimeError("Unexpected EOF while calculating CRC")
+        crc = zlib.crc32(chunk, crc)
+        remaining -= len(chunk)
+    return crc
+
+
+def crc32_update_zeros(size, crc):
+    zero_chunk = b'\x00' * min(1024 * 1024, size)
+    remaining = size
+    while remaining > 0:
+        chunk = zero_chunk[:min(len(zero_chunk), remaining)]
+        crc = zlib.crc32(chunk, crc)
+        remaining -= len(chunk)
+    return crc
+
+
+def crc32_calculate(elf_fp, phdr_table):
+    crc = 0
+
+    # SWIV validates loadable segment contents, including BSS as zeros.
+    for curr_phdr in phdr_table:
+        if curr_phdr.p_type != LOAD_TYPE:
+            continue
+
+        crc = crc32_update_file(elf_fp, curr_phdr.p_offset, curr_phdr.p_filesz, crc)
+
+        zero_size = curr_phdr.p_memsz - curr_phdr.p_filesz
+        if zero_size > 0:
+            crc = crc32_update_zeros(zero_size, crc)
+
+    crc &= 0xffffffff
+    logger.info("calculated CRC: %#x", crc)
+
+    return crc
+
+
+def verify_swiv_seg(elf_in_fp, offset):
+
+    elf_in_fp.seek(offset)
+    swiv_seg = SWIV_Seg()
+    unpacked_data = (SWIV_Seg.s).unpack(elf_in_fp.read(SWIV_SEG_SIZE_MIN))
+    swiv_seg.magic = unpacked_data[0]
+
+    return swiv_seg.magic == SWIV_MAGIC
+
+
+def same_output_file(input_elf, output_elf):
+    try:
+        return os.path.exists(output_elf) and os.path.samefile(input_elf, output_elf)
+    except FileNotFoundError:
+        return False
+
+
+def make_output_target(input_elf, output_elf):
+    if not same_output_file(input_elf, output_elf):
+        return output_elf, None
+
+    output_dir = os.path.dirname(os.path.abspath(output_elf))
+    output_base = os.path.basename(output_elf)
+    fd, temp_output = tempfile.mkstemp(prefix=output_base + ".", suffix=".tmp", dir=output_dir)
+    os.close(fd)
+    shutil.copymode(input_elf, temp_output)
+    return temp_output, temp_output
+
+
+def swiv_generate(input_elf, output_elf):
+    output_target, temp_output = make_output_target(input_elf, output_elf)
+
+    try:
+        with open(input_elf, "rb") as elf_in_fp, open(output_target, "wb+") as elf_out_fp:
+            [elf_header, phdr_table] = preprocess_elf_file(input_elf)
+
+            num_phdrs = elf_header.e_phnum
+            num_shdrs = elf_header.e_shnum
+            phdr_size = elf_header.e_phentsize
+
+            if num_phdrs == 0 or num_phdrs > MAX_PHDR_COUNT:
+                raise RuntimeError("Input ELF has invalid number of program headers.")
+
+            if elf_header.e_shoff == 0 or num_shdrs == 0:
+                raise RuntimeError("Input ELF does not contain a section header table.")
+
+            shdr_table_input = process_section_header(input_elf)
+            shdr_table_output = process_section_header(input_elf)
+
+            last_phdr = phdr_table[num_phdrs - 1]
+            if last_phdr.p_memsz == SWIV_SEG_SIZE_MIN and verify_swiv_seg(elf_in_fp, last_phdr.p_offset):
+                raise RuntimeError("A SWIV segment is detected in the input ELF.")
+
+            last_shdr = shdr_table_input[elf_header.e_shnum - 1]
+            if last_shdr.sh_size == SWIV_SEG_SIZE_MIN and verify_swiv_seg(elf_in_fp, last_shdr.sh_offset):
+                raise RuntimeError("A SWIV section is detected in the input ELF.")
+
+            if is_elf64(elf_header):
+                swiv_shdr = Elf64_Shdr(b'\0' * Elf64_Shdr.s.size)
+            else:
+                swiv_shdr = Elf32_Shdr(b'\0' * Elf32_Shdr.s.size)
+
+            swiv_shdr.sh_name = 0
+            swiv_shdr.sh_type = SWIV_SECTION_TYPE
+            swiv_shdr.sh_flags = 0
+            swiv_shdr.sh_addr = 0
+            swiv_shdr.sh_size = SWIV_SEG_SIZE_MIN
+            swiv_shdr.sh_offset = elf_header.e_shoff + (elf_header.e_shnum + 1) * elf_header.e_shentsize
+            shdr_table_output.append(swiv_shdr)
+            elf_header.e_shnum += 1
+
+            if elf_header.e_phoff != 0:
+                for phdr in phdr_table:
+                    file_copy_offset(elf_in_fp, phdr.p_offset, elf_out_fp, phdr.p_offset, phdr.p_filesz)
+
+            elf_out_fp.seek(0)
+            elf_out_fp.write(elf_header.get_packed_data())
+
+            phdr_start = elf_header.e_phoff
+            for phdr in phdr_table:
+                elf_out_fp.seek(phdr_start)
+                elf_out_fp.write(phdr.get_packed_data())
+                phdr_start += phdr_size
+
+            for shdr in shdr_table_input[:num_shdrs]:
+                if shdr.sh_type == SHT_NOBITS or shdr.sh_size == 0:
+                    continue
+                file_copy_offset(elf_in_fp, shdr.sh_offset, elf_out_fp, shdr.sh_offset, shdr.sh_size)
+
+            shdr_start = elf_header.e_shoff
+            for shdr in shdr_table_output:
+                elf_out_fp.seek(shdr_start)
+                elf_out_fp.write(shdr.get_packed_data())
+                shdr_start += elf_header.e_shentsize
+
+            elf_out_fp.flush()
+
+            swiv_seg = SWIV_Seg()
+            swiv_seg.crc = crc32_calculate(elf_out_fp, phdr_table)
+            elf_out_fp.seek(swiv_shdr.sh_offset)
+            elf_out_fp.write(swiv_seg.get_packed_data())
+
+        if temp_output is not None:
+            os.replace(temp_output, output_elf)
+            temp_output = None
+
+        return 0
+    finally:
+        if temp_output is not None and os.path.exists(temp_output):
+            os.unlink(temp_output)
+
+
+def main(argv):
+    parser = argparse.ArgumentParser(description="Add a Software Image Validation (SWIV) section to an HTP ELF file.")
+    parser.add_argument("-i", "--input", dest="input_file", metavar="INPUT", required=True)
+    parser.add_argument("-o", "--output", dest="output_file", metavar="OUTPUT", required=True)
+    parser.add_argument("-v", "--verbose", action="store_true", help="show calculated CRC")
+    args = parser.parse_args(argv)
+
+    logging.basicConfig(
+        level=logging.INFO if args.verbose else logging.WARNING,
+        format="%(message)s",
+    )
+
+    try:
+        return swiv_generate(args.input_file, args.output_file)
+    except RuntimeError as exc:
+        logger.error("ERROR: %s", exc)
+        return 2
+
+
+if __name__ == "__main__":
+    sys.exit(main(sys.argv[1:]))

ggml/src/ggml-hexagon/libggml-htp.inf

@@ -18,6 +18,7 @@ libggml-htp-v73.so = 1
 libggml-htp-v75.so = 1
 libggml-htp-v79.so = 1
 libggml-htp-v81.so = 1
+libggml-htp-v85.so = 1
 
 [ControlFlags]
 ExcludeFromSelect = *
@@ -30,6 +31,7 @@ libggml-htp-v73.so,,,0x10 ;COPYFLG_NO_OVERWRITE
 libggml-htp-v75.so,,,0x10 ;COPYFLG_NO_OVERWRITE
 libggml-htp-v79.so,,,0x10 ;COPYFLG_NO_OVERWRITE
 libggml-htp-v81.so,,,0x10 ;COPYFLG_NO_OVERWRITE
+libggml-htp-v85.so,,,0x10 ;COPYFLG_NO_OVERWRITE
 
 [Strings]
 GGML   = 'GGML'

ggml/src/ggml-opencl/CMakeLists.txt

@@ -140,6 +140,9 @@ set(GGML_OPENCL_KERNELS
     gemm_noshuffle_q5_0_q8_1_dp4a
     gemm_noshuffle_iq4_nl_q8_1_dp4a
     gemm_noshuffle_q4_0_q8_1_dp4a
+    gemm_coopk_q4_0_q8_1_dp4a
+    gemm_coopk_q4_k_q8_1_dp4a
+    gemm_coopk_q6_k_q8_1_dp4a
     quant_a_q8_1
     gemv_moe_q4_k_f32_ns
     gemm_moe_q5_k_f32_ns

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

@@ -868,7 +868,7 @@ struct ggml_backend_opencl_context {
 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
     ggml_cl_adreno_xmem_attn_state adreno_xmem_attn;
 #endif
-    cl_kernel kernel_get_rows_f32, kernel_get_rows_f16, kernel_get_rows_q4_0;
+    cl_kernel kernel_get_rows_f32, kernel_get_rows_f16, kernel_get_rows_i32, kernel_get_rows_q4_0;
     cl_kernel kernel_set_rows_f32_i64, kernel_set_rows_f32_i32, kernel_set_rows_f16_i64, kernel_set_rows_f16_i32;
     cl_kernel kernel_set_rows_q8_0_i64, kernel_set_rows_q8_0_i32;
     cl_kernel kernel_set_rows_q8_0_soa_i64, kernel_set_rows_q8_0_soa_i32;
@@ -876,7 +876,7 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_set_rows_q4_0_soa_i64, kernel_set_rows_q4_0_soa_i32;
     cl_kernel kernel_rope_norm_f32, kernel_rope_norm_f16, kernel_rope_neox_f32, kernel_rope_neox_f16;
     cl_kernel kernel_rope_multi_f32, kernel_rope_multi_f16, kernel_rope_vision_f32, kernel_rope_vision_f16;
-    cl_kernel kernel_cpy_f16_f16, kernel_cpy_f16_f32, kernel_cpy_f32_f16, kernel_cpy_f32_f32, kernel_cpy_i32_i32;
+    cl_kernel kernel_cpy_f16_f16, kernel_cpy_f16_f32, kernel_cpy_f32_f16, kernel_cpy_f32_f32, kernel_cpy_i32_i32, kernel_cpy_i32_f32, kernel_cpy_f32_i32;
     cl_kernel kernel_cpy_f32_f32_pack = nullptr;
     cl_kernel kernel_cpy_f32_f32_flat = nullptr;
     cl_kernel kernel_mul_mat_f32_f32;
@@ -1296,6 +1296,22 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_gemm_noshuffle_iq4_nl_f32;
     cl_kernel kernel_gemm_noshuffle_iq4_nl_q8_1_dp4a = nullptr;  // dp4a (int8) dense IQ4_NL prefill GEMM
     cl_kernel kernel_gemm_noshuffle_q4_0_q8_1_dp4a = nullptr;  // dp4a (int8) dense q4_0 prefill GEMM
+    cl_kernel kernel_gemm_coopk_q4_0_q8_1_dp4a_c4 = nullptr;
+    cl_kernel kernel_gemm_coopk_q4_0_q8_1_dp4a_c2 = nullptr;
+    int       nsg_coopk_dp4a_q4_0  = 8;
+    int       rows_coopk_dp4a_q4_0 = 4;
+    cl_kernel kernel_gemm_coopk_q4_k_q8_1_dp4a_c4 = nullptr;
+    cl_kernel kernel_gemm_coopk_q4_k_q8_1_dp4a_c2 = nullptr;
+    int       nsg_coopk_dp4a_q4_k  = 8;
+    int       rows_coopk_dp4a_q4_k = 4;
+    cl_kernel kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c4 = nullptr;
+    cl_kernel kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c2 = nullptr;
+    int       nsg_coopk_dp4a_q4_k_bin_c4 = 8;
+    int       nsg_coopk_dp4a_q4_k_bin_c2 = 8;
+    cl_kernel kernel_gemm_coopk_q6_k_q8_1_dp4a_c4 = nullptr;
+    cl_kernel kernel_gemm_coopk_q6_k_q8_1_dp4a_c2 = nullptr;
+    int       nsg_coopk_dp4a_q6_k  = 8;
+    int       rows_coopk_dp4a_q6_k = 4;
 #endif // GGML_OPENCL_USE_ADRENO_KERNELS
 
     void free() {
@@ -1492,6 +1508,180 @@ static bool use_adreno_bin_kernels(ggml_backend_opencl_context * backend_ctx) {
 #endif // GGML_OPENCL_USE_ADRENO_BIN_KERNELS
 }
 
+#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
+static bool ggml_cl_coopk_dp4a_enabled(const ggml_backend_opencl_context * backend_ctx) {
+    static const char * const env = getenv("GGML_OPENCL_COOPK_DP4A");
+    if (env && *env && atoi(env) == 0) {
+        return false;
+    }
+    if (!backend_ctx->has_integer_dot) {
+        return false;
+    }
+    if (backend_ctx->gpu_family == GPU_FAMILY::ADRENO) {
+        if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::ADRENO_UNKNOWN ||
+            backend_ctx->adreno_gen == ADRENO_GPU_GEN::A6X ||
+            backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X) {
+            return false;
+        }
+        if (backend_ctx->adreno_cl_compiler_version.type == ADRENO_CL_COMPILER_TYPE::E17) {
+            return false;
+        }
+    }
+    return true;
+}
+
+static bool ggml_cl_coopk_dp4a_narrow_on(const ggml_backend_opencl_context * backend_ctx,
+                                         const ggml_tensor * op, int rows) {
+    if (!ggml_cl_coopk_dp4a_enabled(backend_ctx)) {
+        return false;
+    }
+    // Each lane handles rows output rows; subgroups split K for the same rows.
+    // Full row groups and 32-value K blocks avoid tail checks in the kernels.
+    return op->ne[1] >= 2 && op->ne[1] <= 4
+        && rows > 0 && (op->src[0]->ne[1] % (64 * rows)) == 0
+        && (op->src[0]->ne[0] % 32) == 0;
+}
+
+// COOPK_NSG sizes the local reduction buffer, so a smaller workgroup requires a rebuild.
+static cl_program ggml_cl_build_program_coopk(ggml_backend_opencl_context * backend_ctx,
+                                              const char * kernel_src, const char * kernel_name,
+                                              const std::string & opts_base, int & nsg) {
+    while (nsg >= 2) {
+        const std::string opts = opts_base + " -DCOOPK_NSG=" + std::to_string(nsg);
+        cl_program program = build_program_from_source(backend_ctx, kernel_src, opts);
+        cl_int err;
+        cl_kernel kernel = clCreateKernel(program, kernel_name, &err);
+        CL_CHECK(err);
+        size_t max_workgroup_size = 0;
+        CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE,
+                                        sizeof(max_workgroup_size), &max_workgroup_size, NULL));
+        CL_CHECK(clReleaseKernel(kernel));
+        if ((size_t) 64 * nsg <= max_workgroup_size) {
+            return program;
+        }
+        CL_CHECK(clReleaseProgram(program));
+        nsg /= 2;
+    }
+    return nullptr;
+}
+
+static void ggml_cl_coopk_build_narrow(ggml_backend_opencl_context * backend_ctx,
+                                       const std::string & kernel_src, const char * kernel_name,
+                                       int rows, int & nsg, cl_kernel & kernel_c4, cl_kernel & kernel_c2) {
+    const std::string base_opts = backend_ctx->kernel_compile_opts
+        + " -DCOOPK_ROWS=" + std::to_string(rows);
+    for (int cols = 4; cols >= 2; cols -= 2) {
+        int nsg_eff = nsg;
+        cl_program program = ggml_cl_build_program_coopk(
+            backend_ctx, kernel_src.c_str(), kernel_name,
+            base_opts + " -DCOOPK_COLS=" + std::to_string(cols), nsg_eff);
+        if (program == nullptr) {
+            return;
+        }
+        cl_int err;
+        cl_kernel kernel = clCreateKernel(program, kernel_name, &err);
+        if (err != CL_SUCCESS) { kernel = nullptr; }
+        CL_CHECK(clReleaseProgram(program));
+        if (cols == 4) {
+            kernel_c4 = kernel;
+            nsg  = nsg_eff;
+            if (kernel == nullptr) {
+                return;
+            }
+        } else {
+            // Both widths share one stored subgroup count for launch.
+            if (kernel != nullptr && nsg_eff != nsg) {
+                CL_CHECK(clReleaseKernel(kernel));
+                kernel = nullptr;
+            }
+            kernel_c2 = kernel;
+        }
+    }
+}
+
+static void load_cl_kernels_gemm_coopk(ggml_backend_opencl_context * backend_ctx) {
+    if (backend_ctx->gpu_family != GPU_FAMILY::ADRENO || !ggml_cl_coopk_dp4a_enabled(backend_ctx)) {
+        return;
+    }
+
+    // Q4_0
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "gemm_coopk_q4_0_q8_1_dp4a.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("gemm_coopk_q4_0_q8_1_dp4a.cl");
+#endif
+        ggml_cl_coopk_build_narrow(backend_ctx, kernel_src, "kernel_gemm_coopk_q4_0_q8_1_dp4a",
+                                 backend_ctx->rows_coopk_dp4a_q4_0, backend_ctx->nsg_coopk_dp4a_q4_0,
+                                 backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c4,
+                                 backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c2);
+        GGML_LOG_CONT(".");
+    }
+
+    // Q4_K
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "gemm_coopk_q4_k_q8_1_dp4a.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("gemm_coopk_q4_k_q8_1_dp4a.cl");
+#endif
+        ggml_cl_coopk_build_narrow(backend_ctx, kernel_src, "kernel_gemm_coopk_q4_k_q8_1_dp4a",
+                                 backend_ctx->rows_coopk_dp4a_q4_k, backend_ctx->nsg_coopk_dp4a_q4_k,
+                                 backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c4,
+                                 backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c2);
+
+        const char * env = getenv("GGML_OPENCL_COOPK_Q4_K_BIN");
+        if (use_adreno_bin_kernels(backend_ctx) && (!env || !*env || atoi(env) != 0)) {
+            const std::string base_opts = backend_ctx->kernel_compile_opts
+                + " -DCOOPK_Q_BIN=1 -DCOOPK_ROWS=" + std::to_string(backend_ctx->rows_coopk_dp4a_q4_k);
+            for (int cols = 4; cols >= 2; cols -= 2) {
+                int & nsg = (cols == 4) ? backend_ctx->nsg_coopk_dp4a_q4_k_bin_c4
+                                        : backend_ctx->nsg_coopk_dp4a_q4_k_bin_c2;
+                int nsg_eff = nsg;
+                cl_program program = ggml_cl_build_program_coopk(
+                    backend_ctx, kernel_src.c_str(), "kernel_gemm_coopk_q4_k_q8_1_dp4a",
+                    base_opts + " -DCOOPK_COLS=" + std::to_string(cols), nsg_eff);
+                if (program == nullptr) {
+                    continue;
+                }
+                cl_int err;
+                cl_kernel kernel = clCreateKernel(program, "kernel_gemm_coopk_q4_k_q8_1_dp4a", &err);
+                if (err != CL_SUCCESS) { kernel = nullptr; }
+                CL_CHECK(clReleaseProgram(program));
+                if (cols == 4) {
+                    backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c4 = kernel;
+                } else {
+                    backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c2 = kernel;
+                }
+                nsg = nsg_eff;
+            }
+        }
+        GGML_LOG_CONT(".");
+    }
+
+    // Q6_K
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "gemm_coopk_q6_k_q8_1_dp4a.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("gemm_coopk_q6_k_q8_1_dp4a.cl");
+#endif
+        ggml_cl_coopk_build_narrow(backend_ctx, kernel_src, "kernel_gemm_coopk_q6_k_q8_1_dp4a",
+                                 backend_ctx->rows_coopk_dp4a_q6_k, backend_ctx->nsg_coopk_dp4a_q6_k,
+                                 backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c4,
+                                 backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c2);
+        GGML_LOG_CONT(".");
+    }
+}
+
+#endif
+
 static std::string ggml_opencl_make_compile_opts(const ggml_backend_opencl_context * backend_ctx) {
     auto opencl_c_std =
         std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor);
@@ -1652,6 +1842,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
             }
         }
         CL_CHECK((backend_ctx->kernel_cpy_i32_i32 = clCreateKernel(prog, "kernel_cpy_i32_i32", &err), err));
+        CL_CHECK((backend_ctx->kernel_cpy_i32_f32 = clCreateKernel(prog, "kernel_cpy_i32_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_cpy_f32_i32 = clCreateKernel(prog, "kernel_cpy_f32_i32", &err), err));
         GGML_LOG_CONT(".");
     }
 
@@ -1840,6 +2032,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
             build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
 
         CL_CHECK((backend_ctx->kernel_get_rows_f32  = clCreateKernel(backend_ctx->program_get_rows, "kernel_get_rows_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_get_rows_i32  = clCreateKernel(backend_ctx->program_get_rows, "kernel_get_rows_i32", &err), err));
         CL_CHECK((backend_ctx->kernel_get_rows_f16  = clCreateKernel(backend_ctx->program_get_rows, "kernel_get_rows_f16", &err), err));
         CL_CHECK((backend_ctx->kernel_get_rows_q4_0 = clCreateKernel(backend_ctx->program_get_rows, "kernel_get_rows_q4_0", &err), err));
         GGML_LOG_CONT(".");
@@ -5437,6 +5630,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
             }
         }
     }
+    load_cl_kernels_gemm_coopk(backend_ctx);
 #endif // GGML_OPENCL_USE_ADRENO_KERNELS
     GGML_LOG_CONT("\n");
     backend_ctx->kernels_loaded = true;
@@ -9102,6 +9296,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             return true;
         case GGML_OP_GET_ROWS:
             switch (op->src[0]->type) {
+                case GGML_TYPE_I32:
+                    return op->type == GGML_TYPE_I32;
                 case GGML_TYPE_F32:
                 case GGML_TYPE_F16:
                     return true;
@@ -9139,6 +9335,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             switch (op->src[0]->type) {
                 case GGML_TYPE_F32:
                     switch (op->type) {
+                        case GGML_TYPE_I32:
                         case GGML_TYPE_F16:
                         case GGML_TYPE_F32:
                             return true;
@@ -9155,6 +9352,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
                     }
                 case GGML_TYPE_I32:
                     switch (op->type) {
+                        case GGML_TYPE_F32:
                         case GGML_TYPE_I32:
                             return true;
                         default:
@@ -13619,6 +13817,9 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c
         case GGML_TYPE_F32:
             kernel = backend_ctx->kernel_get_rows_f32;
             break;
+        case GGML_TYPE_I32:
+            kernel = backend_ctx->kernel_get_rows_i32;
+            break;
         case GGML_TYPE_F16:
             kernel = backend_ctx->kernel_get_rows_f16;
             break;
@@ -13654,7 +13855,7 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c
     }
 
     int nchunks = 1;
-    if (src0->type == GGML_TYPE_F32) {
+    if (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_I32) {
         const int chunk_target = nth * 4;
         nchunks = (ne00 + chunk_target - 1) / chunk_target;
         nchunks = MAX(1, MIN(nchunks, 64));
@@ -20042,6 +20243,7 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
                 warned = true;
             }
         }
+        // This bin layout is not supported by CoopK.
         ggml_cl_mul_mat_q4_0_f32_adreno_ila(backend, src0, src1, dst);
         return;
     }
@@ -20129,6 +20331,71 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
         CL_CHECK(clReleaseMemObject(b_sub_buf));
         CL_CHECK(clReleaseMemObject(b_img));
     } else {
+        // CoopK uses the same untiled weight layout as gemm_noshuffle_q4_0_f32, with K-major Q8_1 activations (K contiguous).
+        if (ggml_cl_coopk_dp4a_narrow_on(backend_ctx, dst, backend_ctx->rows_coopk_dp4a_q4_0)
+            && backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c4 != nullptr
+            && ((N <= 2) ? backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c2
+                         : backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c4) != nullptr) {
+            cl_mem b_sub_buf = nullptr;
+            region.origin = offset1;
+            region.size   = (size_t)K * N * sizeof(float);
+            CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
+
+            const int    cols            = (N <= 2) ? 2 : 4;
+            const size_t n_blocks_padded = (size_t)cols * (K / 32);
+            const size_t size_qa         = (size_t)cols * K * sizeof(cl_char);
+            const size_t size_qa_prev    = backend_ctx->prealloc_moe_qa.size;
+            const size_t size_da_prev    = backend_ctx->prealloc_moe_da.size;
+            backend_ctx->prealloc_moe_qa.allocate(context, size_qa);
+            backend_ctx->prealloc_moe_da.allocate(context, n_blocks_padded * sizeof(cl_half));
+            backend_ctx->prealloc_moe_sa.allocate(context, n_blocks_padded * sizeof(cl_half));
+            if (backend_ctx->prealloc_moe_qa.size != size_qa_prev) {
+                const cl_char z8 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_qa.buffer,
+                                             &z8, sizeof(z8), 0, backend_ctx->prealloc_moe_qa.size, 0, NULL, NULL));
+            }
+            if (backend_ctx->prealloc_moe_da.size != size_da_prev) {
+                const cl_half z16 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_da.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_da.size, 0, NULL, NULL));
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_sa.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_sa.size, 0, NULL, NULL));
+            }
+
+            cl_int n_blocks = (cl_int)((size_t)N * (K / 32));
+            cl_kernel kernel_quant = backend_ctx->kernel_quant_a_q8_1;
+            CL_CHECK(clSetKernelArg(kernel_quant, 0, sizeof(cl_mem), &b_sub_buf));
+            CL_CHECK(clSetKernelArg(kernel_quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 4, sizeof(cl_int), &n_blocks));
+            size_t local_work_size_quant[1]  = { 64 };
+            size_t global_work_size_quant[1] = { (size_t)((((size_t)n_blocks + 63) / 64) * 64) };
+            backend_ctx->enqueue_ndrange_kernel(kernel_quant, 1, global_work_size_quant, local_work_size_quant, dst);
+
+            cl_kernel kernel_coopk = (N <= 2) ? backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c2
+                                              : backend_ctx->kernel_gemm_coopk_q4_0_q8_1_dp4a_c4;
+            const cl_int n_padded = cols;
+            CL_CHECK(clSetKernelArg(kernel_coopk, 0, sizeof(cl_mem),   &extra0_q4_0->q));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 1, sizeof(cl_mem),   &extra0_q4_0->d));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 2, sizeof(cl_mem),   &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 3, sizeof(cl_mem),   &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 4, sizeof(cl_mem),   &extrad->data_device));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 5, sizeof(cl_ulong), &offsetd));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 6, sizeof(cl_int),   &M));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 7, sizeof(cl_int),   &n_padded));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 8, sizeof(cl_int),   &K));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 9, sizeof(cl_int),   &N));
+
+            const size_t nsg = (size_t)backend_ctx->nsg_coopk_dp4a_q4_0;
+            size_t local_work_size[3]  = { 64, nsg, 1 };
+            size_t global_work_size[3] = { (size_t)(M / backend_ctx->rows_coopk_dp4a_q4_0), nsg, 1 };
+            backend_ctx->enqueue_ndrange_kernel(kernel_coopk, 3, global_work_size, local_work_size, dst);
+
+            CL_CHECK(clReleaseMemObject(b_sub_buf));
+            return;
+        }
+
         // dp4a (int8) dense prefill GEMM, default off
         static const char * q4_0_dense_dp4a_env = getenv("GGML_OPENCL_Q4_0_DENSE_DP4A");
         bool q4_0_dense_dp4a_on = q4_0_dense_dp4a_env
@@ -21701,6 +21968,86 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno_ila(ggml_backend_t backend, const gg
     int N = ne1;
     int K = ne00;
 
+    // CoopK handles with _bin use the bin weight layout (uint-transposed q), with K-major Q8_1 activations (K contiguous).
+    // This switch disables only this bin-layout path; weight conversion is unchanged.
+    static const char * env_coopk_bin = getenv("GGML_OPENCL_COOPK_Q4_K_BIN");
+    const bool use_coopk_bin = !env_coopk_bin || !*env_coopk_bin || atoi(env_coopk_bin) != 0;
+    if (use_coopk_bin && ne1 >= 2 && ne1 <= 4 &&
+        ggml_cl_coopk_dp4a_narrow_on(backend_ctx, dst, backend_ctx->rows_coopk_dp4a_q4_k) &&
+        backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c4 != nullptr &&
+        ((ne1 <= 2) ? backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c2
+                    : backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c4) != nullptr) {
+        cl_mem b_sub_buf = nullptr;
+        region.origin = offset1;
+        region.size   = (size_t)K * N * sizeof(float);
+        CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
+
+        // Allocate all kernel columns; padded columns have independent accumulators and are not stored.
+        const int    cols            = (ne1 <= 2) ? 2 : 4;
+        const size_t size_qa         = (size_t)cols * K * sizeof(cl_char);
+        const size_t n_blocks_padded = (size_t)cols * (K / 32);
+        const size_t size_qa_prev    = backend_ctx->prealloc_moe_qa.size;
+        const size_t size_da_prev    = backend_ctx->prealloc_moe_da.size;
+        backend_ctx->prealloc_moe_qa.allocate(context, size_qa);
+        backend_ctx->prealloc_moe_da.allocate(context, n_blocks_padded * sizeof(cl_half));
+        backend_ctx->prealloc_moe_sa.allocate(context, n_blocks_padded * sizeof(cl_half));
+        if (backend_ctx->prealloc_moe_qa.size != size_qa_prev) {
+            const cl_char z8 = 0;
+            CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_qa.buffer,
+                                         &z8, sizeof(z8), 0, backend_ctx->prealloc_moe_qa.size, 0, NULL, NULL));
+        }
+        if (backend_ctx->prealloc_moe_da.size != size_da_prev) {
+            const cl_half z16 = 0;
+            CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_da.buffer,
+                                         &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_da.size, 0, NULL, NULL));
+            CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_sa.buffer,
+                                         &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_sa.size, 0, NULL, NULL));
+        }
+
+        cl_int n_blocks = (cl_int)((size_t)N * (K / 32));
+        cl_kernel kernel_quant = backend_ctx->kernel_quant_a_q8_1;
+        CL_CHECK(clSetKernelArg(kernel_quant, 0, sizeof(cl_mem), &b_sub_buf));
+        CL_CHECK(clSetKernelArg(kernel_quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
+        CL_CHECK(clSetKernelArg(kernel_quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
+        CL_CHECK(clSetKernelArg(kernel_quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
+        CL_CHECK(clSetKernelArg(kernel_quant, 4, sizeof(cl_int), &n_blocks));
+        size_t local_work_size_quant[1]  = { 64 };
+        size_t global_work_size_quant[1] = { (size_t)((((size_t)n_blocks + 63) / 64) * 64) };
+        backend_ctx->enqueue_ndrange_kernel(kernel_quant, 1, global_work_size_quant, local_work_size_quant, dst);
+
+        cl_kernel kernel_coopk = (ne1 <= 2) ? backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c2
+                                            : backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_bin_c4;
+        const int nsg_selected = (ne1 <= 2) ? backend_ctx->nsg_coopk_dp4a_q4_k_bin_c2
+                                            : backend_ctx->nsg_coopk_dp4a_q4_k_bin_c4;
+        const cl_uchar mask_d6 = 0x3F, mask_d4 = 0x0F, mask_hi2 = 0xC0;
+        const cl_int   n_padded = cols;
+        CL_CHECK(clSetKernelArg(kernel_coopk,  0, sizeof(cl_mem),   &extra0_q4_k->q));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  1, sizeof(cl_mem),   &extra0_q4_k->s));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  2, sizeof(cl_mem),   &extra0_q4_k->d));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  3, sizeof(cl_mem),   &extra0_q4_k->dm));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  4, sizeof(cl_mem),   &backend_ctx->prealloc_moe_qa.buffer));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  5, sizeof(cl_mem),   &backend_ctx->prealloc_moe_da.buffer));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  6, sizeof(cl_mem),   &backend_ctx->prealloc_moe_sa.buffer));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  7, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  8, sizeof(cl_ulong), &offsetd));
+        CL_CHECK(clSetKernelArg(kernel_coopk,  9, sizeof(cl_int),   &ne01));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 10, sizeof(cl_int),   &n_padded));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 11, sizeof(cl_int),   &ne00));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 12, sizeof(cl_int),   &ne1));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 13, sizeof(cl_uchar), &mask_d6));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 14, sizeof(cl_uchar), &mask_d4));
+        CL_CHECK(clSetKernelArg(kernel_coopk, 15, sizeof(cl_uchar), &mask_hi2));
+
+        // Each width can compile with a different subgroup count.
+        const size_t nsg = (size_t)nsg_selected;
+        size_t local_work_size[3]  = { 64, nsg, 1 };
+        size_t global_work_size[3] = { (size_t)(ne01 / backend_ctx->rows_coopk_dp4a_q4_k), nsg, 1 };
+        backend_ctx->enqueue_ndrange_kernel(kernel_coopk, 3, global_work_size, local_work_size, dst);
+
+        CL_CHECK(clReleaseMemObject(b_sub_buf));
+        return;
+    }
+
     if (ne1 == 1) {
         cl_mem b_sub_buf = nullptr;
         cl_mem b_img     = nullptr;
@@ -22148,6 +22495,74 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
         region.size = K * N * sizeof(float);
         CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
 
+        // CoopK handles without _bin use the same untiled weight layout as gemm_noshuffle_q4_k_f32.
+        // Activations use K-major Q8_1 SoA (K contiguous).
+        if (ggml_cl_coopk_dp4a_narrow_on(backend_ctx, dst, backend_ctx->rows_coopk_dp4a_q4_k)
+            && backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c4 != nullptr
+            && ((N <= 2) ? backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c2
+                         : backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c4) != nullptr) {
+            const int    cols            = (N <= 2) ? 2 : 4;
+            const size_t size_qa         = (size_t)cols * K * sizeof(cl_char);
+            const size_t n_blocks_padded = (size_t)cols * (K / 32);
+            const size_t size_qa_prev    = backend_ctx->prealloc_moe_qa.size;
+            const size_t size_da_prev    = backend_ctx->prealloc_moe_da.size;
+            backend_ctx->prealloc_moe_qa.allocate(context, size_qa);
+            backend_ctx->prealloc_moe_da.allocate(context, n_blocks_padded * sizeof(cl_half));
+            backend_ctx->prealloc_moe_sa.allocate(context, n_blocks_padded * sizeof(cl_half));
+            if (backend_ctx->prealloc_moe_qa.size != size_qa_prev) {
+                const cl_char z8 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_qa.buffer,
+                                             &z8, sizeof(z8), 0, backend_ctx->prealloc_moe_qa.size, 0, NULL, NULL));
+            }
+            if (backend_ctx->prealloc_moe_da.size != size_da_prev) {
+                const cl_half z16 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_da.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_da.size, 0, NULL, NULL));
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_sa.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_sa.size, 0, NULL, NULL));
+            }
+
+            cl_int n_blocks = (cl_int)((size_t)N * (K / 32));
+            cl_kernel kernel_quant = backend_ctx->kernel_quant_a_q8_1;
+            CL_CHECK(clSetKernelArg(kernel_quant, 0, sizeof(cl_mem), &b_sub_buf));
+            CL_CHECK(clSetKernelArg(kernel_quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 4, sizeof(cl_int), &n_blocks));
+            size_t local_work_size_quant[1]  = { 64 };
+            size_t global_work_size_quant[1] = { (size_t)((((size_t)n_blocks + 63) / 64) * 64) };
+            backend_ctx->enqueue_ndrange_kernel(kernel_quant, 1, global_work_size_quant, local_work_size_quant, dst);
+
+            cl_kernel kernel_coopk = (N <= 2) ? backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c2
+                                              : backend_ctx->kernel_gemm_coopk_q4_k_q8_1_dp4a_c4;
+            const cl_uchar mask_d6 = 0x3F, mask_d4 = 0x0F, mask_hi2 = 0xC0;
+            const cl_int   n_padded = cols;
+            CL_CHECK(clSetKernelArg(kernel_coopk,  0, sizeof(cl_mem),   &extra0_q4_k->q));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  1, sizeof(cl_mem),   &extra0_q4_k->s));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  2, sizeof(cl_mem),   &extra0_q4_k->d));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  3, sizeof(cl_mem),   &extra0_q4_k->dm));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  4, sizeof(cl_mem),   &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  5, sizeof(cl_mem),   &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  6, sizeof(cl_mem),   &backend_ctx->prealloc_moe_sa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  7, sizeof(cl_mem),   &extrad->data_device));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  8, sizeof(cl_ulong), &offsetd));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  9, sizeof(cl_int),   &M));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 10, sizeof(cl_int),   &n_padded));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 11, sizeof(cl_int),   &K));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 12, sizeof(cl_int),   &N));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 13, sizeof(cl_uchar), &mask_d6));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 14, sizeof(cl_uchar), &mask_d4));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 15, sizeof(cl_uchar), &mask_hi2));
+
+            const size_t nsg = (size_t)backend_ctx->nsg_coopk_dp4a_q4_k;
+            size_t local_work_size[3]  = { 64, nsg, 1 };
+            size_t global_work_size[3] = { (size_t)(M / backend_ctx->rows_coopk_dp4a_q4_k), nsg, 1 };
+            backend_ctx->enqueue_ndrange_kernel(kernel_coopk, 3, global_work_size, local_work_size, dst);
+
+            CL_CHECK(clReleaseMemObject(b_sub_buf));
+            return;
+        }
+
         // image for activations
         img_fmt = {CL_RGBA, CL_FLOAT};
         memset(&img_desc, 0, sizeof(img_desc));
@@ -22662,6 +23077,7 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
                 warned = true;
             }
         }
+        // This bin layout is not supported by CoopK.
         ggml_cl_mul_mat_q6_K_f32_adreno_ila(backend, src0, src1, dst);
         return;
     }
@@ -22817,6 +23233,77 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
         region.size = ne00 * ne1 * sizeof(float);
         CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
 
+        // CoopK uses the same untiled weight layout as gemm_noshuffle_q6_K_f32, with K-major Q8_1 activations (K contiguous).
+        // Keep output and embedding weights on the existing path for logit stability.
+        if (ggml_cl_coopk_dp4a_narrow_on(backend_ctx, dst, backend_ctx->rows_coopk_dp4a_q6_k)
+            && backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c4 != nullptr
+            && strncmp(src0->name, "output", 6) != 0
+            && strncmp(src0->name, "token_embd", 10) != 0
+            && ((ne1 <= 2) ? backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c2
+                           : backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c4) != nullptr) {
+            const int    N   = (int)ne1, K = (int)ne00;
+            const int    cols            = (N <= 2) ? 2 : 4;
+            const size_t n_blocks_padded = (size_t)cols * (K / 32);
+            const size_t size_qa         = (size_t)cols * K * sizeof(cl_char);
+            const size_t size_qa_prev    = backend_ctx->prealloc_moe_qa.size;
+            const size_t size_da_prev    = backend_ctx->prealloc_moe_da.size;
+            backend_ctx->prealloc_moe_qa.allocate(context, size_qa);
+            backend_ctx->prealloc_moe_da.allocate(context, n_blocks_padded * sizeof(cl_half));
+            backend_ctx->prealloc_moe_sa.allocate(context, n_blocks_padded * sizeof(cl_half));
+            if (backend_ctx->prealloc_moe_qa.size != size_qa_prev) {
+                const cl_char z8 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_qa.buffer,
+                                             &z8, sizeof(z8), 0, backend_ctx->prealloc_moe_qa.size, 0, NULL, NULL));
+            }
+            if (backend_ctx->prealloc_moe_da.size != size_da_prev) {
+                const cl_half z16 = 0;
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_da.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_da.size, 0, NULL, NULL));
+                CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, backend_ctx->prealloc_moe_sa.buffer,
+                                             &z16, sizeof(z16), 0, backend_ctx->prealloc_moe_sa.size, 0, NULL, NULL));
+            }
+
+            cl_int n_blocks = (cl_int)((size_t)N * (K / 32));
+            cl_kernel kernel_quant = backend_ctx->kernel_quant_a_q8_1;
+            CL_CHECK(clSetKernelArg(kernel_quant, 0, sizeof(cl_mem), &b_sub_buf));
+            CL_CHECK(clSetKernelArg(kernel_quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_quant, 4, sizeof(cl_int), &n_blocks));
+            size_t local_work_size_quant[1]  = { 64 };
+            size_t global_work_size_quant[1] = { (size_t)((((size_t)n_blocks + 63) / 64) * 64) };
+            backend_ctx->enqueue_ndrange_kernel(kernel_quant, 1, global_work_size_quant, local_work_size_quant, dst);
+
+            cl_kernel kernel_coopk = (N <= 2) ? backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c2
+                                              : backend_ctx->kernel_gemm_coopk_q6_k_q8_1_dp4a_c4;
+            const cl_ushort mask_f000 = 0xF000;
+            const cl_uchar  mask_c0 = 0xC0;
+            const cl_int    n_padded = cols;
+            const cl_int    m = (cl_int)ne01, k = (cl_int)ne00, n = (cl_int)ne1;
+            CL_CHECK(clSetKernelArg(kernel_coopk,  0, sizeof(cl_mem),    &extra0_q6_K->ql));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  1, sizeof(cl_mem),    &extra0_q6_K->qh));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  2, sizeof(cl_mem),    &extra0_q6_K->s));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  3, sizeof(cl_mem),    &extra0_q6_K->d));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  4, sizeof(cl_mem),    &backend_ctx->prealloc_moe_qa.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  5, sizeof(cl_mem),    &backend_ctx->prealloc_moe_da.buffer));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  6, sizeof(cl_mem),    &extrad->data_device));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  7, sizeof(cl_ulong),  &offsetd));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  8, sizeof(cl_int),    &m));
+            CL_CHECK(clSetKernelArg(kernel_coopk,  9, sizeof(cl_int),    &n_padded));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 10, sizeof(cl_int),    &k));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 11, sizeof(cl_int),    &n));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 12, sizeof(cl_ushort), &mask_f000));
+            CL_CHECK(clSetKernelArg(kernel_coopk, 13, sizeof(cl_uchar),  &mask_c0));
+
+            const size_t nsg = (size_t)backend_ctx->nsg_coopk_dp4a_q6_k;
+            size_t local_work_size[3]  = { 64, nsg, 1 };
+            size_t global_work_size[3] = { (size_t)(ne01 / backend_ctx->rows_coopk_dp4a_q6_k), nsg, 1 };
+            backend_ctx->enqueue_ndrange_kernel(kernel_coopk, 3, global_work_size, local_work_size, dst);
+
+            CL_CHECK(clReleaseMemObject(b_sub_buf));
+            return;
+        }
+
         // dp4a (int8) dense q6_K prefill GEMM
         static const char * q6k_dense_dp4a_env = getenv("GGML_OPENCL_Q6K_DENSE_DP4A");
                      bool   q6k_dense_dp4a_on  = (q6k_dense_dp4a_env != nullptr)
@@ -28726,6 +29213,9 @@ static void ggml_cl_cpy(ggml_backend_t backend, const ggml_tensor * src0, const
     switch (src0t) {
         case GGML_TYPE_F32:
             switch (src1t) {
+                case GGML_TYPE_I32:
+                    kernel = backend_ctx->kernel_cpy_f32_i32;
+                    break;
                 case GGML_TYPE_F16:
                     kernel = backend_ctx->kernel_cpy_f32_f16;
                     break;
@@ -28752,6 +29242,9 @@ static void ggml_cl_cpy(ggml_backend_t backend, const ggml_tensor * src0, const
             break;
         case GGML_TYPE_I32:
             switch (src1t) {
+                case GGML_TYPE_F32:
+                    kernel = backend_ctx->kernel_cpy_i32_f32;
+                    break;
                 case GGML_TYPE_I32:
                     kernel = backend_ctx->kernel_cpy_i32_i32;
                     break;

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

@@ -313,3 +313,88 @@ kernel void kernel_cpy_f32_f32_flat(
         }
     }
 }
+
+kernel void kernel_cpy_i32_f32(
+        global int * src0,
+        ulong offset0,
+        global float * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        int ne02,
+        int ne03,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne0,
+        int ne1,
+        int ne2,
+        int ne3,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = (global int *) ((global char *) src0 + offset0);
+    dst  = (global float *) ((global char *) dst  + offsetd);
+
+    const ulong i01 = get_group_id(0);
+    const ulong i02 = get_group_id(1);
+    const ulong i03 = get_group_id(2);
+
+    for (ulong i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
+        const ulong n  = ((i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+        const ulong i0 = n % ne0;
+        const ulong i1 = n / ne0 % ne1;
+        const ulong i2 = n / ne0 / ne1 % ne2;
+        const ulong i3 = n / ne0 / ne1 / ne2;
+
+        global const int * src = (global const int *) ((global char *) src0 + i00*nb00 + i01*nb01 + i02*nb02 + i03*nb03);
+        global float * out = (global float *) ((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+        out[0] = convert_float(src[0]);
+    }
+}
+
+
+kernel void kernel_cpy_f32_i32(
+        global float * src0,
+        ulong offset0,
+        global int * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        int ne02,
+        int ne03,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne0,
+        int ne1,
+        int ne2,
+        int ne3,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = (global float *) ((global char *) src0 + offset0);
+    dst  = (global int *) ((global char *) dst  + offsetd);
+
+    const ulong i01 = get_group_id(0);
+    const ulong i02 = get_group_id(1);
+    const ulong i03 = get_group_id(2);
+
+    for (ulong i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
+        const ulong n  = ((i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+        const ulong i0 = n % ne0;
+        const ulong i1 = n / ne0 % ne1;
+        const ulong i2 = n / ne0 / ne1 % ne2;
+        const ulong i3 = n / ne0 / ne1 / ne2;
+
+        global const float * src = (global const float *) ((global char *) src0 + i00*nb00 + i01*nb01 + i02*nb02 + i03*nb03);
+        global int * out = (global int *) ((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+        out[0] = convert_int_rtz(src[0]);
+    }
+}

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

@@ -0,0 +1,205 @@
+// Cooperative-K DP4A GEMM for 2 to 4 columns with Q4_0 weights and Q8_1 activations.
+// Named vector components avoid dynamic private-array indexing, which can cause register spills.
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#ifdef cl_khr_integer_dot_product
+#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
+#endif
+
+#ifndef COOPK_NSG
+#define COOPK_NSG 4
+#endif
+#define COOPK_SG  64
+
+#ifndef COOPK_ROWS
+#define COOPK_ROWS 4
+#endif
+#ifndef COOPK_COLS
+#define COOPK_COLS 4
+#endif
+
+#if COOPK_COLS == 4
+typedef float4 coopk_accv;
+typedef int4   coopk_dotv;
+#define COOPK_CONVF convert_float4
+#else
+typedef float2 coopk_accv;
+typedef int2   coopk_dotv;
+#define COOPK_CONVF convert_float2
+#endif
+
+// Pack four Q4_0 weights as signed bytes after subtracting 8.
+#define EXP40(l)                                                                      \
+    ( (((uint)((int)(( (l)         & 0x000Fu)       - 8) & 0xFF)))                    \
+    | (((uint)((int)((((l) & 0x00F0u) >>  4)        - 8) & 0xFF)) <<  8)              \
+    | (((uint)((int)((((l) & 0x0F00u) >>  8)        - 8) & 0xFF)) << 16)              \
+    | (((uint)((int)((((l) & 0xF000u) >> 12)        - 8) & 0xFF)) << 24) )
+
+#if COOPK_ROWS == 4
+#define COOPK_DOT(ci, t)                                               \
+    d0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, d0.s##ci); \
+    d1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, d1.s##ci); \
+    d2.s##ci = dot_acc_sat_4x8packed_ss_int(w2, A##ci.s##t, d2.s##ci); \
+    d3.s##ci = dot_acc_sat_4x8packed_ss_int(w3, A##ci.s##t, d3.s##ci);
+#else
+#define COOPK_DOT(ci, t)                                               \
+    d0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, d0.s##ci); \
+    d1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, d1.s##ci);
+#endif
+
+#if COOPK_COLS == 4
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t) COOPK_DOT(2,t) COOPK_DOT(3,t)
+#define COOPK_FOR_COLS(F) F(0) F(1) F(2) F(3)
+#else
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t)
+#define COOPK_FOR_COLS(F) F(0) F(1)
+#endif
+
+#if COOPK_ROWS == 4
+#define COOPK_KSTEP(t)                                     \
+    {                                                      \
+    ushort4 bl = vload4(0, src0_q + row0 + (ku0 + t) * m); \
+    const uint w0 = EXP40(bl.s0);                          \
+    const uint w1 = EXP40(bl.s1);                          \
+    const uint w2 = EXP40(bl.s2);                          \
+    const uint w3 = EXP40(bl.s3);                          \
+    COOPK_DOTS_AT(t)                                       \
+    }
+#else
+#define COOPK_KSTEP(t)                                     \
+    {                                                      \
+    ushort2 bl = vload2(0, src0_q + row0 + (ku0 + t) * m); \
+    const uint w0 = EXP40(bl.s0);                          \
+    const uint w1 = EXP40(bl.s1);                          \
+    COOPK_DOTS_AT(t)                                       \
+    }
+#endif
+
+// Weights use the untiled Adreno noshuffle SoA layout, with the output row contiguous in each plane.
+// Q8_1 activation planes are K-major (K contiguous); no activation transpose is used.
+kernel void kernel_gemm_coopk_q4_0_q8_1_dp4a(
+    global const ushort * src0_q,     // four consecutive K nibbles per ushort, [k/4][m]
+    global const half   * src0_d,     // block scales, [k/32][m]
+    global const uint   * src1_qa,    // four signed activation bytes per uint, [n][k/4]
+    global const half   * src1_da,    // activation scales, [n][k/32]
+    global float * dst,
+    ulong offsetd,
+    int m,
+    int n,
+    int k,
+    int n_no_padding
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const int gx   = get_global_id(0);
+    const int sg   = get_local_id(1);
+    const int lane = get_local_id(0);
+
+    const int row0 = gx * COOPK_ROWS;
+    const int num_32blk = k / 32;
+    const int k_u = k >> 2;
+    const int k_b = k >> 5;
+
+    // The host pads activations to COOPK_COLS; stores mask columns beyond n_no_padding.
+    const int c0 = 0, c1 = 1;
+#if COOPK_COLS == 4
+    const int c2 = 2, c3 = 3;
+#endif
+
+    coopk_accv acc0 = (coopk_accv)(0.0f), acc1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv acc2 = (coopk_accv)(0.0f), acc3 = (coopk_accv)(0.0f);
+#endif
+
+    for (int blk = sg; blk < num_32blk; blk += COOPK_NSG) {
+#if COOPK_ROWS == 4
+        half4 scale = vload4(0, src0_d + row0 + blk * m);
+#else
+        half2 scale = vload2(0, src0_d + row0 + blk * m);
+#endif
+
+        coopk_accv da;
+        da.s0 = (float)src1_da[c0*k_b + blk];
+        da.s1 = (float)src1_da[c1*k_b + blk];
+#if COOPK_COLS == 4
+        da.s2 = (float)src1_da[c2*k_b + blk];
+        da.s3 = (float)src1_da[c3*k_b + blk];
+#endif
+
+        // Q4_0 uses one weight scale per 32 K values.
+        coopk_dotv d0 = (coopk_dotv)(0), d1 = (coopk_dotv)(0);
+#if COOPK_ROWS == 4
+        coopk_dotv d2 = (coopk_dotv)(0), d3 = (coopk_dotv)(0);
+#endif
+        for (int half_idx = 0; half_idx < 2; ++half_idx) {
+            const int ku0 = (blk << 3) + half_idx * 4;
+
+            uint4 A0 = vload4(0, src1_qa + (uint)c0 * k_u + ku0);
+            uint4 A1 = vload4(0, src1_qa + (uint)c1 * k_u + ku0);
+#if COOPK_COLS == 4
+            uint4 A2 = vload4(0, src1_qa + (uint)c2 * k_u + ku0);
+            uint4 A3 = vload4(0, src1_qa + (uint)c3 * k_u + ku0);
+#endif
+            COOPK_KSTEP(0)
+            COOPK_KSTEP(1)
+            COOPK_KSTEP(2)
+            COOPK_KSTEP(3)
+        }
+
+        acc0 += (float)scale.s0 * da * COOPK_CONVF(d0);
+        acc1 += (float)scale.s1 * da * COOPK_CONVF(d1);
+#if COOPK_ROWS == 4
+        acc2 += (float)scale.s2 * da * COOPK_CONVF(d2);
+        acc3 += (float)scale.s3 * da * COOPK_CONVF(d3);
+#endif
+    }
+
+    // Reuse one local buffer for all rows; the first barrier protects the previous row's reads.
+    local coopk_accv reduceLM[COOPK_SG * (COOPK_NSG - 1)];
+    coopk_accv out0 = (coopk_accv)(0.0f), out1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv out2 = (coopk_accv)(0.0f), out3 = (coopk_accv)(0.0f);
+#endif
+
+#define COOPK_REDUCE(accv, outv)                                   \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg > 0) { reduceLM[(sg - 1) * COOPK_SG + lane] = (accv); } \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg == 0) {                                                 \
+        coopk_accv sum = (accv);                                   \
+        for (int s = 0; s < COOPK_NSG - 1; s++) {                  \
+            sum += reduceLM[s * COOPK_SG + lane];                  \
+        }                                                          \
+        (outv) = sum;                                              \
+    }
+
+    COOPK_REDUCE(acc0, out0)
+    COOPK_REDUCE(acc1, out1)
+#if COOPK_ROWS == 4
+    COOPK_REDUCE(acc2, out2)
+    COOPK_REDUCE(acc3, out3)
+#endif
+
+#undef COOPK_REDUCE
+
+#if COOPK_ROWS == 4
+#define COOPK_STORE_COL(ci)                                                              \
+    if (idx < m*n_no_padding) {                                                          \
+        vstore4((float4)(out0.s##ci, out1.s##ci, out2.s##ci, out3.s##ci), 0, dst + idx); \
+        idx += m;                                                                        \
+    }
+#else
+#define COOPK_STORE_COL(ci)                                      \
+    if (idx < m*n_no_padding) {                                  \
+        vstore2((float2)(out0.s##ci, out1.s##ci), 0, dst + idx); \
+        idx += m;                                                \
+    }
+#endif
+
+    if (sg == 0) {
+        int idx = row0;
+        COOPK_FOR_COLS(COOPK_STORE_COL)
+    }
+
+#undef COOPK_STORE_COL
+}

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

@@ -0,0 +1,264 @@
+// Cooperative-K DP4A GEMM for 2 to 4 columns with Q4_K weights and Q8_1 activations.
+// Named vector components avoid dynamic private-array indexing, which can cause register spills.
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#ifdef cl_khr_integer_dot_product
+#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
+#endif
+
+#define QK_K          256
+#define K_SCALE_SIZE   12
+
+#ifndef COOPK_NSG
+#define COOPK_NSG 4
+#endif
+#define COOPK_SG  64
+
+#ifndef COOPK_ROWS
+#define COOPK_ROWS 4
+#endif
+#ifndef COOPK_COLS
+#define COOPK_COLS 4
+#endif
+
+#if COOPK_COLS == 4
+typedef float4 coopk_accv;
+typedef int4   coopk_dotv;
+#define COOPK_CONVF convert_float4
+#else
+typedef float2 coopk_accv;
+typedef int2   coopk_dotv;
+#define COOPK_CONVF convert_float2
+#endif
+
+// Expand four Q4_K nibbles into bytes for DP4A.
+#define EXP4(u)  ( ((uint)((u) & 0x000Fu))        | \
+                  (((uint)((u) & 0x00F0u)) << 4)  | \
+                  (((uint)((u) & 0x0F00u)) << 8)  | \
+                  (((uint)((u) & 0xF000u)) << 12) )
+
+#if COOPK_ROWS == 4
+#define COOPK_DOT(ci, t)                                               \
+    s0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, s0.s##ci); \
+    s1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, s1.s##ci); \
+    s2.s##ci = dot_acc_sat_4x8packed_ss_int(w2, A##ci.s##t, s2.s##ci); \
+    s3.s##ci = dot_acc_sat_4x8packed_ss_int(w3, A##ci.s##t, s3.s##ci);
+#else
+#define COOPK_DOT(ci, t)                                               \
+    s0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, s0.s##ci); \
+    s1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, s1.s##ci);
+#endif
+
+#if COOPK_COLS == 4
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t) COOPK_DOT(2,t) COOPK_DOT(3,t)
+#define COOPK_FOR_COLS(F) F(0) F(1) F(2) F(3)
+#else
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t)
+#define COOPK_FOR_COLS(F) F(0) F(1)
+#endif
+
+// The bin layout packs eight K values in each uint, with four nibbles in each 16-bit half.
+// Since ku0 is a multiple of four, t / 2 selects the uint and t & 1 selects its half.
+#ifdef COOPK_Q_BIN
+#define COOPK_QBITS4(t) convert_ushort4(((t) & 1) ? (vload4(0, src0_q + row0 + ((ku0 >> 1) + ((t) >> 1)) * m) >> 16) \
+                                                : (vload4(0, src0_q + row0 + ((ku0 >> 1) + ((t) >> 1)) * m) & 0xFFFFu))
+#define COOPK_QBITS2(t) convert_ushort2(((t) & 1) ? (vload2(0, src0_q + row0 + ((ku0 >> 1) + ((t) >> 1)) * m) >> 16) \
+                                                : (vload2(0, src0_q + row0 + ((ku0 >> 1) + ((t) >> 1)) * m) & 0xFFFFu))
+#else
+#define COOPK_QBITS4(t) vload4(0, src0_q + row0 + (ku0 + t) * m)
+#define COOPK_QBITS2(t) vload2(0, src0_q + row0 + (ku0 + t) * m)
+#endif
+
+#if COOPK_ROWS == 4
+#define COOPK_KSTEP(t)              \
+    {                               \
+    ushort4 bits = COOPK_QBITS4(t); \
+    const uint w0 = EXP4(bits.s0);  \
+    const uint w1 = EXP4(bits.s1);  \
+    const uint w2 = EXP4(bits.s2);  \
+    const uint w3 = EXP4(bits.s3);  \
+    COOPK_DOTS_AT(t)                \
+    }
+#else
+#define COOPK_KSTEP(t)              \
+    {                               \
+    ushort2 bits = COOPK_QBITS2(t); \
+    const uint w0 = EXP4(bits.s0);  \
+    const uint w1 = EXP4(bits.s1);  \
+    COOPK_DOTS_AT(t)                \
+    }
+#endif
+
+inline void get_scale_min_k4_c(int j, global const uchar * q, int stride,
+                               uchar * d, uchar * m,
+                               uchar mask_d6, uchar mask_d4, uchar mask_hi2) {
+    if (j < 4) {
+        *d = q[j*stride] & mask_d6;
+        *m = q[(j + 4)*stride] & mask_d6;
+    } else {
+        *d = (q[(j + 4)*stride] & mask_d4) | ((q[(j - 4)*stride] >> 6) << 4);
+        *m = (q[(j + 4)*stride] >>   4)    | ((q[(j    )*stride] >> 6) << 4);
+    }
+}
+
+// Weights use the untiled Adreno noshuffle SoA layout; COOPK_Q_BIN selects uint-transposed q instead of ushort-transposed q.
+// Q8_1 activation planes are K-major (K contiguous); no activation transpose is used.
+kernel void kernel_gemm_coopk_q4_k_q8_1_dp4a(
+#ifdef COOPK_Q_BIN
+    global const uint   * src0_q,     // eight consecutive K nibbles per uint, [k/8][m]
+#else
+    global const ushort * src0_q,     // four consecutive K nibbles per ushort, [k/4][m]
+#endif
+    global const uchar  * src0_s,     // packed 6-bit scales/mins, [k/256][12][m]
+    global const half   * src0_d,     // super-block scales, [k/256][m]
+    global const half   * src0_dm,    // super-block mins, [k/256][m]
+    global const uint   * src1_qa,    // four signed activation bytes per uint, [n][k/4]
+    global const half   * src1_da,    // activation scales, [n][k/32]
+    global const half   * src1_sa,    // scaled activation sums, [n][k/32]
+    global float * dst,
+    ulong offsetd,
+    int m,
+    int n,
+    int k,
+    int n_no_padding,
+    uchar mask_d6,
+    uchar mask_d4,
+    uchar mask_hi2
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const int gx   = get_global_id(0);   // row group
+    const int sg   = get_local_id(1);    // K-split subgroup
+    const int lane = get_local_id(0);
+
+    const int row0 = gx * COOPK_ROWS;
+    const int num_32blk = k / 32;
+    const int k_u = k >> 2;              // K in uint (int8x4) units
+    const int k_b = k >> 5;              // 32-blocks along K
+
+    // The host pads activations to COOPK_COLS; stores mask columns beyond n_no_padding.
+    const int nl = n - 1;
+    const int c0 = 0;
+    const int c1 = (1 < n) ? 1 : nl;
+#if COOPK_COLS >= 4
+    const int c2 = (2 < n) ? 2 : nl;
+    const int c3 = (3 < n) ? 3 : nl;
+#endif
+
+    coopk_accv acc0 = (coopk_accv)(0.0f), acc1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv acc2 = (coopk_accv)(0.0f), acc3 = (coopk_accv)(0.0f);
+#endif
+
+    for (int blk = sg; blk < num_32blk; blk += COOPK_NSG) {
+        const int i       = blk << 5;
+        const int sb_idx  = blk >> 3;
+        const int sub_idx = blk & 7;
+
+        global const uchar * sc = src0_s + sb_idx * K_SCALE_SIZE * m + row0;
+        uchar sv0, mn0, sv1, mn1;
+        get_scale_min_k4_c(sub_idx, sc + 0, m, &sv0, &mn0, mask_d6, mask_d4, mask_hi2);
+        get_scale_min_k4_c(sub_idx, sc + 1, m, &sv1, &mn1, mask_d6, mask_d4, mask_hi2);
+#if COOPK_ROWS == 4
+        half4 dd  = vload4(0, src0_d  + row0 + sb_idx * m);
+        half4 dmm = vload4(0, src0_dm + row0 + sb_idx * m);
+        uchar sv2, mn2, sv3, mn3;
+        get_scale_min_k4_c(sub_idx, sc + 2, m, &sv2, &mn2, mask_d6, mask_d4, mask_hi2);
+        get_scale_min_k4_c(sub_idx, sc + 3, m, &sv3, &mn3, mask_d6, mask_d4, mask_hi2);
+        const float sc2 = (float)dd.s2 * (float)sv2, mv2 = (float)dmm.s2 * (float)mn2;
+        const float sc3 = (float)dd.s3 * (float)sv3, mv3 = (float)dmm.s3 * (float)mn3;
+#else
+        half2 dd  = vload2(0, src0_d  + row0 + sb_idx * m);
+        half2 dmm = vload2(0, src0_dm + row0 + sb_idx * m);
+#endif
+        const float sc0 = (float)dd.s0 * (float)sv0, mv0 = (float)dmm.s0 * (float)mn0;
+        const float sc1 = (float)dd.s1 * (float)sv1, mv1 = (float)dmm.s1 * (float)mn1;
+
+        // Q4_K scale and min values change every 32 K values.
+        coopk_dotv s0 = (coopk_dotv)(0), s1 = (coopk_dotv)(0);
+#if COOPK_ROWS == 4
+        coopk_dotv s2 = (coopk_dotv)(0), s3 = (coopk_dotv)(0);
+#endif
+
+        for (int uq = 0; uq < 2; ++uq) {
+            const int ku0 = (i >> 2) + uq * 4;
+
+            uint4 A0 = vload4(0, src1_qa + (uint)c0 * k_u + ku0);
+            uint4 A1 = vload4(0, src1_qa + (uint)c1 * k_u + ku0);
+#if COOPK_COLS >= 4
+            uint4 A2 = vload4(0, src1_qa + (uint)c2 * k_u + ku0);
+            uint4 A3 = vload4(0, src1_qa + (uint)c3 * k_u + ku0);
+#endif
+            COOPK_KSTEP(0)
+            COOPK_KSTEP(1)
+            COOPK_KSTEP(2)
+            COOPK_KSTEP(3)
+        }
+
+        // Q4_K stores q * scale - min for each 32-value block.
+        // src1_sa already includes the activation scale, so the min term needs no extra factor of da.
+        coopk_accv da, sa;
+        da.s0 = (float)src1_da[c0*k_b + blk];  sa.s0 = (float)src1_sa[c0*k_b + blk];
+        da.s1 = (float)src1_da[c1*k_b + blk];  sa.s1 = (float)src1_sa[c1*k_b + blk];
+#if COOPK_COLS >= 4
+        da.s2 = (float)src1_da[c2*k_b + blk];  sa.s2 = (float)src1_sa[c2*k_b + blk];
+        da.s3 = (float)src1_da[c3*k_b + blk];  sa.s3 = (float)src1_sa[c3*k_b + blk];
+#endif
+
+        acc0 += sc0 * da * COOPK_CONVF(s0) - mv0 * sa;
+        acc1 += sc1 * da * COOPK_CONVF(s1) - mv1 * sa;
+#if COOPK_ROWS == 4
+        acc2 += sc2 * da * COOPK_CONVF(s2) - mv2 * sa;
+        acc3 += sc3 * da * COOPK_CONVF(s3) - mv3 * sa;
+#endif
+    }
+
+    // Reuse one local buffer for all rows; the first barrier protects the previous row's reads.
+    local coopk_accv reduceLM[COOPK_SG * (COOPK_NSG - 1)];
+    coopk_accv out0 = (coopk_accv)(0.0f), out1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv out2 = (coopk_accv)(0.0f), out3 = (coopk_accv)(0.0f);
+#endif
+
+#define COOPK_REDUCE(accv, outv)                                   \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg > 0) { reduceLM[(sg - 1) * COOPK_SG + lane] = (accv); } \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg == 0) {                                                 \
+        coopk_accv sum = (accv);                                   \
+        for (int s = 0; s < COOPK_NSG - 1; s++) {                  \
+            sum += reduceLM[s * COOPK_SG + lane];                  \
+        }                                                          \
+        (outv) = sum;                                              \
+    }
+
+    COOPK_REDUCE(acc0, out0)
+    COOPK_REDUCE(acc1, out1)
+#if COOPK_ROWS == 4
+    COOPK_REDUCE(acc2, out2)
+    COOPK_REDUCE(acc3, out3)
+#endif
+
+#undef COOPK_REDUCE
+
+#if COOPK_ROWS == 4
+#define COOPK_STORE_COL(ci)                                                              \
+    if (idx < m*n_no_padding) {                                                          \
+        vstore4((float4)(out0.s##ci, out1.s##ci, out2.s##ci, out3.s##ci), 0, dst + idx); \
+        idx += m;                                                                        \
+    }
+#else
+#define COOPK_STORE_COL(ci)                                      \
+    if (idx < m*n_no_padding) {                                  \
+        vstore2((float2)(out0.s##ci, out1.s##ci), 0, dst + idx); \
+        idx += m;                                                \
+    }
+#endif
+
+    if (sg == 0) {
+        int idx = row0;
+        COOPK_FOR_COLS(COOPK_STORE_COL)
+    }
+
+#undef COOPK_STORE_COL
+}

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

@@ -0,0 +1,216 @@
+// Cooperative-K DP4A GEMM for 2 to 4 columns with Q6_K weights and Q8_1 activations.
+// Named vector components avoid dynamic private-array indexing, which can cause register spills.
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#ifdef cl_khr_integer_dot_product
+#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
+#endif
+
+#ifndef COOPK_NSG
+#define COOPK_NSG 4
+#endif
+#define COOPK_SG  64
+
+#ifndef COOPK_ROWS
+#define COOPK_ROWS 4
+#endif
+#ifndef COOPK_COLS
+#define COOPK_COLS 4
+#endif
+
+#if COOPK_COLS == 4
+typedef float4 coopk_accv;
+typedef int4   coopk_dotv;
+#define COOPK_CONVF convert_float4
+#else
+typedef float2 coopk_accv;
+typedef int2   coopk_dotv;
+#define COOPK_CONVF convert_float2
+#endif
+
+// Combine low nibbles and high bits, subtract 32, and pack four signed bytes for DP4A.
+#define EXP6(l, h, mf, mc)                                                                    \
+    ( (((uint)((int)((( (l)        & 0x000Fu) | (((h) & 0x03u) << 4)) - 32) & 0xFF)))         \
+    | (((uint)((int)(((((l) & 0x00F0u) >> 4)  | (((h) & 0x0Cu) << 2)) - 32) & 0xFF)) <<  8)   \
+    | (((uint)((int)(((((l) & 0x0F00u) >> 8)  |  ((h) & 0x30u))       - 32) & 0xFF)) << 16)   \
+    | (((uint)((int)(((((l) & (mf))   >> 12)  | (((h) & (mc)) >> 2))  - 32) & 0xFF)) << 24) )
+
+#if COOPK_ROWS == 4
+#define COOPK_DOT(ci, t)                                               \
+    d0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, d0.s##ci); \
+    d1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, d1.s##ci); \
+    d2.s##ci = dot_acc_sat_4x8packed_ss_int(w2, A##ci.s##t, d2.s##ci); \
+    d3.s##ci = dot_acc_sat_4x8packed_ss_int(w3, A##ci.s##t, d3.s##ci);
+#else
+#define COOPK_DOT(ci, t)                                               \
+    d0.s##ci = dot_acc_sat_4x8packed_ss_int(w0, A##ci.s##t, d0.s##ci); \
+    d1.s##ci = dot_acc_sat_4x8packed_ss_int(w1, A##ci.s##t, d1.s##ci);
+#endif
+
+#if COOPK_COLS == 4
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t) COOPK_DOT(2,t) COOPK_DOT(3,t)
+#define COOPK_FOR_COLS(F) F(0) F(1) F(2) F(3)
+#else
+#define COOPK_DOTS_AT(t)  COOPK_DOT(0,t) COOPK_DOT(1,t)
+#define COOPK_FOR_COLS(F) F(0) F(1)
+#endif
+
+#if COOPK_ROWS == 4
+#define COOPK_KSTEP(t)                                      \
+    {                                                       \
+    ushort4 bl = vload4(0, src0_ql + row0 + (ku0 + t) * m); \
+    uchar4  bh = vload4(0, src0_qh + row0 + (ku0 + t) * m); \
+    const uint w0 = EXP6(bl.s0, bh.s0, mask_f000, mask_c0); \
+    const uint w1 = EXP6(bl.s1, bh.s1, mask_f000, mask_c0); \
+    const uint w2 = EXP6(bl.s2, bh.s2, mask_f000, mask_c0); \
+    const uint w3 = EXP6(bl.s3, bh.s3, mask_f000, mask_c0); \
+    COOPK_DOTS_AT(t)                                        \
+    }
+#else
+#define COOPK_KSTEP(t)                                      \
+    {                                                       \
+    ushort2 bl = vload2(0, src0_ql + row0 + (ku0 + t) * m); \
+    uchar2  bh = vload2(0, src0_qh + row0 + (ku0 + t) * m); \
+    const uint w0 = EXP6(bl.s0, bh.s0, mask_f000, mask_c0); \
+    const uint w1 = EXP6(bl.s1, bh.s1, mask_f000, mask_c0); \
+    COOPK_DOTS_AT(t)                                        \
+    }
+#endif
+
+// Weights use the untiled Adreno noshuffle SoA layout, with the output row contiguous in each plane.
+// Q8_1 activation planes are K-major (K contiguous); no activation transpose is used.
+kernel void kernel_gemm_coopk_q6_k_q8_1_dp4a(
+    global const ushort * src0_ql,    // four consecutive low nibbles per ushort, [k/4][m]
+    global const uchar  * src0_qh,    // four consecutive high 2-bit codes per uchar, [k/4][m]
+    global const ushort * src0_s,     // two signed int8 scales per ushort, [k/32][m]
+    global const half   * src0_d,     // super-block scales, [k/256][m]
+    global const uint   * src1_qa,    // four signed activation bytes per uint, [n][k/4]
+    global const half   * src1_da,    // activation scales, [n][k/32]
+    global float * dst,
+    ulong offsetd,
+    int m,
+    int n,
+    int k,
+    int n_no_padding,
+    ushort mask_f000,
+    uchar  mask_c0
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const int gx   = get_global_id(0);
+    const int sg   = get_local_id(1);
+    const int lane = get_local_id(0);
+
+    const int row0 = gx * COOPK_ROWS;
+    const int num_32blk = k / 32;
+    const int k_u = k >> 2;
+    const int k_b = k >> 5;
+
+    // The host pads activations to COOPK_COLS; stores mask columns beyond n_no_padding.
+    const int c0 = 0, c1 = 1;
+#if COOPK_COLS == 4
+    const int c2 = 2, c3 = 3;
+#endif
+
+    coopk_accv acc0 = (coopk_accv)(0.0f), acc1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv acc2 = (coopk_accv)(0.0f), acc3 = (coopk_accv)(0.0f);
+#endif
+
+    for (int blk = sg; blk < num_32blk; blk += COOPK_NSG) {
+        const int i = blk << 5;
+
+        // Each ushort packs two signed scale bytes, one for each 16-value half of this 32-value block.
+        // The half-precision super-block scale is shared by 256 K values.
+        ushort4 spk = vload4(0, src0_s + row0 + blk * m);
+        half4   scd = vload4(0, src0_d + row0 + (i >> 8) * m);
+
+        coopk_accv da;
+        da.s0 = (float)src1_da[c0*k_b + blk];
+        da.s1 = (float)src1_da[c1*k_b + blk];
+#if COOPK_COLS == 4
+        da.s2 = (float)src1_da[c2*k_b + blk];
+        da.s3 = (float)src1_da[c3*k_b + blk];
+#endif
+
+        // Apply each signed weight scale to its 16-value dot product.
+        for (int half_idx = 0; half_idx < 2; ++half_idx) {
+            const int ku0 = (i >> 2) + half_idx * 4;
+
+            coopk_dotv d0 = (coopk_dotv)(0), d1 = (coopk_dotv)(0);
+#if COOPK_ROWS == 4
+            coopk_dotv d2 = (coopk_dotv)(0), d3 = (coopk_dotv)(0);
+#endif
+            uint4 A0 = vload4(0, src1_qa + (uint)c0 * k_u + ku0);
+            uint4 A1 = vload4(0, src1_qa + (uint)c1 * k_u + ku0);
+#if COOPK_COLS == 4
+            uint4 A2 = vload4(0, src1_qa + (uint)c2 * k_u + ku0);
+            uint4 A3 = vload4(0, src1_qa + (uint)c3 * k_u + ku0);
+#endif
+            COOPK_KSTEP(0)
+            COOPK_KSTEP(1)
+            COOPK_KSTEP(2)
+            COOPK_KSTEP(3)
+
+            const float s0f = (float)(half_idx ? as_char2(spk.s0).s1 : as_char2(spk.s0).s0) * (float)scd.s0;
+            const float s1f = (float)(half_idx ? as_char2(spk.s1).s1 : as_char2(spk.s1).s0) * (float)scd.s1;
+            acc0 += s0f * da * COOPK_CONVF(d0);
+            acc1 += s1f * da * COOPK_CONVF(d1);
+#if COOPK_ROWS == 4
+            const float s2f = (float)(half_idx ? as_char2(spk.s2).s1 : as_char2(spk.s2).s0) * (float)scd.s2;
+            const float s3f = (float)(half_idx ? as_char2(spk.s3).s1 : as_char2(spk.s3).s0) * (float)scd.s3;
+            acc2 += s2f * da * COOPK_CONVF(d2);
+            acc3 += s3f * da * COOPK_CONVF(d3);
+#endif
+        }
+    }
+
+    // Reuse one local buffer for all rows; the first barrier protects the previous row's reads.
+    local coopk_accv reduceLM[COOPK_SG * (COOPK_NSG - 1)];
+    coopk_accv out0 = (coopk_accv)(0.0f), out1 = (coopk_accv)(0.0f);
+#if COOPK_ROWS == 4
+    coopk_accv out2 = (coopk_accv)(0.0f), out3 = (coopk_accv)(0.0f);
+#endif
+
+#define COOPK_REDUCE(accv, outv)                                   \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg > 0) { reduceLM[(sg - 1) * COOPK_SG + lane] = (accv); } \
+    barrier(CLK_LOCAL_MEM_FENCE);                                  \
+    if (sg == 0) {                                                 \
+        coopk_accv sum = (accv);                                   \
+        for (int s = 0; s < COOPK_NSG - 1; s++) {                  \
+            sum += reduceLM[s * COOPK_SG + lane];                  \
+        }                                                          \
+        (outv) = sum;                                              \
+    }
+
+    COOPK_REDUCE(acc0, out0)
+    COOPK_REDUCE(acc1, out1)
+#if COOPK_ROWS == 4
+    COOPK_REDUCE(acc2, out2)
+    COOPK_REDUCE(acc3, out3)
+#endif
+
+#undef COOPK_REDUCE
+
+#if COOPK_ROWS == 4
+#define COOPK_STORE_COL(ci)                                                              \
+    if (idx < m*n_no_padding) {                                                          \
+        vstore4((float4)(out0.s##ci, out1.s##ci, out2.s##ci, out3.s##ci), 0, dst + idx); \
+        idx += m;                                                                        \
+    }
+#else
+#define COOPK_STORE_COL(ci)                                      \
+    if (idx < m*n_no_padding) {                                  \
+        vstore2((float2)(out0.s##ci, out1.s##ci), 0, dst + idx); \
+        idx += m;                                                \
+    }
+#endif
+
+    if (sg == 0) {
+        int idx = row0;
+        COOPK_FOR_COLS(COOPK_STORE_COL)
+    }
+
+#undef COOPK_STORE_COL
+}

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

@@ -106,6 +106,53 @@ kernel void kernel_get_rows_f32(
     }
 }
 
+kernel void kernel_get_rows_i32(
+        global void * src0,
+        ulong offset0,
+        global int * src1,
+        ulong offset1,
+        global int * dst,
+        ulong offsetd,
+        int ne00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne10,
+        ulong nb10,
+        ulong nb11,
+        ulong nb12,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = (global void*)((global char*)src0 + offset0);
+    src1 = (global int*)((global char*)src1 + offset1);
+    dst = (global int*)((global char*)dst + offsetd);
+
+    int nchunks = get_num_groups(0) / ne10;
+    int g       = get_group_id(0);
+    int i10     = g / nchunks;
+    int chunk   = g - i10 * nchunks;
+    int i11     = get_group_id(1);
+    int i12     = get_group_id(2);
+
+    int r = ((global int *) ((global char *) src1 + i12*nb12 + i11*nb11 + i10*nb10))[0];
+
+    int i02 = i11;
+    int i03 = i12;
+
+    global int * dst_row = (global int *) ((global char *) dst  + i12*nb3 + i11*nb2 + i10*nb1);
+    global int * src_row = (global int *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03);
+
+    int span  = (ne00 + nchunks - 1) / nchunks;
+    int start = chunk * span;
+    int end   = min(start + span, ne00);
+
+    for (int ind = start + get_local_id(0); ind < end; ind += get_local_size(0)) {
+        dst_row[ind] = src_row[ind];
+    }
+}
+
 kernel void kernel_get_rows_f16(
         global void * src0,
         ulong offset0,

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

@@ -1,9 +1,7 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-// Quantize a contiguous [N, K] f32 activation buffer (token-major, K contiguous
-// per token) into q8_1 blocks of 32: int8 quants + per-block scale d + per-block
-// sum s (= d * Sum(qs)). Consumed by kernel_gemm_noshuffle_q4_k_q8_1_dp4a for the
-// dp4a (int8) dense q4_K prefill GEMM. One work-item per 32-element block.
+// Quantize a contiguous K-major [N, K] f32 activation buffer (K contiguous) into 32-value Q8_1 SoA blocks.
+// Store int8 quants, per-block scale d, and scaled sum s = d * sum(qs); one work-item handles each block.
 __kernel void kernel_quant_a_q8_1(
         __global const float * src,   // [N * K]
         __global       char  * qa,    // [N * K]

ggml/src/ggml-sycl/CMakeLists.txt

@@ -44,7 +44,7 @@ if (WIN32)
         if(DEFINED ENV{LEVEL_ZERO_V1_SDK_PATH})
             set(LEVEL_ZERO_V1_SDK_PATH $ENV{LEVEL_ZERO_V1_SDK_PATH})
             if(EXISTS "${LEVEL_ZERO_V1_SDK_PATH}")
-                target_include_directories(ggml-sycl PRIVATE "${LEVEL_ZERO_V1_SDK_PATH}/include")
+                target_include_directories(ggml-sycl SYSTEM PRIVATE "${LEVEL_ZERO_V1_SDK_PATH}/include")
                 set(LEVEL_ZERO_V1_SDK_LIB_PATH "${LEVEL_ZERO_V1_SDK_PATH}/lib")
             else()
                 message(WARNING "LEVEL_ZERO_V1_SDK_PATH set but folder not found: ${LEVEL_ZERO_V1_SDK_PATH}")

ggml/src/ggml-sycl/concat.cpp

@@ -504,6 +504,5 @@ void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
     default:
         fprintf(stderr, "%s: unsupported types: dst: %s\n", __func__, ggml_type_name(dst->type));
         GGML_ASSERT(false);
-    break;
     }
 }

ggml/src/ggml-sycl/convert.cpp

@@ -85,7 +85,7 @@ static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int
     stream->parallel_for(
         sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
         [=](sycl::nd_item<3> item_ct1) {
-            dequantize_block_q2_K_reorder(vx, y, item_ct1, nb);
+            dequantize_block_q2_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
         });
 }
 
@@ -129,7 +129,7 @@ static void dequantize_row_q3_K_sycl_reorder(const void *vx, dst_t *y, const int
     stream->parallel_for(
         sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
         [=](sycl::nd_item<3> item_ct1) {
-            dequantize_block_q3_K_reorder(vx, y, item_ct1, nb);
+            dequantize_block_q3_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
         });
 }
 
@@ -231,7 +231,7 @@ static void dequantize_row_q4_K_sycl_reorder(const void * vx, dst_t * y, const i
 
         cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(global_size), sycl::range<1>(local_size)),
                          [=](sycl::nd_item<1> item_ct1) {
-                             dequantize_block_q4_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
+                             dequantize_block_q4_K_reorder(vx, y + item_ct1.get_group(0) * QK_K, item_ct1.get_group(0), get_pointer(scale_local_acc), item_ct1, nb);
                          });
     });
 }
@@ -258,7 +258,7 @@ static void dequantize_row_q5_K_sycl_reorder(const void * vx, dst_t * y, const i
         cgh.parallel_for(
             sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
             [=](sycl::nd_item<3> item_ct1) {
-                dequantize_block_q5_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
+                dequantize_block_q5_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), get_pointer(scale_local_acc), item_ct1, nb);
             });
     });
 }
@@ -303,7 +303,7 @@ static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const i
 
     stream->parallel_for(
         sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
-        [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
+        [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb); });
 }
 
 template <typename dst_t>
@@ -757,7 +757,6 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
 #endif
         default:
             GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
-            return nullptr;
     }
 }
 
@@ -861,7 +860,6 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
 #endif
         default:
             GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
-            return nullptr;
     }
 }
 
@@ -877,7 +875,6 @@ to_bf16_sycl_t ggml_get_to_bf16_sycl(ggml_type type, ggml_tensor * /*dst*/) {
             return convert_unary_sycl<sycl::ext::oneapi::bfloat16>;
         default:
             GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
-            return nullptr;
     }
 }
 #endif

ggml/src/ggml-sycl/count-equal.cpp

@@ -74,6 +74,5 @@ void ggml_sycl_count_equal(ggml_backend_sycl_context &ctx, ggml_tensor *dst) {
     } break;
     default:
         GGML_ASSERT(false);
-        break;
     }
 }

ggml/src/ggml-sycl/dequantize.hpp

@@ -926,10 +926,9 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri
 }
 
 template<typename dst_t>
-static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
                                           const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
 #if QK_K == 256
-    const int64_t i = item_ct1.get_group(2);
     if (i >= n_blocks) {
         return;
     }
@@ -949,7 +948,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
     const int64_t is   = 8 * n + l / 16;
 
     const uint8_t q = qs[32 * n + l];
-    dst_t * y = yy + i * QK_K + 128 * n;
+    dst_t * y = yy + 128 * n;
 
     const float dall = (*dm)[0];
     const float dmin = (*dm)[1];
@@ -960,6 +959,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
 #else
     GGML_UNUSED(vx);
     GGML_UNUSED(yy);
+    GGML_UNUSED(i);
     GGML_UNUSED(item_ct1);
     GGML_UNUSED(n_blocks);
     GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256");
@@ -1022,10 +1022,9 @@ static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restri
 }
 
 template<typename dst_t>
-static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
                                           const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
 #if QK_K == 256
-    const int64_t i = item_ct1.get_group(2);
     if (i >= n_blocks) {
         return;
     }
@@ -1062,7 +1061,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
 
     const float dl = d_all * (us - 32);
 
-    dst_t * y = yy + i * QK_K + 128 * n + 32 * j;
+    dst_t * y = yy + 128 * n + 32 * j;
     const uint8_t * q  = qs + 32 * n;
     const uint8_t * hm = hmask;
 
@@ -1072,6 +1071,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
 #else
     GGML_UNUSED(vx);
     GGML_UNUSED(yy);
+    GGML_UNUSED(i);
     GGML_UNUSED(item_ct1);
     GGML_UNUSED(n_blocks);
     GGML_ABORT("Q3_K reorder dequantize not supported for QK_K != 256");
@@ -1160,14 +1160,13 @@ inline void dequantize_q4_K_common(dst_t * __restrict__ y, const uint8_t * __res
 }
 
 template <typename dst_t>
-static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, uint8_t * scales_local,
-                                          const sycl::nd_item<1> & item_ct1, int64_t nb) {
-    const int64_t i   = item_ct1.get_group(0);     // block index
+static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
+                                          uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) {
     const int64_t tid = item_ct1.get_local_id(0);  // thread index within block
     const int64_t il  = tid / 8;
     const int64_t ir  = tid % 8;
 
-    dst_t * y = yy + i * QK_K + 64 * il + 4 * ir;
+    dst_t * y = yy + 64 * il + 4 * ir;
 
     const uint8_t * base          = static_cast<const uint8_t *>(vx);
     const size_t    qs_offset     = i * (QK_K / 2);
@@ -1237,9 +1236,8 @@ static void dequantize_block_q5_K_wide(const void * __restrict__ vx, dst_t * __r
 }
 
 template <typename dst_t>
-static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
                                           uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
-    const int64_t ib = item_ct1.get_group(2);
 
 #if QK_K == 256
     // assume 64 threads
@@ -1248,7 +1246,7 @@ static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t *
     const int64_t ir  = tid % 16;   // 0...15
     const int64_t is  = 2 * il;
 
-    dst_t * y = yy + ib * QK_K + 64 * il + 2 * ir;
+    dst_t * y = yy + 64 * il + 2 * ir;
 
     const uint8_t * base = static_cast<const uint8_t *>(vx);
 
@@ -1340,9 +1338,8 @@ static void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restri
 }
 
 template <typename dst_t>
-static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
                                           const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
-    const int64_t ib = item_ct1.get_group(2);
 
     const int64_t tid = item_ct1.get_local_id(2);
     const int64_t ip  = tid / 32;       // ip is 0 or 1
@@ -1359,7 +1356,7 @@ static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t *
     const uint8_t *   scales_ptr         = base_ptr + base_scales_offset;
     const ggml_half * d                  = (const ggml_half *) (base_ptr + base_d_offset) + ib;
 
-    dst_t * y = yy + ib * QK_K + 128 * ip + il;
+    dst_t * y = yy + 128 * ip + il;
 
     const uint8_t * ql = ql_ptr + 64 * ip + il;
     const uint8_t   qh = *(qh_ptr + 32 * ip + il);

ggml/src/ggml-sycl/dpct/helper.hpp

@@ -3159,7 +3159,7 @@ namespace dpct
             _local_mem_size = local_mem_size;
 
 
-        };
+        }
         static inline std::mutex kernel_function_ptr_map_mutex;
 
       public:

ggml/src/ggml-sycl/element_wise.cpp

@@ -1169,7 +1169,7 @@ static void swiglu_oai_sycl(const T *       x,
                          });
 }
 
-void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+static void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     const ggml_tensor * src0 = dst->src[0];
     const ggml_tensor * src1 = dst->src[1];
     void * src0_d = src0->data;

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

@@ -832,6 +832,29 @@ static void flash_attn_combine_results(const float * __restrict__ VKQ_parts,
     dst[tid] = VKQ_numerator / VKQ_denominator;
 }
 
+// Carries the sub-group size (and grf_size_automatic when use_large_grf) as
+// kernel properties via get(properties_tag), in place of the
+// reqd_sub_group_size attribute and the deprecated parallel_for overload.
+template <typename KernelFunc, uint32_t Size, bool UseGrf>
+struct props_kernel {
+    KernelFunc kernel;
+
+    void operator()(sycl::nd_item<3> item) const {
+        kernel(item);
+    }
+
+    auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+        if constexpr (UseGrf) {
+            return sycl::ext::oneapi::experimental::properties{
+                sycl::ext::intel::experimental::grf_size_automatic,
+                sycl::ext::oneapi::experimental::sub_group_size<Size> };
+        } else {
+            return sycl::ext::oneapi::experimental::properties{
+                sycl::ext::oneapi::experimental::sub_group_size<Size> };
+        }
+    }
+};
+
 template <fattn_kernel_t fattn_kernel, int warp_size, bool use_large_grf = false>
 static void lauch_kernel(
     dpct::dim3 group_range,
@@ -881,7 +904,7 @@ static void lauch_kernel(
         static_cast<sycl::range<3>>(group_range * local_range),
         static_cast<sycl::range<3>>(local_range));
 
-    const auto kernel = [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(warp_size)]] {
+    const auto kernel = [=](sycl::nd_item<3> item_ct1) {
         GGML_UNUSED(item_ct1);
         fattn_kernel(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
                      max_bias, m0, m1, n_head_log2, logit_softcap, ne00,
@@ -892,14 +915,7 @@ static void lauch_kernel(
 
     q->submit([&](sycl::handler &cgh) {
         // grf_size_automatic lets the compiler use the large register file when the kernel needs it. grf_size<256> would say the same but is undefined on devices that do not have it.
-        if constexpr (use_large_grf) {
-            cgh.parallel_for(
-                rng,
-                sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::grf_size_automatic },
-                kernel);
-        } else {
-            cgh.parallel_for(rng, kernel);
-        }
+        cgh.parallel_for(rng, props_kernel<decltype(kernel), warp_size, use_large_grf>{kernel});
     });
 }
 

ggml/src/ggml-sycl/fattn-tile.cpp

@@ -54,6 +54,6 @@ void ggml_sycl_flash_attn_ext_tile(ggml_backend_sycl_context & ctx, ggml_tensor
         } break;
         default: {
             GGML_ABORT("Unsupported head size");
-        } break;
+        }
     }
 }

ggml/src/ggml-sycl/getrows.cpp

@@ -137,6 +137,47 @@ static void k_get_rows_float(
     dst_row[i00] = src0_row[i00];
 }
 
+// src0 is in reorder layout: all qs of the tensor, then all d
+template<int qk, int qr, dequantize_kernel_t_reorder dequantize_kernel, typename dst_t>
+static void k_get_rows_reorder(
+            const void * src0_qs, const void * src0_d, const int32_t * src1, dst_t * dst,
+            int64_t ne00, int64_t ne12,
+            size_t s1, size_t s2, size_t s3,
+            size_t s10, size_t s11, size_t s12,
+            const sycl::nd_item<3> &item_ct1) {
+
+    const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
+                     item_ct1.get_local_id(2)) *
+                    2;
+    const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
+                    item_ct1.get_local_id(1);
+    const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
+                     item_ct1.get_local_id(0)) /
+                    ne12;
+    const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
+                     item_ct1.get_local_id(0)) %
+                    ne12;
+
+    if (i00 >= ne00) {
+        return;
+    }
+
+    const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
+
+    dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
+
+    const int64_t ib = (int64_t) i01*(ne00/qk) + i00/qk; // block index in the whole tensor
+    const int iqs = (i00%qk)/qr;
+    const int iybs = i00 - i00%qk;
+    const int y_offset = qr == 1 ? 1 : qk/2;
+
+    dfloat2 v;
+    dequantize_kernel(src0_d, ib, src0_qs, ib*(qk/qr) + iqs, v);
+
+    dst_row[iybs + iqs + 0] = v.x();
+    dst_row[iybs + iqs + y_offset] = v.y();
+}
+
 template <int qk, int qr, dequantize_kernel_t dq>
 static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
                           ggml_tensor *dst, const void *src0_dd,
@@ -173,6 +214,93 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr
     GGML_UNUSED(ctx);
 }
 
+template <int qk, int qr, dequantize_kernel_t_reorder dq>
+static void get_rows_sycl_reorder(const ggml_tensor *src0, const ggml_tensor *src1,
+                                  ggml_tensor *dst, const void *src0_dd,
+                                  const int32_t *src1_dd, float *dst_dd,
+                                  queue_ptr stream) {
+
+    GGML_TENSOR_BINARY_OP_LOCALS
+
+    // reorder is done only on 2D weights
+    GGML_ASSERT(ne02 == 1 && ne03 == 1);
+    GGML_ASSERT(ne00 % 2 == 0);
+
+    const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
+    const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
+    const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
+
+    const size_t s1 = nb1 / ggml_element_size(dst);
+    const size_t s2 = nb2 / ggml_element_size(dst);
+    const size_t s3 = nb3 / ggml_element_size(dst);
+
+    const size_t s10 = nb10 / ggml_element_size(src1);
+    const size_t s11 = nb11 / ggml_element_size(src1);
+    const size_t s12 = nb12 / ggml_element_size(src1);
+
+    const void * src0_d = (const char *) src0_dd + ne00*ne01/qr;
+
+    stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
+                         [=](sycl::nd_item<3> item_ct1) {
+                             k_get_rows_reorder<qk, qr, dq>(
+                                 src0_dd, src0_d, src1_dd, dst_dd, ne00, ne12, s1, s2,
+                                 s3, s10, s11, s12, item_ct1);
+                         });
+
+    GGML_UNUSED(src0);
+}
+
+// K-quants in reorder layout: one work-group dequantizes one QK_K block of one row
+template <int dims, int nthreads, int nlocal, typename F>
+static void get_rows_sycl_reorder_k(const ggml_tensor *src0, const ggml_tensor *src1,
+                                    ggml_tensor *dst, const int32_t *src1_dd, float *dst_dd,
+                                    queue_ptr stream, F dequantize_block) {
+
+    GGML_TENSOR_BINARY_OP_LOCALS
+
+    // reorder is done only on 2D weights
+    GGML_ASSERT(ne02 == 1 && ne03 == 1);
+    GGML_ASSERT(ne00 % QK_K == 0);
+
+    const int64_t nblk = ne00 / QK_K;
+    const int64_t nb = nblk * ne01; // blocks in the whole tensor
+    const int64_t ngroups = ne10 * ne11 * ne12 * nblk;
+
+    const size_t s1 = nb1 / ggml_element_size(dst);
+    const size_t s2 = nb2 / ggml_element_size(dst);
+    const size_t s3 = nb3 / ggml_element_size(dst);
+
+    const size_t s10 = nb10 / ggml_element_size(src1);
+    const size_t s11 = nb11 / ggml_element_size(src1);
+    const size_t s12 = nb12 / ggml_element_size(src1);
+
+    stream->submit([&](sycl::handler & cgh) {
+        sycl::local_accessor<uint8_t, 1> scales_local(sycl::range<1>(nlocal), cgh);
+
+        auto kernel = [=](sycl::nd_item<dims> item_ct1) {
+            const int64_t g   = item_ct1.get_group(dims - 1);
+            const int64_t b   = g % nblk;
+            const int64_t r   = g / nblk;
+            const int64_t i10 = r % ne10;
+            const int64_t i11 = (r / ne10) % ne11;
+            const int64_t i12 = r / (ne10 * ne11);
+
+            const int64_t i01 = src1_dd[i10*s10 + i11*s11 + i12*s12];
+
+            float * y = dst_dd + i10*s1 + i11*s2 + i12*s3 + b*QK_K;
+            dequantize_block(y, i01*nblk + b, get_pointer(scales_local), item_ct1, nb);
+        };
+
+        if constexpr (dims == 1) {
+            cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(ngroups * nthreads), sycl::range<1>(nthreads)), kernel);
+        } else {
+            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, ngroups * nthreads), sycl::range<3>(1, 1, nthreads)), kernel);
+        }
+    });
+
+    GGML_UNUSED(src0);
+}
+
 template <int qk, int qr, dequantize_kernel_f32_t dq>
 static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
                               ggml_tensor *dst, const void *src0_dd,
@@ -332,6 +460,9 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type));
 
     const int32_t * src1_i32 = (const int32_t *) dst->src[1]->data;
+    // a weight shared with MUL_MAT (e.g. tied embeddings) can be in reorder layout
+    const bool src0_reordered = dst->src[0]->extra &&
+        ((const ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder;
     /* TODO: Refactor and remove duplicates */
     switch (dst->src[0]->type) {
         case GGML_TYPE_F16:
@@ -399,14 +530,37 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q2_K:
+            if (src0_reordered) {
+                const void * vx = dst->src[0]->data;
+                get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+                    [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+                        dequantize_block_q2_K_reorder(vx, y, ib, item_ct1, nb);
+                        GGML_UNUSED(scales_local);
+                    });
+                break;
+            }
             get_rows_sycl_f32<QK_K, 1, dequantize_q2_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q3_K:
+            if (src0_reordered) {
+                const void * vx = dst->src[0]->data;
+                get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+                    [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+                        dequantize_block_q3_K_reorder(vx, y, ib, item_ct1, nb);
+                        GGML_UNUSED(scales_local);
+                    });
+                break;
+            }
             get_rows_sycl<QK_K, 1, dequantize_q3_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q4_0:
+            if (src0_reordered) {
+                get_rows_sycl_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(dst->src[0], dst->src[1], dst, dst->src[0]->data,
+                src1_i32, (float *)dst->data, ctx.stream());
+                break;
+            }
             get_rows_sycl<QK4_0, QR4_0, dequantize_q4_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
@@ -415,6 +569,14 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q4_K:
+            if (src0_reordered) {
+                const void * vx = dst->src[0]->data;
+                get_rows_sycl_reorder_k<1, 32, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+                    [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) {
+                        dequantize_block_q4_K_reorder(vx, y, ib, scales_local, item_ct1, nb);
+                    });
+                break;
+            }
             get_rows_sycl_f32<QK_K, 1, dequantize_q4_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
@@ -427,14 +589,36 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q5_K:
+            if (src0_reordered) {
+                const void * vx = dst->src[0]->data;
+                get_rows_sycl_reorder_k<3, 64, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+                    [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+                        dequantize_block_q5_K_reorder(vx, y, ib, scales_local, item_ct1, nb);
+                    });
+                break;
+            }
             get_rows_sycl_f32<QK_K, 1, dequantize_q5_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q6_K:
+            if (src0_reordered) {
+                const void * vx = dst->src[0]->data;
+                get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+                    [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+                        dequantize_block_q6_K_reorder(vx, y, ib, item_ct1, nb);
+                        GGML_UNUSED(scales_local);
+                    });
+                break;
+            }
             get_rows_sycl<QK_K, 1, dequantize_q6_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
         case GGML_TYPE_Q8_0:
+            if (src0_reordered) {
+                get_rows_sycl_reorder<QK8_0, QR8_0, dequantize_q8_0_reorder>(dst->src[0], dst->src[1], dst, dst->src[0]->data,
+                src1_i32, (float *)dst->data, ctx.stream());
+                break;
+            }
             get_rows_sycl<QK8_0, QR8_0, dequantize_q8_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
             src1_i32, (float *)dst->data, ctx.stream());
             break;
@@ -468,6 +652,5 @@ void ggml_sycl_op_get_rows_back(ggml_backend_sycl_context & ctx, ggml_tensor * d
             break;
         default:
             GGML_ABORT("%s: unsupported src0 type: %s\n", __func__, ggml_type_name(src0->type));
-            break;
     }
 }

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

@@ -388,14 +388,14 @@ static ze_result_t init_zes() {
     return res;
 }
 
-ze_result_t get_zes_init_res() {
+static ze_result_t get_zes_init_res() {
     static ze_result_t zes_init_res = init_zes();
     GGML_SYCL_DEBUG("[SYCL] call %s: zesInit result: %d\n", __func__, (int) zes_init_res);
     return zes_init_res;
 }
 #endif
 
-void initialize_sycl_begining() {
+static void initialize_sycl_begining() {
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
     //must be called in initialization stage, before any other Level Zero API calls
     GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
@@ -1173,7 +1173,7 @@ static size_t ggml_backend_sycl_buffer_type_get_alignment(ggml_backend_buffer_ty
     GGML_UNUSED(buft);
 }
 
-bool is_bmg_g31_arch(int device) {
+static bool is_bmg_g31_arch(int device) {
     return ggml_sycl_info().devices[device].hw_info.arch == gpu_arch::intel_gpu_bmg_g31;
 }
 
@@ -6326,7 +6326,7 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
-bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
+static bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
     GGML_SYCL_DEBUG("[SYCL] [%s] g_ggml_sycl_get_mem_api=%d\n",
         __func__, g_ggml_sycl_get_mem_api);
 
@@ -6952,7 +6952,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                 default:
                     return false;
             }
-            break;
         case GGML_OP_MUL_MAT:
         case GGML_OP_MUL_MAT_ID:
             {
@@ -7045,7 +7044,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                            (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;

ggml/src/ggml-sycl/gla.cpp

@@ -1,6 +1,7 @@
 #include <sycl/sycl.hpp>
 
 #include "common.hpp"
+#include "gla.hpp"
 
 template <u_int HEAD_SIZE>
 static void gated_linear_attn_f32_kernel(const dpct::queue_ptr stream, u_int B, u_int T, u_int C, u_int H, float scale,

ggml/src/ggml-sycl/mem.cpp

@@ -27,7 +27,7 @@ const char * mem_api_int2str(int mem_api) {
 /*
 * Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail.
 */
-bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+static bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
     GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__);
 
     free_bytes  = 0;
@@ -94,7 +94,7 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
 }
 #endif
 
-bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+static bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
     GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using SYCL API.\n", __func__);
     total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
 

ggml/src/ggml-sycl/topk-moe.cpp

@@ -296,7 +296,6 @@ static void ggml_sycl_op_topk_moe(ggml_backend_sycl_context &     ctx,
             break;
         default:
             GGML_ASSERT(false && "fatal error");
-            break;
     }
 }
 

ggml/src/ggml-sycl/upscale.cpp

@@ -183,8 +183,8 @@ static void upscale_f32_bilinear_antialias(const float * src0,
 }
 
 namespace bicubic_interpolation {
-static float weight1(float x, const float &a) { return ((a + 2) * x - (a + 3)) * x * x + 1; };
-static float weight2(float x, const float &a) { return ((a * x - 5 * a) * x + 8 * a) * x - 4 * a; };
+static float weight1(float x, const float &a) { return ((a + 2) * x - (a + 3)) * x * x + 1; }
+static float weight2(float x, const float &a) { return ((a * x - 5 * a) * x + 8 * a) * x - 4 * a; }
 
 static float bicubic(float p0, float p1, float p2, float p3, float x, float a) {
     const float w0 = weight2(x + 1, a);
@@ -192,7 +192,7 @@ static float bicubic(float p0, float p1, float p2, float p3, float x, float a) {
     const float w2 = weight1(1 - x, a);
     const float w3 = weight2(2 - x, a);
     return p0 * w0 + p1 * w1 + p2 * w2 + p3 * w3;
-};
+}
 
 }
 
@@ -358,7 +358,7 @@ static void upscale_f32_bicubic_sycl(const float *   x,
     }
 }
 
-void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+static void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     const ggml_tensor * src0 = dst->src[0];
     const float * src0_d = (const float *)src0->data;
     float * dst_d = (float *)dst->data;

ggml/src/ggml-sycl/vecdotq.hpp

@@ -442,7 +442,7 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0> {
         }
 
         return vec_dot_q4_0_q8_1_impl(v, u, d, *q8_1_ds);
-    };
+    }
 };
 
 // Load four contiguous dwords per operand instead of loading each value separately.

gguf-py/gguf/constants.py

@@ -366,6 +366,7 @@ class Keys:
         ADD_BOS              = "tokenizer.ggml.add_bos_token"
         ADD_EOS              = "tokenizer.ggml.add_eos_token"
         ADD_SEP              = "tokenizer.ggml.add_sep_token"
+        DNA_AUTO_TAGS        = "tokenizer.ggml.dna_auto_tags"
         ADD_PREFIX           = "tokenizer.ggml.add_space_prefix"
         REMOVE_EXTRA_WS      = "tokenizer.ggml.remove_extra_whitespaces"
         PRECOMPILED_CHARSMAP = "tokenizer.ggml.precompiled_charsmap"

gguf-py/gguf/gguf_writer.py

@@ -1314,6 +1314,9 @@ class GGUFWriter:
     def add_add_sep_token(self, value: bool) -> None:
         self.add_bool(Keys.Tokenizer.ADD_SEP, value)
 
+    def add_dna_auto_tags(self, value: bool) -> None:
+        self.add_bool(Keys.Tokenizer.DNA_AUTO_TAGS, value)
+
     def add_add_space_prefix(self, value: bool) -> None:
         self.add_bool(Keys.Tokenizer.ADD_PREFIX, value)
 

scripts/connect-download.cmake

@@ -0,0 +1,86 @@
+# download the prebuilt llama-connect binary for the given platform
+# the release assets are made by: https://github.com/ggml-org/llama-connect/blob/master/.github/workflows/build.yml
+
+cmake_minimum_required(VERSION 3.18)
+
+set(REPO     "" CACHE STRING "GitHub repository to download from (owner/name)")
+set(VERSION  "" CACHE STRING "Release tag to download, or 'latest'")
+set(PLATFORM "" CACHE STRING "Platform suffix of the release asset (ex: linux-x64)")
+set(OUT_DIR  "" CACHE STRING "Directory to place the binary into")
+set(WORK_DIR "" CACHE STRING "Scratch directory used for the download")
+
+if (PLATFORM MATCHES "^win-")
+    set(ASSET  "llama-connect-${PLATFORM}.zip")
+    set(BINARY "llama-connect.exe")
+else()
+    set(ASSET  "llama-connect-${PLATFORM}.tar.gz")
+    set(BINARY "llama-connect")
+endif()
+
+if (VERSION STREQUAL "latest")
+    set(URL "https://github.com/${REPO}/releases/latest/download/${ASSET}")
+else()
+    set(URL "https://github.com/${REPO}/releases/download/${VERSION}/${ASSET}")
+endif()
+
+set(DST     "${OUT_DIR}/${BINARY}")
+set(ARCHIVE "${WORK_DIR}/${ASSET}")
+set(STAMP   "${WORK_DIR}/${BINARY}.url")
+
+# skip if already downloaded from the same URL
+if (EXISTS "${DST}" AND EXISTS "${STAMP}")
+    file(READ "${STAMP}" prev_url)
+    if (prev_url STREQUAL "${URL}")
+        return()
+    endif()
+endif()
+
+file(REMOVE "${STAMP}")
+file(MAKE_DIRECTORY "${WORK_DIR}")
+
+message(STATUS "llama-connect: downloading ${URL}")
+
+# retry a few times, this runs in CI where a transient failure must not break the build
+foreach (attempt RANGE 1 3)
+    file(DOWNLOAD "${URL}" "${ARCHIVE}" STATUS status TLS_VERIFY ON INACTIVITY_TIMEOUT 60)
+
+    list(GET status 0 code)
+    list(GET status 1 msg)
+
+    if (code EQUAL 0)
+        break()
+    endif()
+
+    message(STATUS "llama-connect: download attempt ${attempt} failed (${code}: ${msg})")
+    file(REMOVE "${ARCHIVE}")
+endforeach()
+
+if (NOT code EQUAL 0)
+    file(REMOVE "${ARCHIVE}")
+    message(FATAL_ERROR
+        "llama-connect: failed to download ${URL} (${code}: ${msg})\n"
+        "  use -DLLAMA_CONNECT_VERSION=<tag> to select another release, "
+        "or -DLLAMA_CONNECT=OFF to skip this tool")
+endif()
+
+set(EXTRACT_DIR "${WORK_DIR}/extract")
+
+file(REMOVE_RECURSE "${EXTRACT_DIR}")
+file(MAKE_DIRECTORY "${EXTRACT_DIR}")
+file(ARCHIVE_EXTRACT INPUT "${ARCHIVE}" DESTINATION "${EXTRACT_DIR}")
+
+if (NOT EXISTS "${EXTRACT_DIR}/${BINARY}")
+    message(FATAL_ERROR "llama-connect: ${ASSET} does not contain ${BINARY}")
+endif()
+
+file(MAKE_DIRECTORY "${OUT_DIR}")
+file(COPY "${EXTRACT_DIR}/${BINARY}" DESTINATION "${OUT_DIR}"
+    FILE_PERMISSIONS
+        OWNER_READ OWNER_WRITE OWNER_EXECUTE
+        GROUP_READ GROUP_EXECUTE
+        WORLD_READ WORLD_EXECUTE
+)
+
+file(WRITE "${STAMP}" "${URL}")
+
+message(STATUS "llama-connect: ready at ${DST}")

scripts/snapdragon/build.py

@@ -11,6 +11,9 @@ import platform
 import shutil
 import logging
 import glob
+import shlex
+import tarfile
+import tempfile
 
 from sdk import validate_windows_sdks
 
@@ -19,19 +22,70 @@ logger = logging.getLogger("build")
 
 def parse_target(target_str):
     if not target_str:
-        return None, None
-    if target_str.startswith("adb") or target_str.startswith("android"):
-        parts = target_str.split(":", 1)
-        serial = parts[1] if len(parts) > 1 else None
-        return "android", serial
-    elif target_str.startswith("lnx") or target_str.startswith("linux") or target_str.startswith("ubuntu"):
-        parts = target_str.split(":", 1)
-        host = parts[1] if len(parts) > 1 else None
-        return "linux", host
+        return None, None, None
+
+    parts = target_str.split(":")
+    target_name = parts[0]
+
+    if target_name in ("adb", "android"):
+        serial = ":".join(parts[1:]) if len(parts) > 1 else None
+        return "android", serial, "adb"
+    elif target_name in ("lnx", "linux", "ubuntu"):
+        if len(parts) > 1 and parts[1] == "adb":
+            serial = ":".join(parts[2:]) if len(parts) > 2 else None
+            return "linux", serial, "adb"
+        host = ":".join(parts[1:]) if len(parts) > 1 else None
+        return "linux", host, "ssh"
     elif target_str in ("wos", "windows"):
-        return "windows", None
+        return "windows", None, None
     else:
-        return None, None
+        return None, None, None
+
+
+def default_target_dir(target_type, target_transport, home_dir):
+    if target_type == "android":
+        return "/data/local/tmp/llama.cpp"
+    if target_type == "linux":
+        if home_dir:
+            return home_dir.rstrip("/") + "/llama.cpp"
+        if target_transport == "adb":
+            return "/home/llama.cpp"
+        return "~/llama.cpp"
+    return "~/llama.cpp"
+
+
+def normalize_target_dir(target_dir):
+    target_dir = target_dir.rstrip("/")
+    if not target_dir or target_dir == "/":
+        raise RuntimeError("target directory must not be the filesystem root")
+    return target_dir
+
+
+def adb_root(adb_cmd):
+    root_cmd = adb_cmd + ["root"]
+    logger.info(f"+ {' '.join(root_cmd)}")
+    res = subprocess.run(root_cmd)
+    if res.returncode != 0:
+        logger.error("ADB root failed.")
+        sys.exit(res.returncode)
+
+    wait_cmd = adb_cmd + ["wait-for-device"]
+    logger.info(f"+ {' '.join(wait_cmd)}")
+    res = subprocess.run(wait_cmd)
+    if res.returncode != 0:
+        logger.error("ADB wait-for-device failed.")
+        sys.exit(res.returncode)
+
+
+def make_package_archive(src_path):
+    fd, archive_path = tempfile.mkstemp(prefix="llama-cpp-", suffix=".tar.gz")
+    os.close(fd)
+    with tarfile.open(archive_path, "w:gz", dereference=False) as tar:
+        for item in os.listdir(src_path):
+            if item.startswith("."):
+                continue
+            tar.add(os.path.join(src_path, item), arcname=item, recursive=True)
+    return archive_path
 
 
 def get_uid_gid():
@@ -45,7 +99,7 @@ def main():
     parser = argparse.ArgumentParser(
         description="Build llama.cpp for Snapdragon using cross-compilation docker containers or natively."
     )
-    parser.add_argument("--target", default="android", help="Compilation target and deployment definition (e.g. android[:serial]/adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, windows/wos) (default: android)")
+    parser.add_argument("--target", default="android", help="Compilation target and deployment definition (e.g. android[:serial]/adb[:serial], ubuntu:adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, windows/wos) (default: android)")
     parser.add_argument("--build-dir", help="Build directory name (defaults to build-TARGET[-dbg], e.g. build-android)")
     parser.add_argument("--install-dir", help="Install directory name (defaults to pkg-TARGET[-dbg], e.g. pkg-android)")
     parser.add_argument("--jobs", "-j", type=int, help="Number of build jobs (defaults to CPU thread count)")
@@ -53,21 +107,22 @@ def main():
     parser.add_argument("--preset", help="Override the CMake preset to use")
     parser.add_argument("--debug", action="store_true", help="Build in debug mode (uses -debug presets instead of -release)")
     parser.add_argument("--hex-only", "--hexagon-only", action="store_true", help="Only build and push ggml-hexagon and libggml-htp libraries")
-    parser.add_argument("--hex-arch", choices=["v73", "v75", "v79", "v81", "73", "75", "79", "81"], help="Specific Hexagon DSP architecture to build with --hex-only (default: all)")
+    parser.add_argument("--hex-arch", choices=["v73", "v75", "v79", "v81", "v85", "73", "75", "79", "81", "85"], help="Specific Hexagon DSP architecture to build with --hex-only (default: all)")
 
     # Push options
     parser.add_argument("--push", action="store_true", help="Push built package to the target device via ADB or SSH/SCP")
-    parser.add_argument("--target-dir", help="Target directory on the device (default: /data/local/tmp/llama.cpp for Android, ~/llama.cpp for Linux)")
+    parser.add_argument("--target-dir", help="Target directory on the device (default: /data/local/tmp/llama.cpp for Android, ~/llama.cpp for Linux, /home/llama.cpp for Linux via ADB)")
+    parser.add_argument("--home-dir", help="Home directory for Linux targets when deriving the default target directory")
 
     # Toolchain options
-    parser.add_argument("--toolchain-version", default="v0.7", help="Docker toolchain image version/tag (default: v0.7)")
+    parser.add_argument("--toolchain-version", default="v0.8", help="Docker toolchain image version/tag (default: v0.8)")
     parser.add_argument("--toolchain-url", default="ghcr.io/snapdragon-toolchain", help="Docker toolchain registry URL/namespace (default: ghcr.io/snapdragon-toolchain)")
 
     args = parser.parse_args()
 
-    target_type, target_val = parse_target(args.target)
+    target_type, target_val, target_transport = parse_target(args.target)
     if not target_type:
-        logger.error(f"Error: Invalid target format '{args.target}'. Must be android[:serial]/adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, or windows/wos.")
+        logger.error(f"Error: Invalid target format '{args.target}'. Must be android[:serial]/adb[:serial], ubuntu:adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, or windows/wos.")
         sys.exit(1)
 
     if target_type == "windows":
@@ -140,7 +195,7 @@ def main():
         if hex_arch:
             hex_targets = ["ggml-hexagon", f"htp-{hex_arch}"]
         else:
-            hex_targets = ["ggml-hexagon", "htp-v73", "htp-v75", "htp-v79", "htp-v81"]
+            hex_targets = ["ggml-hexagon", "htp-v73", "htp-v75", "htp-v79", "htp-v81", "htp-v85"]
 
     if args.no_docker:
         # Native/local host build
@@ -218,8 +273,8 @@ def main():
         # Resolve target directory on device
         target_dir = args.target_dir
         if not target_dir:
-            target_dir = "/data/local/tmp/llama.cpp" if target_type == "android" else "~/llama.cpp"
-        target_dir = target_dir.rstrip("/")
+            target_dir = default_target_dir(target_type, target_transport, args.home_dir)
+        target_dir = normalize_target_dir(target_dir)
 
         sub_items = [item for item in os.listdir(src_path) if not item.startswith(".")]
 
@@ -235,14 +290,23 @@ def main():
                 logger.error(f"Error: no hexagon libraries found in {lib_dir} to push.")
                 sys.exit(1)
 
-        if target_type == "android":
+        if target_transport == "adb":
             adb_cmd = ["adb"]
             if target_val: # serial
                 adb_cmd += ["-s", target_val]
 
+            if target_type == "linux":
+                adb_root(adb_cmd)
+
             if args.hex_only:
-                logger.info("\nPushing Hexagon libraries to Android device via ADB...")
+                logger.info("\nPushing Hexagon libraries to target device via ADB...")
                 target_lib_dir = f"{target_dir}/lib"
+                mkdir_cmd = adb_cmd + ["shell", f"mkdir -p {shlex.quote(target_lib_dir)}"]
+                logger.info(f"+ {' '.join(mkdir_cmd)}")
+                res = subprocess.run(mkdir_cmd)
+                if res.returncode != 0:
+                    logger.error("ADB mkdir failed.")
+                    sys.exit(res.returncode)
                 for lib in hex_libs:
                     push_cmd = adb_cmd + ["push", lib, f"{target_lib_dir}/"]
                     logger.info(f"+ {' '.join(push_cmd)}")
@@ -251,6 +315,44 @@ def main():
                         logger.error("ADB push failed.")
                         sys.exit(res.returncode)
                 logger.info("Hexagon libraries deployed successfully!")
+            elif target_type == "linux":
+                logger.info("\nDeploying Linux package to target device via ADB tar...")
+                remote_archive = f"/tmp/llama.cpp-package-{os.getpid()}.tar.gz"
+                archive_path = make_package_archive(src_path)
+                try:
+                    mkdir_cmd = adb_cmd + ["shell", f"mkdir -p {shlex.quote(target_dir)}"]
+                    logger.info(f"+ {' '.join(mkdir_cmd)}")
+                    res = subprocess.run(mkdir_cmd)
+                    if res.returncode != 0:
+                        logger.error("ADB mkdir failed.")
+                        sys.exit(res.returncode)
+
+                    if sub_items:
+                        clean_paths = " ".join(shlex.quote(f"{target_dir}/{item}") for item in sub_items)
+                        clean_cmd = adb_cmd + ["shell", f"rm -rf {clean_paths}"]
+                        logger.info(f"+ {' '.join(clean_cmd)}")
+                        subprocess.run(clean_cmd)
+
+                    push_cmd = adb_cmd + ["push", archive_path, remote_archive]
+                    logger.info(f"+ {' '.join(push_cmd)}")
+                    res = subprocess.run(push_cmd)
+                    if res.returncode != 0:
+                        logger.error("ADB push failed.")
+                        sys.exit(res.returncode)
+
+                    extract_cmd = adb_cmd + ["shell", f"tar -xzf {shlex.quote(remote_archive)} -C {shlex.quote(target_dir)} && rm -f {shlex.quote(remote_archive)}"]
+                    logger.info(f"+ {' '.join(extract_cmd)}")
+                    res = subprocess.run(extract_cmd)
+                    if res.returncode != 0:
+                        logger.error("ADB untar failed.")
+                        sys.exit(res.returncode)
+
+                    chmod_cmd = adb_cmd + ["shell", f"chmod -R 755 {shlex.quote(target_dir)}/bin 2>/dev/null || true"]
+                    logger.info(f"+ {' '.join(chmod_cmd)}")
+                    subprocess.run(chmod_cmd)
+                finally:
+                    os.unlink(archive_path)
+                logger.info("ADB tar deploy completed successfully!")
             else:
                 # Clean stale package files on device
                 if sub_items:

scripts/snapdragon/ggml-hexagon-inspect.py

@@ -255,7 +255,7 @@ class HexagonToolchain:
         use_docker: bool = False,
         image_url: str = "ghcr.io/snapdragon-toolchain",
         image_name: str = "arm64-android",
-        image_ver: str = "v0.7",
+        image_ver: str = "v0.8",
     ):
         self.repo_root = repo_root
         self.image = f"{image_url}/{image_name}:{image_ver}"
@@ -322,15 +322,14 @@ class HexagonToolchain:
         if not self.docker_bin:
             raise RuntimeError("Docker is required but not installed or found on PATH.")
 
-        container_tools_dir = "/opt/hexagon/6.6.0.0/tools/HEXAGON_Tools/19.0.07/Tools/bin"
         if "objdump" in tool_name:
-            container_tool = f"{container_tools_dir}/hexagon-llvm-objdump"
+            tool_bin = "hexagon-llvm-objdump"
         elif "addr2line" in tool_name:
-            container_tool = f"{container_tools_dir}/hexagon-addr2line"
+            tool_bin = "hexagon-addr2line"
         elif "nm" in tool_name:
-            container_tool = f"{container_tools_dir}/llvm-nm"
+            tool_bin = "llvm-nm"
         else:
-            container_tool = f"{container_tools_dir}/{tool_name}"
+            tool_bin = tool_name
 
         # Translate file paths from host to /workspace
         translated_args = [translate_container_arg(arg, self.repo_root) for arg in args]
@@ -350,7 +349,9 @@ class HexagonToolchain:
         if platform.system() != "Windows":
             docker_cmd += ["-u", f"{os.getuid()}:{os.getgid()}"]
 
-        docker_cmd += [self.image, container_tool] + translated_args
+        default_tools_root = "/opt/hexagon/6.7.0.0/tools/HEXAGON_Tools/21.0.04"
+        sh_cmd = f'exec "${{HEXAGON_TOOLS_ROOT:-{default_tools_root}}}/Tools/bin/{tool_bin}" "$@"'
+        docker_cmd += [self.image, "sh", "-c", sh_cmd, "--"] + translated_args
 
         res = subprocess.run(docker_cmd, capture_output=True, text=True, input=stdin_data)
         if res.returncode != 0:
@@ -1457,8 +1458,8 @@ def main():
     )
     parser.add_argument(
         "--toolchain-version",
-        default="v0.7",
-        help="Docker toolchain tag (default: v0.7).",
+        default="v0.8",
+        help="Docker toolchain tag (default: v0.8).",
     )
     parser.add_argument(
         "--toolchain-url",

scripts/snapdragon/run.py

@@ -49,27 +49,46 @@ MANAGED_ENV_NAMES = (
     "GGML_OPENCL_ADRENO_XMEM_GEMM",
     "GGML_OPENCL_ADRENO_USE_LARGE_BUFFER",
     "GGML_SCHED_DEBUG",
+    "ADSP_LIBRARY_PATH",
+    "CDSP_LIBRARY_PATH",
+    "CDSP1_LIBRARY_PATH",
+    "CDSP2_LIBRARY_PATH",
+    "CDSP3_LIBRARY_PATH",
     "MTMD_BACKEND_DEVICE",
     "D",
     "DEVICE",
 )
 
 
+DSP_LIBRARY_ENV_NAMES = (
+    "ADSP_LIBRARY_PATH",
+    "CDSP_LIBRARY_PATH",
+    "CDSP1_LIBRARY_PATH",
+    "CDSP2_LIBRARY_PATH",
+    "CDSP3_LIBRARY_PATH",
+)
+
+
 def parse_target(target_str):
     if not target_str:
-        return None, None
-    if target_str.startswith("adb") or target_str.startswith("android"):
-        parts = target_str.split(":", 1)
-        serial = parts[1] if len(parts) > 1 else None
-        return "android", serial
-    elif target_str.startswith("lnx") or target_str.startswith("linux") or target_str.startswith("ubuntu"):
-        parts = target_str.split(":", 1)
-        host = parts[1] if len(parts) > 1 else None
-        return "linux", host
+        return None, None, None
+
+    parts = target_str.split(":")
+    target_name = parts[0]
+
+    if target_name in ("adb", "android"):
+        serial = ":".join(parts[1:]) if len(parts) > 1 else None
+        return "android", serial, "adb"
+    elif target_name in ("lnx", "linux", "ubuntu"):
+        if len(parts) > 1 and parts[1] == "adb":
+            serial = ":".join(parts[2:]) if len(parts) > 2 else None
+            return "linux", serial, "adb"
+        host = ":".join(parts[1:]) if len(parts) > 1 else None
+        return "linux", host, "ssh"
     elif target_str in ("wos", "windows"):
-        return "windows", None
+        return "windows", None, None
     else:
-        return None, None
+        return None, None, None
 
 
 def shlex_join(args_list):
@@ -79,6 +98,41 @@ def shlex_join(args_list):
     return " ".join(pipes.quote(x) for x in args_list)
 
 
+def default_target_dir(target_type, target_transport, home_dir):
+    if target_type == "android":
+        return "/data/local/tmp/llama.cpp"
+    if target_type == "linux":
+        if home_dir:
+            return home_dir.rstrip("/") + "/llama.cpp"
+        if target_transport == "adb":
+            return "/home/llama.cpp"
+        return "~/llama.cpp"
+    return "~/llama.cpp"
+
+
+def library_env_parts(lib_dir):
+    env_parts = [f"LD_LIBRARY_PATH={lib_dir}"]
+    for name in DSP_LIBRARY_ENV_NAMES:
+        env_parts.append(f"{name}={lib_dir}")
+    return env_parts
+
+
+def adb_root(adb_base):
+    root_cmd = adb_base + ["root"]
+    logger.info(f"+ {' '.join(root_cmd)}")
+    res = subprocess.run(root_cmd)
+    if res.returncode != 0:
+        logger.error("ADB root failed.")
+        sys.exit(res.returncode)
+
+    wait_cmd = adb_base + ["wait-for-device"]
+    logger.info(f"+ {' '.join(wait_cmd)}")
+    res = subprocess.run(wait_cmd)
+    if res.returncode != 0:
+        logger.error("ADB wait-for-device failed.")
+        sys.exit(res.returncode)
+
+
 def split_device_list(devices):
     parts = []
     curr = []
@@ -144,8 +198,9 @@ def main():
     parser = argparse.ArgumentParser(
         description="Unified runner for llama.cpp tools on Snapdragon (natively, via ADB, or via SSH)."
     )
-    parser.add_argument("--target", help="Execution target (e.g. android[:serial]/adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, windows/wos) (default: local run)")
-    parser.add_argument("--target-dir", help="Target directory on the device (default: /data/local/tmp/llama.cpp for Android, ~/llama.cpp for Linux)")
+    parser.add_argument("--target", help="Execution target (e.g. android[:serial]/adb[:serial], ubuntu:adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, windows/wos) (default: local run)")
+    parser.add_argument("--target-dir", help="Target directory on the device (default: /data/local/tmp/llama.cpp for Android, ~/llama.cpp for Linux, /home/llama.cpp for Linux via ADB)")
+    parser.add_argument("--home-dir", help="Home directory for Linux targets when deriving the default target directory")
     parser.add_argument("--install-dir", help="Install directory name (defaults to pkg-TARGET or pkg-TARGET-dbg prefix based on target)")
     parser.add_argument("--debug", action="store_true", help="Use debug build (defaults to pkg-TARGET-dbg folder)")
     parser.add_argument("--devices", "--device", "-d", help="Select execution devices (split into NPU and OpenCL GPUs automatically, default: HTP0)")
@@ -201,11 +256,12 @@ def main():
 
     target_type = None
     target_val = None
+    target_transport = None
     target_prefix = None
     if args.target:
-        target_type, target_val = parse_target(args.target)
+        target_type, target_val, target_transport = parse_target(args.target)
         if not target_type:
-            logger.error(f"Error: Invalid target format '{args.target}'. Must be android[:serial]/adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, or windows/wos.")
+            logger.error(f"Error: Invalid target format '{args.target}'. Must be android[:serial]/adb[:serial], ubuntu:adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, or windows/wos.")
             sys.exit(1)
         target_prefix = args.target.split(":", 1)[0]
 
@@ -435,18 +491,19 @@ def main():
     # Resolve target directory on device
     target_dir = args.target_dir
     if not target_dir:
-        target_dir = "/data/local/tmp/llama.cpp" if target_type == "android" else "~/llama.cpp"
+        target_dir = default_target_dir(target_type, target_transport, args.home_dir)
 
-    if target_type == "android":
-        # Run via ADB
+    if target_transport == "adb":
         adb_base = ["adb"]
         if target_val: # serial
             adb_base += ["-s", target_val]
 
-        env_parts = [
-            "LD_LIBRARY_PATH=./lib",
-            "ADSP_LIBRARY_PATH=./lib"
-        ]
+        if target_type == "linux":
+            adb_root(adb_base)
+
+        env_parts = library_env_parts("./lib")
+        if target_type == "linux":
+            env_parts.insert(0, f"HOME={args.home_dir or '/home'}")
         for k, v in env_vars.items():
             env_parts.append(f"{k}={v}")
         env_str = " ".join(env_parts)
@@ -466,10 +523,7 @@ def main():
             sys.exit(1)
 
         # Linux remote run via SSH
-        env_parts = [
-            "LD_LIBRARY_PATH=./lib",
-            "ADSP_LIBRARY_PATH=./lib"
-        ]
+        env_parts = library_env_parts("./lib")
         for k, v in env_vars.items():
             env_parts.append(f"{k}={v}")
         env_str = " ".join(env_parts)
@@ -490,7 +544,6 @@ def main():
         # Run locally
         local_env = os.environ.copy()
         lib_dir = os.path.normpath(os.path.join(package_path, "lib"))
-        local_env["ADSP_LIBRARY_PATH"] = lib_dir
         if platform.system() == "Windows":
             local_env["PATH"] = lib_dir + os.path.pathsep + local_env.get("PATH", "")
         else:
@@ -498,6 +551,8 @@ def main():
 
         for k in MANAGED_ENV_NAMES:
             local_env.pop(k, None)
+        for name in DSP_LIBRARY_ENV_NAMES:
+            local_env[name] = lib_dir
         for k, v in env_vars.items():
             local_env[k] = v
 

scripts/snapdragon/sdk.py

@@ -6,7 +6,7 @@ SDK_CONFIGS = (
     {
         "name": "Hexagon SDK",
         "repo": "snapdragon-toolchain/hexagon-sdk",
-        "default_version": "6.6.0.0",
+        "default_version": "6.7.0.0",
         "parent_dir": "Hexagon_SDK",
         "archive_prefix": "hexagon-sdk-v",
         "markers": ("hexagon_sdk.json",),

skills/slice-model/slice_model.py

@@ -223,6 +223,7 @@ def patch_config(cfg: dict, families: dict[str, Family], new_vocab: int | None,
                 lfam = next((f for f in families.values() if f.kind == "layer"), None)
                 if lfam is None:
                     continue
+                assert lfam.kept is not None
                 d[key] = sum(1 for i in lfam.kept if i < val)
             elif "expert" in key and fam is not None and fam.kind == "expert" and not any(x in key for x in EXPERT_KEY_EXCLUDE):
                 d[key] = fam.new_count
@@ -236,6 +237,7 @@ def patch_config(cfg: dict, families: dict[str, Family], new_vocab: int | None,
             # per-layer lists, e.g. layer_types
             fam = by_card.get(len(val))
             if fam is not None and fam.kind == "layer" and len(val) in (fam.card, fam.regular_total):
+                assert fam.kept is not None
                 d[key] = [val[i] for i in fam.kept if i < len(val)]
                 log.append(f"{key}: list[{len(val)}] -> list[{len(d[key])}]")
     return log

src/llama-arch.cpp

@@ -430,6 +430,7 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
     { LLM_KV_TOKENIZER_ADD_BOS,                  "tokenizer.ggml.add_bos_token"            },
     { LLM_KV_TOKENIZER_ADD_EOS,                  "tokenizer.ggml.add_eos_token"            },
     { LLM_KV_TOKENIZER_ADD_SEP,                  "tokenizer.ggml.add_sep_token"            },
+    { LLM_KV_TOKENIZER_DNA_AUTO_TAGS,            "tokenizer.ggml.dna_auto_tags"            },
     { LLM_KV_TOKENIZER_ADD_PREFIX,               "tokenizer.ggml.add_space_prefix"         },
     { LLM_KV_TOKENIZER_REMOVE_EXTRA_WS,          "tokenizer.ggml.remove_extra_whitespaces" },
     { LLM_KV_TOKENIZER_PRECOMPILED_CHARSMAP,     "tokenizer.ggml.precompiled_charsmap"     },

src/llama-arch.h

@@ -394,6 +394,7 @@ enum llm_kv {
     LLM_KV_TOKENIZER_ADD_BOS,
     LLM_KV_TOKENIZER_ADD_EOS,
     LLM_KV_TOKENIZER_ADD_SEP,
+    LLM_KV_TOKENIZER_DNA_AUTO_TAGS,
     LLM_KV_TOKENIZER_ADD_PREFIX,
     LLM_KV_TOKENIZER_REMOVE_EXTRA_WS,
     LLM_KV_TOKENIZER_PRECOMPILED_CHARSMAP,

src/llama-vocab.cpp

@@ -1776,16 +1776,15 @@ private:
 // text at load
 static const std::string dna_kmer_marker = "\xee\x80\x80";
 
+// HybridDNATokenizer (HuggingFaceBio/Carbon-*): text inside <dna>...</dna> is cut into fixed 6-mers, everything else goes through the base BPE
+// Mirrors HybridDNATokenizer.encode, with optional tagging of untagged DNA (auto_dna_tags, Carbon-14B) and tail="pad_n" vocab, whose dedicated tokens pad a trailing partial k-mer with N instead of A ("CT" -> "CTNNNN")
 struct llm_tokenizer_hybriddna_session : llm_tokenizer_bpe_session {
     llm_tokenizer_hybriddna_session(const llama_vocab & vocab, const llm_tokenizer_bpe & tokenizer) : llm_tokenizer_bpe_session{vocab, tokenizer}, vocab{vocab} {}
 
     void tokenize(const std::string & text, std::vector<llama_token> & output) override {
-        static const std::string open_tag  = "<dna>";
-        static const std::string close_tag = "</dna>";
-
-        const auto dna_begin_id = vocab.text_to_token(open_tag);
-        const auto dna_end_id   = vocab.text_to_token(close_tag);
-        const auto dna_oov_id   = vocab.text_to_token("<oov>");
+        dna_begin_id = vocab.text_to_token(open_tag);
+        dna_end_id   = vocab.text_to_token(close_tag);
+        dna_oov_id   = vocab.text_to_token("<oov>");
 
         // Fall back to plain BPE if the DNA pieces aren't in the vocab.
         if (dna_begin_id == LLAMA_TOKEN_NULL || dna_end_id == LLAMA_TOKEN_NULL || dna_oov_id == LLAMA_TOKEN_NULL) {
@@ -1793,63 +1792,187 @@ struct llm_tokenizer_hybriddna_session : llm_tokenizer_bpe_session {
             return;
         }
 
-        const size_t k = 6;
-        size_t pos = 0;
+        pad_char = vocab.text_to_token("A" + std::string(k - 1, 'N') + dna_kmer_marker) != LLAMA_TOKEN_NULL ? 'N' : 'A';
 
+        // split on explicit tags as HybridDNATokenizer._split_by_dna_tags does; only the text outside them is auto-tagged
+        size_t pos = 0;
         while (pos < text.size()) {
             const size_t start = text.find(open_tag, pos);
-            if (start == std::string::npos) {
-                if (pos < text.size()) {
-                    llm_tokenizer_bpe_session::tokenize(text.substr(pos), output);
-                }
+            const size_t end   = text.find(close_tag, pos);
+
+            if (start == std::string::npos && end == std::string::npos) {
+                tokenize_text(text.substr(pos), output);
                 break;
             }
-            if (start > pos) {
-                llm_tokenizer_bpe_session::tokenize(text.substr(pos, start - pos), output);
+            if (start == std::string::npos || end < start) {
+                // as in _split_by_dna_tags, text before a </dna> without a matching <dna> is DNA (no begin token)
+                tokenize_dna(text.substr(pos, end - pos), false, true, output);
+                pos = end + close_tag.size();
+                continue;
             }
-            output.push_back(dna_begin_id);
+
+            tokenize_text(text.substr(pos, start - pos), output);
 
             const size_t content_start = start + open_tag.size();
-            const size_t end           = text.find(close_tag, content_start);
-            const size_t content_end   = (end == std::string::npos) ? text.size() : end;
-
-            emit_dna_kmers(text.substr(content_start, content_end - content_start), k, dna_oov_id, output);
-
             if (end == std::string::npos) {
+                // an unterminated <dna> runs to the end of the text (no end token)
+                tokenize_dna(text.substr(content_start), true, false, output);
                 break;
             }
-            output.push_back(dna_end_id);
+            tokenize_dna(text.substr(content_start, end - content_start), true, true, output);
             pos = end + close_tag.size();
         }
     }
 
 private:
-    void emit_dna_kmers(const std::string & raw, size_t k, llama_token oov_id, std::vector<llama_token> & output) {
-        std::string seq = raw;
-        for (char & c : seq) {
-            if (c >= 'a' && c <= 'z') {
-                c = char(c - 32);
+    static bool is_letter(char c) { return (c >= 'A' && c <= 'Z') || (c >= 'a' && c <= 'z'); }
+    static bool is_acgt  (char c) { return c == 'A' || c == 'C' || c == 'G' || c == 'T'; }
+    static char to_upper (char c) { return c >= 'a' && c <= 'z' ? char(c - 32) : c; }
+
+    // IUPAC nucleotide codes, either case (HybridDNATokenizer.IUPAC_RUN_RE)
+    static bool is_iupac(char c) {
+        c = to_upper(c);
+        return std::string("ACGTRYSWKMBDHVN").find(c) != std::string::npos;
+    }
+
+    // HybridDNATokenizer._normalize_spaced_dna: join displays like "ACGTA CGTAC ..." when they hold at least twelve bases
+    // (at least six chunks of 1-5 uppercase ACGT separated by spaces or tabs, not touching a letter)
+    static std::string normalize_spaced_dna(const std::string & text) {
+        const auto is_blank = [](char c) { return c == ' ' || c == '\t'; };
+        std::string out;
+        size_t p = 0;
+        while (p < text.size()) {
+            if (is_acgt(text[p]) && (p == 0 || !is_letter(text[p - 1]))) {
+                size_t match_end = std::string::npos;
+                size_t chunk = p;
+                for (int n_chunks = 1; ; n_chunks++) {
+                    size_t chunk_end = chunk;
+                    while (chunk_end < text.size() && chunk_end - chunk < 5 && is_acgt(text[chunk_end])) {
+                        chunk_end++;
+                    }
+                    if (chunk_end == chunk) {
+                        break;
+                    }
+                    if (n_chunks >= 6 && (chunk_end == text.size() || !is_letter(text[chunk_end]))) {
+                        match_end = chunk_end; // the regex is greedy, so the last valid end wins
+                    }
+                    size_t next = chunk_end;
+                    while (next < text.size() && is_blank(text[next])) {
+                        next++;
+                    }
+                    if (next == chunk_end) {
+                        break;
+                    }
+                    chunk = next;
+                }
+                if (match_end != std::string::npos) {
+                    std::string seq;
+                    for (size_t i = p; i < match_end; i++) {
+                        if (!is_blank(text[i])) {
+                            seq += text[i];
+                        }
+                    }
+                    out += seq.size() >= 12 ? seq : text.substr(p, match_end - p);
+                    p = match_end;
+                    continue;
+                }
             }
+            out += text[p++];
+        }
+        return out;
+    }
+
+    // text outside explicit tags: with auto_dna_tags, tag DNA runs as HybridDNATokenizer._auto_tag_text_segment does
+    // (pure ACGT runs of at least 6 bases, or runs of at least 12 with at most 20% IUPAC ambiguity codes)
+    void tokenize_text(const std::string & raw, std::vector<llama_token> & output) {
+        auto bpe = [&](const std::string & piece) {
+            if (!piece.empty()) {
+                llm_tokenizer_bpe_session::tokenize(piece, output);
+            }
+        };
+        if (!vocab.get_dna_auto_tags()) {
+            bpe(raw);
+            return;
         }
 
-        // k-mers carry the reserved marker suffix; a non-ACGT k-mer simply
-        // isn't in the vocab and falls back to <oov>
-        auto kmer_token = [&](const std::string & kmer) {
+        const std::string text = normalize_spaced_dna(raw);
+        size_t prev = 0;
+        size_t i = 0;
+        while (i < text.size()) {
+            if (!is_iupac(text[i])) {
+                i++;
+                continue;
+            }
+            size_t j = i;
+            size_t n_ambiguous = 0;
+            for (; j < text.size() && is_iupac(text[j]); j++) {
+                n_ambiguous += !is_acgt(to_upper(text[j]));
+            }
+            const size_t len = j - i;
+            if ((n_ambiguous == 0 && len >= 6) || (n_ambiguous > 0 && len >= 12 && 5 * n_ambiguous <= len)) {
+                bpe(text.substr(prev, i - prev));
+                tokenize_dna(text.substr(i, len), true, true, output);
+                prev = j;
+            }
+            i = j;
+        }
+        bpe(text.substr(prev));
+    }
+
+    // HybridDNATokenizer._process_dna_sequence: strip, uppercase, then one token per k bases
+    // a k-mer with anything but ACGT is <oov>, and a trailing partial k-mer is right-padded with pad_char
+    void tokenize_dna(const std::string & content, bool has_start, bool has_end, std::vector<llama_token> & output) {
+        if (has_start) {
+            output.push_back(dna_begin_id);
+        }
+
+        auto cpts = unicode_cpts_from_utf8(content);
+        size_t first = 0;
+        size_t last  = cpts.size();
+        while (first < last && unicode_cpt_flags_from_cpt(cpts[first]).is_whitespace) {
+            first++;
+        }
+        while (last > first && unicode_cpt_flags_from_cpt(cpts[last - 1]).is_whitespace) {
+            last--;
+        }
+
+        // k-mers carry the reserved marker suffix, so they never resolve to a base-vocab token
+        auto kmer_token = [&](size_t from, size_t to) {
+            std::string kmer;
+            for (size_t i = from; i < to; i++) {
+                const uint32_t c = cpts[i] < 128 ? (uint32_t) to_upper((char) cpts[i]) : cpts[i];
+                if (c > 127 || !is_acgt((char) c)) {
+                    return dna_oov_id;
+                }
+                kmer += (char) c;
+            }
+            kmer.append(k - kmer.size(), pad_char);
             const auto tok = vocab.text_to_token(kmer + dna_kmer_marker);
-            return tok != LLAMA_TOKEN_NULL ? tok : oov_id;
+            return tok != LLAMA_TOKEN_NULL ? tok : dna_oov_id;
         };
 
-        size_t i = 0;
-        for (; i + k <= seq.size(); i += k) {
-            output.push_back(kmer_token(seq.substr(i, k)));
+        size_t i = first;
+        for (; i + k <= last; i += k) {
+            output.push_back(kmer_token(i, i + k));
         }
-        if (i < seq.size()) {
-            std::string kmer = seq.substr(i);
-            kmer.append(k - kmer.size(), 'A');
-            output.push_back(kmer_token(kmer));
+        if (i < last) {
+            output.push_back(kmer_token(i, last));
+        }
+
+        if (has_end) {
+            output.push_back(dna_end_id);
         }
     }
 
+    const std::string open_tag  = "<dna>";
+    const std::string close_tag = "</dna>";
+    const size_t k = 6;
+
+    llama_token dna_begin_id = LLAMA_TOKEN_NULL;
+    llama_token dna_end_id   = LLAMA_TOKEN_NULL;
+    llama_token dna_oov_id   = LLAMA_TOKEN_NULL;
+    char pad_char = 'A';
+
     const llama_vocab & vocab;
 };
 
@@ -1957,6 +2080,7 @@ struct llama_vocab::impl {
     bool add_bos                    = false;
     bool add_eos                    = false;
     bool add_sep                    = false;
+    bool dna_auto_tags              = false; // hybriddna: tag untagged DNA runs before tokenizing
     bool ignore_merges              = false;
     bool clean_spaces               = false;  // clean_up_tokenization_spaces
     bool remove_extra_whitespaces   = false;
@@ -2689,6 +2813,12 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
                 if (text.size() > dna_kmer_marker.size()
                         && text.compare(text.size() - dna_kmer_marker.size(), dna_kmer_marker.size(), dna_kmer_marker) == 0) {
                     text.erase(text.size() - dna_kmer_marker.size());
+                    // tail="pad_n": "CTNNNN" detokenizes to its real bases "CT"
+                    // full k-mers are ACGT only, so only N-padded tail tokens end in N
+                    const size_t last_base = text.find_last_not_of('N');
+                    if (last_base != std::string::npos) {
+                        text.erase(last_base + 1);
+                    }
                 }
             }
         }
@@ -2781,6 +2911,9 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
             if (ml.get_key(LLM_KV_TOKENIZER_ADD_SEP, temp, false)) {
                 add_sep = temp;
             }
+            if (ml.get_key(LLM_KV_TOKENIZER_DNA_AUTO_TAGS, temp, false)) {
+                dna_auto_tags = temp;
+            }
 
             // workaround for Gemma 4
             // ref: https://github.com/ggml-org/llama.cpp/pull/21500
@@ -4312,6 +4445,10 @@ bool llama_vocab::get_add_sep() const {
     return pimpl->add_sep;
 }
 
+bool llama_vocab::get_dna_auto_tags() const {
+    return pimpl->dna_auto_tags;
+}
+
 bool llama_vocab::get_ignore_merges() const {
     return pimpl->ignore_merges;
 }

src/llama-vocab.h

@@ -149,6 +149,7 @@ struct llama_vocab {
     bool get_add_bos                   () const;
     bool get_add_eos                   () const;
     bool get_add_sep                   () const;
+    bool get_dna_auto_tags             () const;
     bool get_ignore_merges             () const;
     bool get_clean_spaces              () const;
     bool get_remove_extra_whitespaces  () const;

tests/test-backend-ops.cpp

@@ -2453,6 +2453,63 @@ struct test_get_rows : public test_case {
     }
 };
 
+// GGML_OP_GET_ROWS on a weight also used by GGML_OP_MUL_MAT (tied embeddings)
+// the mul_mat runs first, so a backend that changes the weight layout must still return the right rows
+struct test_get_rows_after_mul_mat : public test_case {
+    const ggml_type type;
+    const int n; // cols
+    const int m; // rows
+    const int r; // rows to get
+
+    std::string vars() override {
+        return VARS_TO_STR4(type, n, m, r);
+    }
+
+    test_get_rows_after_mul_mat(ggml_type type = GGML_TYPE_Q8_0, int n = 256, int m = 16, int r = 4)
+        : type(type), n(n), m(m), r(r) {}
+
+    double max_nmse_err() override {
+        return 5e-4; // same as test_mul_mat
+    }
+
+    ggml_tensor * build_graph(ggml_context * ctx) override {
+        ggml_tensor * w = ggml_new_tensor_2d(ctx, type, n, m);
+        ggml_set_name(w, "w");
+
+        ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, n, 1);
+        ggml_set_name(x, "x");
+
+        ggml_tensor * rows = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, r);
+        ggml_set_name(rows, "rows");
+
+        ggml_tensor * mm = ggml_mul_mat(ctx, w, x);
+        ggml_set_name(mm, "mm");
+
+        ggml_tensor * gr = ggml_get_rows(ctx, w, rows);
+        ggml_set_name(gr, "gr");
+
+        // repeat(mm) is the first source, so mul_mat is computed before get_rows
+        ggml_tensor * out = ggml_add(ctx, ggml_repeat(ctx, mm, gr), gr);
+        ggml_set_name(out, "out");
+
+        return out;
+    }
+
+    void initialize_tensors(ggml_context * ctx) override {
+        for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+            if (t->type == GGML_TYPE_I32) {
+                std::vector<int> data(ggml_nelements(t));
+                for (size_t i = 0; i < data.size(); i++) {
+                    data[i] = rand() % m;
+                }
+                ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(int));
+            } else {
+                init_tensor_uniform(t);
+            }
+        }
+    }
+};
+
 // GGML_OP_GET_ROWS_BACK
 struct test_get_rows_back : public test_case {
     const ggml_type type;
@@ -9828,6 +9885,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     }
 
     test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false));
+    for (ggml_type type : {GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K}) {
+        test_cases.emplace_back(new test_get_rows_after_mul_mat(type));
+    }
     for (ggml_type type : all_types) {
         for (int b : {1, 7}) {
             for (bool v : {false, true}) {

tests/test-jinja.cpp

@@ -747,6 +747,62 @@ static void test_filters(testing & t) {
         "321"
     );
 
+    test_template(t, "unique",
+        "{{ items|unique|join(',') }}",
+        {{"items", json::array({3, 1, 3, 2, 1})}},
+        "3,1,2"
+    );
+
+    test_template(t, "unique case sensitive",
+        "{{ items|unique(case_sensitive=true)|join(',') }}",
+        {{"items", json::array({"a", "A", "b", "B", "a"})}},
+        "a,A,b,B"
+    );
+
+    test_template(t, "unique case sensitive positional",
+        "{{ items|unique(true)|list|length }}",
+        {{"items", json::array({"a", "A", "a"})}},
+        "2"
+    );
+
+    test_template(t, "unique with attribute",
+        "{{ items|unique(attribute='id')|join(attribute='v') }}",
+        {{"items", json::array({
+            json({{"id", "x"}, {"v", 1}}),
+            json({{"id", "y"}, {"v", 2}}),
+            json({{"id", "x"}, {"v", 3}}),
+        })}},
+        "12"
+    );
+
+    test_template(t, "unique with attribute case sensitive",
+        "{{ items|unique(case_sensitive=true, attribute='id')|join(attribute='v') }}",
+        {{"items", json::array({
+            json({{"id", "x"}, {"v", 1}}),
+            json({{"id", "X"}, {"v", 2}}),
+            json({{"id", "x"}, {"v", 3}}),
+        })}},
+        "12"
+    );
+
+    test_template(t, "unique mixed types",
+        "{{ items|unique|list|length }}",
+        {{"items", json::array({1, 1.0, 2, "a", "a"})}},
+        "3"
+    );
+
+    test_template(t, "unique empty",
+        "{{ items|unique|list|length }}",
+        {{"items", json::array()}},
+        "0"
+    );
+
+    test_template(t, "unique length comparison",
+        "{{ (items|unique(case_sensitive=true)|list|length) != (items|length) }}",
+        {{"items", json::array({"call_1", "call_2", "call_1"})}},
+        "True"
+    );
+
     test_template(t, "sort",
         "{% for i in items|sort %}{{ i }}{% endfor %}",
         {{"items", json::array({3, 1, 2})}},

tests/test-llama-archs.cpp

@@ -85,10 +85,18 @@ static void fill_resid_gate(std::vector<float> & dst, const char * name, const t
     }
 }
 
+// a random router bias can tie within the noise between the backends, a ramp keeps the top-k selection the same
+static void fill_exp_probs_b(std::vector<float> & dst, const char *, const tensor_data_params &) {
+    for (size_t i = 0; i < dst.size(); i++) {
+        dst[i] = 2.0f*i;
+    }
+}
+
 // a tensor whose values hold a property of their own, matched on a substring of its name, first match wins
 static const struct { const char * match; tensor_filler fill; } TENSOR_FILLERS[] = {
-    { "ssm_a",   fill_ssm_a      },
-    { "_resid_", fill_resid_gate },
+    { "ssm_a",       fill_ssm_a       },
+    { "_resid_",     fill_resid_gate  },
+    { "exp_probs_b", fill_exp_probs_b },
 };
 
 static void set_tensor_data(struct ggml_tensor * tensor, void * userdata) {
@@ -934,6 +942,18 @@ static int save_models(const std::string & arch_filter, const size_t seed, const
     return 0;
 }
 
+// the CPU and the GPU of an Apple SoC carry the same description, the name tells them apart (CPU, MTL0)
+static std::string get_device_label(ggml_backend_dev_t dev) {
+    const char * desc = ggml_backend_dev_description(dev);
+    const char * name = ggml_backend_dev_name(dev);
+
+    if (strcmp(desc, name) == 0) {
+        return desc;
+    }
+
+    return std::string(desc) + " (" + name + ")";
+}
+
 static int test_backends(const std::string & arch_filter, const size_t seed, const float stdev, const int verbosity, const char * target_backend) {
     struct user_data_t {
         struct {
@@ -986,7 +1006,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
                 if (target_backend != nullptr && strcmp(target_backend, ggml_backend_dev_name(dev)) != 0) {
                     continue;
                 }
-                dev_configs.emplace_back(std::vector<ggml_backend_dev_t>{dev}, ggml_backend_dev_description(dev), LLAMA_SPLIT_MODE_LAYER);
+                dev_configs.emplace_back(std::vector<ggml_backend_dev_t>{dev}, get_device_label(dev), LLAMA_SPLIT_MODE_LAYER);
                 max_device_label_length = std::max(max_device_label_length, dev_configs.back().label.length());
 
                 // cpu-based devices cannot be used in tensor split mode

tests/test-mtmd-impl.cpp

@@ -66,6 +66,31 @@ MAKE_TEST(test_image_preprocessor_longest_edge) {
     }
 }
 
+MAKE_TEST(test_image_preprocessor_dyn_size) {
+    // { patch_size, image size, expected size }
+    // a side that is an even number plus .5 merged patches rounds down, like Python round() in smart_resize
+    const std::vector<std::tuple<int, clip_image_size, clip_image_size>> cases = {
+        { 16, {  768, 400 }, {  768, 384 } },
+        { 16, {  800, 336 }, {  800, 320 } },
+        { 16, {  400, 400 }, {  384, 384 } },
+        { 16, {  640, 240 }, {  640, 256 } },
+        { 14, {  966, 560 }, {  952, 560 } },
+    };
+
+    for (const auto & [patch_size, input, expected] : cases) {
+        clip_hparams hparams;
+        hparams.patch_size = patch_size;
+        hparams.n_merge = 2;
+        hparams.image_round_half_even = true;
+        hparams.set_limit_image_tokens(8, 4096);
+
+        const clip_image_size actual = mtmd_image_preprocessor_dyn_size::get_target_size(hparams, input);
+        const std::string name = std::to_string(input.width) + "x" + std::to_string(input.height);
+        t.assert_equal(name + " width", expected.width, actual.width);
+        t.assert_equal(name + " height", expected.height, actual.height);
+    }
+}
+
 MAKE_TEST(test_image_preprocessor_lfm2) {
     clip_hparams hparams;
     hparams.patch_size = 16;

tools/CMakeLists.txt

@@ -32,6 +32,9 @@ else()
     if (GGML_RPC)
         add_subdirectory(rpc)
     endif()
+    if (LLAMA_CONNECT)
+        add_subdirectory(connect)
+    endif()
     if (NOT GGML_BACKEND_DL AND GGML_CPU)
         # these tools use backends directly (no dynamic loading) and depend on CPU backend symbols
         add_subdirectory(cvector-generator)

tools/cli/README.md

@@ -177,7 +177,7 @@
 | `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
 | `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
 | `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
-| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
+| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
 | `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
 | `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
 | `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |

tools/completion/README.md

@@ -255,7 +255,7 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1
 | `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
 | `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
 | `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
-| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
+| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
 | `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
 | `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
 | `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |

tools/connect/CMakeLists.txt

@@ -0,0 +1,85 @@
+# llama-connect is developed in a separate repo, nothing is compiled here
+# we only download the prebuilt binary, see scripts/connect-download.cmake
+
+set(TARGET llama-connect)
+
+set(LLAMA_CONNECT_REPO    "ggml-org/llama-connect" CACHE STRING "llama-connect: GitHub repository to download from")
+set(LLAMA_CONNECT_VERSION "latest"                 CACHE STRING "llama-connect: release tag to download (ex: v0.0.1), or 'latest'")
+
+# map the target platform to the release asset name
+set(LLAMA_CONNECT_PLATFORM "")
+
+if (APPLE)
+    set(_arch "${CMAKE_SYSTEM_PROCESSOR}")
+    if (CMAKE_OSX_ARCHITECTURES)
+        list(LENGTH CMAKE_OSX_ARCHITECTURES _narch)
+        if (_narch EQUAL 1)
+            set(_arch "${CMAKE_OSX_ARCHITECTURES}")
+        else()
+            set(_arch "universal") # no such asset, so the tool is skipped below
+        endif()
+    endif()
+
+    # note: there is no macos-x64 asset, it is not built upstream (no Intel macOS runners)
+    if (_arch MATCHES "^(arm64|aarch64)$")
+        set(LLAMA_CONNECT_PLATFORM "macos-arm64")
+    endif()
+elseif (WIN32)
+    # with Visual Studio generators, the target arch comes from -A
+    set(_arch "${CMAKE_SYSTEM_PROCESSOR}")
+    if (CMAKE_GENERATOR_PLATFORM)
+        set(_arch "${CMAKE_GENERATOR_PLATFORM}")
+    endif()
+
+    if (_arch MATCHES "^(x86_64|AMD64|amd64|x64)$")
+        set(LLAMA_CONNECT_PLATFORM "win-x64")
+    endif()
+elseif (CMAKE_SYSTEM_NAME STREQUAL "Linux")
+    if (CMAKE_SYSTEM_PROCESSOR MATCHES "^(x86_64|AMD64|amd64|x64)$")
+        set(LLAMA_CONNECT_PLATFORM "linux-x64")
+    elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "^(arm64|aarch64)$")
+        set(LLAMA_CONNECT_PLATFORM "linux-arm64")
+    endif()
+endif()
+
+if (LLAMA_CONNECT_PLATFORM STREQUAL "")
+    message(WARNING "llama-connect: no prebuilt binary for ${CMAKE_SYSTEM_NAME}/${CMAKE_SYSTEM_PROCESSOR}, skipping")
+    return()
+endif()
+
+if (LLAMA_CONNECT_PLATFORM MATCHES "^win-")
+    set(LLAMA_CONNECT_BINARY "llama-connect.exe")
+else()
+    set(LLAMA_CONNECT_BINARY "llama-connect")
+endif()
+
+# multi-config generators put the binaries in a per-config subdir
+get_property(LLAMA_CONNECT_MULTI_CONFIG GLOBAL PROPERTY GENERATOR_IS_MULTI_CONFIG)
+if (LLAMA_CONNECT_MULTI_CONFIG)
+    set(LLAMA_CONNECT_OUT_DIR "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/$<CONFIG>")
+else()
+    set(LLAMA_CONNECT_OUT_DIR "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}")
+endif()
+
+# the script is a no-op if the binary is already downloaded
+add_custom_target(${TARGET} ALL
+    BYPRODUCTS "${LLAMA_CONNECT_OUT_DIR}/${LLAMA_CONNECT_BINARY}"
+    COMMAND ${CMAKE_COMMAND}
+        "-DREPO=${LLAMA_CONNECT_REPO}"
+        "-DVERSION=${LLAMA_CONNECT_VERSION}"
+        "-DPLATFORM=${LLAMA_CONNECT_PLATFORM}"
+        "-DOUT_DIR=${LLAMA_CONNECT_OUT_DIR}"
+        "-DWORK_DIR=${CMAKE_CURRENT_BINARY_DIR}"
+        -P "${PROJECT_SOURCE_DIR}/scripts/connect-download.cmake"
+    COMMENT "Downloading llama-connect (${LLAMA_CONNECT_PLATFORM}, ${LLAMA_CONNECT_VERSION})"
+    VERBATIM
+)
+
+# building llama-server should also pull the binary (for convenience)
+if (TARGET llama-server-impl)
+    add_dependencies(llama-server-impl ${TARGET})
+endif()
+
+if (LLAMA_TOOLS_INSTALL)
+    install(PROGRAMS "${LLAMA_CONNECT_OUT_DIR}/${LLAMA_CONNECT_BINARY}" TYPE BIN)
+endif()

tools/mtmd/clip-model.h

@@ -65,6 +65,7 @@ struct clip_hparams {
     int32_t image_max_pixels = -1;
     resize_algo image_resize_algo = RESIZE_ALGO_BICUBIC;
     pad_style image_resize_pad = PAD_CEIL; // padding style when resizing
+    bool image_round_half_even = false;    // smart_resize rounds .5 to even, like Python round() in the HF processor
     std::array<uint8_t, 3> image_pad_color = {0, 0, 0};
 
     // (preprocessor) for llava-uhd style models

tools/mtmd/clip.cpp

@@ -1634,6 +1634,8 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_DOTS_OCR:
                     {
                         hparams.rope_theta = 10000.0f;
+                        hparams.image_resize_pad  = PAD_NONE; // the HF processor resizes to the smart_resize size, no padding
+                        hparams.image_round_half_even = true; // smart_resize uses Python round()
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge);
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
                         get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
@@ -1750,6 +1752,8 @@ struct clip_model_loader {
                     {
                         hparams.n_merge = 2; // default value for Qwen 2 and 2.5
                         hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_pad  = PAD_NONE; // the HF processor resizes to the smart_resize size, no padding
+                        hparams.image_round_half_even = true; // smart_resize uses Python round()
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         get_u32(KEY_WIN_ATTN_PATTERN, hparams.n_wa_pattern, model.proj_type == PROJECTOR_TYPE_QWEN25VL); // only 2.5 requires it
                         // ref: https://huggingface.co/Qwen/Qwen2.5-VL-7B-Instruct/blob/main/preprocessor_config.json
@@ -1825,6 +1829,8 @@ struct clip_model_loader {
                     {
                         hparams.n_merge = 2; // spatial_merge_size
                         hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_pad  = PAD_NONE; // the HF processor resizes to the smart_resize size, no padding
+                        hparams.image_round_half_even = true; // smart_resize uses Python round()
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         get_u32(string_format(KEY_N_HEAD_KV, "vision"), hparams.n_head_kv);
                         // 1D banded sliding-window radius (visual_token_window_size); required
@@ -2092,6 +2098,8 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_EXAONE4_5:
                     {
                         hparams.n_merge = 2;
+                        hparams.image_resize_pad  = PAD_NONE; // the HF processor resizes to the smart_resize size, no padding
+                        hparams.image_round_half_even = true; // smart_resize uses Python round()
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         get_u32(KEY_WIN_ATTN_PATTERN, hparams.n_wa_pattern, false);
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);

tools/mtmd/mtmd-image.cpp

@@ -115,6 +115,7 @@ struct img_tool {
         int max_pixels = 0;   // 0 = disabled
         // applied before min/max_pixels, so min_pixels can push an edge back above longest_edge
         int longest_edge = 0; // 0 = disabled
+        bool round_half_even = false; // Python round() instead of std::round
     };
 
     // calculate the size of the **resized** image, while preserving the aspect ratio and
@@ -127,7 +128,10 @@ struct img_tool {
             return {0, 0};
         }
 
-        auto round_by_factor = [f = opts.align_size](float x) { return static_cast<int>(std::round(x / static_cast<float>(f))) * f; };
+        auto round_by_factor = [f = opts.align_size, even = opts.round_half_even](float x) {
+            const float q = x / static_cast<float>(f);
+            return static_cast<int>(even ? std::nearbyint(q) : std::round(q)) * f;
+        };
         auto ceil_by_factor  = [f = opts.align_size](float x) { return static_cast<int>(std::ceil(x / static_cast<float>(f))) * f; };
         auto floor_by_factor = [f = opts.align_size](float x) { return static_cast<int>(std::floor(x / static_cast<float>(f))) * f; };
 
@@ -767,20 +771,24 @@ mtmd_image_preproc_out mtmd_image_preprocessor_fixed_size::preprocess(const clip
 // mtmd_image_preprocessor_dyn_size
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_dyn_size::preprocess(const clip_image_u8 & img) const {
+clip_image_size mtmd_image_preprocessor_dyn_size::get_target_size(const clip_hparams & hparams, const clip_image_size & original_size) {
     GGML_ASSERT(hparams.image_min_pixels > 0 && hparams.image_max_pixels > 0);
-    clip_image_u8 resized_image;
-    const clip_image_size original_size = img.get_size();
     // the original pixtral model doesn't have n_merge
     const int cur_merge = hparams.n_merge;
-    const clip_image_size target_size = img_tool::calc_size_preserved_ratio(
+    return img_tool::calc_size_preserved_ratio(
         original_size,
         {
-            /* align_size   */ hparams.patch_size * cur_merge,
-            /* min_pixels   */ hparams.image_min_pixels,
-            /* max_pixels   */ hparams.image_max_pixels,
-            /* longest_edge */ 0,
+            /* align_size      */ hparams.patch_size * cur_merge,
+            /* min_pixels      */ hparams.image_min_pixels,
+            /* max_pixels      */ hparams.image_max_pixels,
+            /* longest_edge    */ 0,
+            /* round_half_even */ hparams.image_round_half_even,
         });
+}
+
+mtmd_image_preproc_out mtmd_image_preprocessor_dyn_size::preprocess(const clip_image_u8 & img) const {
+    clip_image_u8 resized_image;
+    const clip_image_size target_size = get_target_size(hparams, img.get_size());
     img_tool::resize(img, resized_image, target_size,
                         hparams.image_resize_algo,
                         hparams.image_resize_pad,

tools/mtmd/mtmd-image.h

@@ -131,6 +131,7 @@ struct mtmd_image_preprocessor_fixed_size : mtmd_image_preprocessor {
 // this is used by models with native support for dynamic image size, for example: Qwen-VL, Pixtral, Kimi-VL, etc
 struct mtmd_image_preprocessor_dyn_size : mtmd_image_preprocessor {
     mtmd_image_preprocessor_dyn_size(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
+    static clip_image_size get_target_size(const clip_hparams & hparams, const clip_image_size & original_size);
     mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 

tools/server/CMakeLists.txt

@@ -41,6 +41,8 @@ set(TARGET llama-server-impl)
 
 add_library(${TARGET}
     server.cpp
+    server-connect.cpp
+    server-connect.h
     server-http.cpp
     server-http.h
     server-models.cpp

tools/server/README.md

@@ -208,6 +208,8 @@ For the full list of features, please refer to [server's changelog](https://gith
 | `--mcp-servers-json JSON` | experimental: inline JSON with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_MCP_SERVERS_JSON) |
 | `-ag, --agent, -no-ag, --no-agent` | whether to enable CORS proxy and all built-in tools - do not enable in untrusted environments (default: disabled)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_AGENT) |
 | `--ui, --webui, --no-ui, --no-webui` | whether to enable the Web UI (default: enabled)<br/>(env: LLAMA_ARG_UI) |
+| `--connect` | (experimental) open a peer-to-peer tunnel for server using llama-connect (default: disabled)<br/>(env: LLAMA_ARG_CONNECT) |
+| `--connect-code CODE` | manually specify a 40-character code for --connect (default: generate a new one for each run)<br/>(env: LLAMA_ARG_CONNECT_CODE) |
 | `--embedding, --embeddings` | restrict to only support embedding use case; use only with dedicated embedding models (default: disabled)<br/>(env: LLAMA_ARG_EMBEDDINGS) |
 | `--rerank, --reranking` | enable reranking endpoint on server (default: disabled)<br/>(env: LLAMA_ARG_RERANKING) |
 | `--api-key KEY` | API key to use for authentication, multiple keys can be provided as a comma-separated list (default: none)<br/>(env: LLAMA_API_KEY) |
@@ -235,7 +237,7 @@ For the full list of features, please refer to [server's changelog](https://gith
 | `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
 | `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
 | `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
-| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
+| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
 | `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
 | `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
 | `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |

tools/server/server-common.cpp

@@ -9,6 +9,7 @@
 
 #include "server-common.h"
 
+#include <filesystem>
 #include <random>
 #include <sstream>
 #include <fstream>
@@ -33,6 +34,60 @@
 #include <unistd.h>
 #endif
 
+#if defined(_WIN32)
+#define WIN32_LEAN_AND_MEAN
+#ifndef NOMINMAX
+#   define NOMINMAX
+#endif
+#include <windows.h>
+#elif defined(__APPLE__) && defined(__MACH__)
+#include <mach-o/dyld.h>
+#include <limits.h>
+#else
+#include <unistd.h>
+#endif
+
+std::filesystem::path get_server_exec_path() {
+#if defined(_WIN32)
+    wchar_t buf[32768] = { 0 };  // Large buffer to handle long paths
+    DWORD len = GetModuleFileNameW(nullptr, buf, _countof(buf));
+    if (len == 0 || len >= _countof(buf)) {
+        throw std::runtime_error("GetModuleFileNameW failed or path too long");
+    }
+    return std::filesystem::path(buf);
+#elif defined(__APPLE__) && defined(__MACH__)
+    char small_path[PATH_MAX];
+    uint32_t size = sizeof(small_path);
+
+    if (_NSGetExecutablePath(small_path, &size) == 0) {
+        // resolve any symlinks to get absolute path
+        try {
+            return std::filesystem::canonical(std::filesystem::path(small_path));
+        } catch (...) {
+            return std::filesystem::path(small_path);
+        }
+    } else {
+        // buffer was too small, allocate required size and call again
+        std::vector<char> buf(size);
+        if (_NSGetExecutablePath(buf.data(), &size) == 0) {
+            try {
+                return std::filesystem::canonical(std::filesystem::path(buf.data()));
+            } catch (...) {
+                return std::filesystem::path(buf.data());
+            }
+        }
+        throw std::runtime_error("_NSGetExecutablePath failed after buffer resize");
+    }
+#else
+    char path[FILENAME_MAX];
+    ssize_t count = readlink("/proc/self/exe", path, FILENAME_MAX);
+    if (count <= 0) {
+        throw std::runtime_error("failed to resolve /proc/self/exe");
+    }
+    return std::filesystem::path(std::string(path, count));
+#endif
+}
+
 json format_error_response(const std::string & message, const enum error_type type) {
     std::string type_str;
     int code = 500;

tools/server/server-common.h

@@ -15,6 +15,7 @@
 #include <condition_variable>
 #include <cinttypes>
 #include <cstdio>
+#include <filesystem>
 #include <functional>
 #include <mutex>
 #include <queue>
@@ -94,6 +95,9 @@ struct server_grammar_trigger {
 
 json format_error_response(const std::string & message, const enum error_type type);
 
+// path of the running llama-server binary, used to spawn siblings. throws on failure
+std::filesystem::path get_server_exec_path();
+
 //
 // random string / id
 //

tools/server/server-connect.cpp

@@ -0,0 +1,297 @@
+#include "server-connect.h"
+
+#include "server-common.h"
+#include "subproc.h"
+
+#include <cctype>
+#include <chrono>
+#include <exception>
+#include <filesystem>
+#include <random>
+#include <string_view>
+#include <system_error>
+#include <thread>
+#include <vector>
+
+// share code = room code + pass code, must match llama-connect and the Web UI
+// ref: https://github.com/ggml-org/llama-connect/blob/master/src/protocol.rs
+static constexpr size_t CONNECT_ROOM_CODE_LEN = 8;
+static constexpr size_t CONNECT_PASS_CODE_LEN = 32;
+
+static const std::string CONNECT_CODE_CHARS = "ABCDEFGHJKMNPQRSTUVWXYZabcdefghjkmnpqrstuvwxyz23456789";
+
+#if defined(_WIN32)
+static const std::string CONNECT_EXE_NAME = "llama-connect.exe";
+static constexpr char    PATH_SEPARATOR   = ';';
+#else
+static const std::string CONNECT_EXE_NAME = "llama-connect";
+static constexpr char    PATH_SEPARATOR   = ':';
+#endif
+
+// how long to wait for the child to notice the closed stdin before killing it
+static constexpr int CONNECT_STOP_TIMEOUT_MS = 3000;
+
+// the pass code guards the tunnel, so do not use random_string(): its mt19937 is predictable
+static std::string gen_share_code() {
+    std::random_device rd;
+    std::uniform_int_distribution<size_t> dist(0, CONNECT_CODE_CHARS.size() - 1);
+
+    std::string code(CONNECT_ROOM_CODE_LEN + CONNECT_PASS_CODE_LEN, ' ');
+    for (char & c : code) {
+        c = CONNECT_CODE_CHARS[dist(rd)];
+    }
+
+    return code;
+}
+
+static std::string format_share_code(const std::string & code) {
+    std::string out;
+    for (size_t i = 0; i < code.size(); i += 8) {
+        if (i > 0) {
+            out += ' ';
+        }
+        out += code.substr(i, 8);
+    }
+    return out;
+}
+
+// whitespace is tolerated, the Web UI shows the code in blocks and users paste it back
+// returns an empty string if it is not a valid share code
+static std::string normalize_share_code(const std::string & key) {
+    std::string code;
+    for (char c : key) {
+        if (!std::isspace((unsigned char) c)) {
+            code += c;
+        }
+    }
+
+    if (code.size() != CONNECT_ROOM_CODE_LEN + CONNECT_PASS_CODE_LEN) {
+        return "";
+    }
+
+    if (code.find_first_not_of(CONNECT_CODE_CHARS) != std::string::npos) {
+        return "";
+    }
+
+    return code;
+}
+
+// llama-connect logs as "[LEVEL] text", forward at the same level so our verbosity filter applies
+static void forward_child_log(const std::string & line) {
+    static const std::pair<std::string_view, ggml_log_level> tags[] = {
+        { "[ERROR] ", GGML_LOG_LEVEL_ERROR },
+        { "[WARN] ",  GGML_LOG_LEVEL_WARN  },
+        { "[INFO] ",  GGML_LOG_LEVEL_INFO  },
+        { "[DEBUG] ", GGML_LOG_LEVEL_DEBUG },
+        { "[TRACE] ", GGML_LOG_LEVEL_DEBUG },
+    };
+
+    // untagged lines are the startup banner, show them as info
+    ggml_log_level level = GGML_LOG_LEVEL_INFO;
+    const char * text = line.c_str();
+
+    for (const auto & [tag, tag_level] : tags) {
+        if (string_starts_with(line, tag)) {
+            level = tag_level;
+            text += tag.size();
+            break;
+        }
+    }
+
+    switch (level) {
+        case GGML_LOG_LEVEL_ERROR: LOG_ERR("connect | %s", text); break;
+        case GGML_LOG_LEVEL_WARN:  LOG_WRN("connect | %s", text); break;
+        case GGML_LOG_LEVEL_DEBUG: LOG_DBG("connect | %s", text); break;
+        default:                   LOG_INF("connect | %s", text); break;
+    }
+}
+
+// the child dials over TCP: a wildcard bind is reached on loopback, a specific address on itself
+// empty when every listener is a unix socket
+static std::string upstream_host(const std::vector<std::string> & hosts) {
+    for (const auto & host : hosts) {
+        if (string_ends_with(host, ".sock")) {
+            continue;
+        }
+        if (host == "0.0.0.0") {
+            return "127.0.0.1";
+        }
+        if (host == "::") {
+            return "[::1]";
+        }
+        // llama-connect builds http://host:port, so an IPv6 literal needs brackets
+        if (host.find(':') != std::string::npos && host[0] != '[') {
+            return "[" + host + "]";
+        }
+        return host;
+    }
+    return "";
+}
+
+static bool path_is_file(const std::filesystem::path & p) {
+    std::error_code ec;
+    return std::filesystem::is_regular_file(p, ec);
+}
+
+std::string server_connect::find_binary() {
+    // prefer the copy shipped next to llama-server over an unrelated one in PATH
+    try {
+        auto sibling = get_server_exec_path().parent_path() / CONNECT_EXE_NAME;
+        if (path_is_file(sibling)) {
+            return sibling.string();
+        }
+    } catch (const std::exception & e) {
+        SRV_WRN("could not resolve the llama-server path (%s), looking for llama-connect in PATH only\n", e.what());
+    }
+
+    const std::string path_env = common_get_env("PATH");
+    size_t start = 0;
+    while (start <= path_env.size()) {
+        size_t end = path_env.find(PATH_SEPARATOR, start);
+        if (end == std::string::npos) {
+            end = path_env.size();
+        }
+        const std::string dir = path_env.substr(start, end - start);
+        if (!dir.empty()) {
+            auto candidate = std::filesystem::path(dir) / CONNECT_EXE_NAME;
+            if (path_is_file(candidate)) {
+                return candidate.string();
+            }
+        }
+        start = end + 1;
+    }
+
+    return "";
+}
+
+std::string server_connect::unavailable_reason(const common_params & params) {
+    if (!common_subproc::is_supported()) {
+        return "this build has subprocess support disabled, rebuild with -DLLAMA_SUBPROCESS=ON";
+    }
+
+    if (!params.server_connect_code.empty() && normalize_share_code(params.server_connect_code).empty()) {
+        return "--connect-code must be " + std::to_string(CONNECT_ROOM_CODE_LEN + CONNECT_PASS_CODE_LEN)
+             + " characters from '" + CONNECT_CODE_CHARS + "'";
+    }
+
+    if (upstream_host(params.hostnames).empty()) {
+        return "llama-connect reaches the server over TCP, add a TCP address to --host";
+    }
+
+    if (find_binary().empty()) {
+        return "could not find '" + CONNECT_EXE_NAME + "' next to llama-server or in PATH.\n"
+               "    it is a separate binary: download it from https://github.com/ggml-org/llama-connect/releases\n"
+               "    or build llama.cpp with -DLLAMA_CONNECT=ON to have it fetched automatically";
+    }
+
+    return "";
+}
+
+std::string server_connect::resolve_code(const std::string & code) {
+    return code.empty() ? gen_share_code() : normalize_share_code(code);
+}
+
+bool server_connect::start(const common_params & params, int port) {
+    const std::string bin = find_binary();
+    if (bin.empty()) {
+        SRV_ERR("%s", "llama-connect binary not found\n");
+        return false;
+    }
+
+    // main() resolves this before anything can read it from /props
+    const std::string & code = params.server_connect_code;
+    GGML_ASSERT(!code.empty());
+
+    const std::vector<std::string> args = {
+        bin,
+        "--host", upstream_host(params.hostnames),
+        "--port", std::to_string(port),
+        "--code", code,
+        // the kernel closes our end of this pipe even if we are killed without cleanup
+        // this is what stops the child from outliving us
+        "--exit-on-stdin-eof",
+    };
+
+    proc = std::make_unique<common_subproc>();
+
+    const int options = subprocess_option_no_window
+                      | subprocess_option_combined_stdout_stderr
+                      | subprocess_option_inherit_environment;
+
+    if (!proc->create(args, options)) {
+        SRV_ERR("failed to spawn '%s'\n", bin.c_str());
+        proc.reset();
+        return false;
+    }
+
+    log_thread = std::thread([this]() {
+        FILE * out = proc->stdout_file();
+        if (out == nullptr) {
+            SRV_ERR("%s", "failed to get stdout of the llama-connect process\n");
+            return;
+        }
+        std::vector<char> buf(4096);
+        while (fgets(buf.data(), (int) buf.size(), out) != nullptr) {
+            forward_child_log(buf.data());
+        }
+        // EOF means the child is gone
+        if (!stopping.load(std::memory_order_acquire)) {
+            SRV_ERR("%s", "llama-connect exited on its own, remote access is no longer available\n");
+        }
+    });
+
+    SRV_INF("%s", "-----------------\n");
+    SRV_INF("%s", "remote access is enabled via llama-connect\n");
+    SRV_INF("share code (enter it in the Web UI under Settings -> Remote Access): %s\n",
+            format_share_code(code).c_str());
+    SRV_WRN("%s", "anyone with this code can use this server, do not share it publicly\n");
+    SRV_INF("%s", "-----------------\n");
+
+    return true;
+}
+
+void server_connect::stop() {
+    if (!proc) {
+        return;
+    }
+
+    SRV_INF("%s", "stopping llama-connect...\n");
+
+    stopping.store(true, std::memory_order_release);
+
+    proc->close_stdin();
+
+    for (int elapsed = 0; elapsed < CONNECT_STOP_TIMEOUT_MS && proc->alive(); elapsed += 100) {
+        std::this_thread::sleep_for(std::chrono::milliseconds(100));
+    }
+
+    // no-op if the child already exited. the log thread leaves fgets() on EOF, which the kernel
+    // raises once every writer of the stdout pipe is gone, so llama-connect must not spawn
+    // children of its own: one of them holding that pipe blocks the join below forever
+    proc->terminate();
+
+    if (log_thread.joinable()) {
+        try {
+            log_thread.join();
+        } catch (const std::system_error & e) {
+            // ~thread() on a still-joinable thread calls std::terminate, detach instead
+            SRV_ERR("failed to join the llama-connect log thread: %s\n", e.what());
+            log_thread.detach();
+        }
+    }
+
+    proc->join(); // reap the zombie
+    proc.reset();
+}
+
+server_connect::server_connect() = default;
+
+server_connect::~server_connect() {
+    try {
+        stop();
+    } catch (const std::exception & e) {
+        SRV_ERR("failed to stop llama-connect: %s\n", e.what());
+    } catch (...) {
+        SRV_ERR("%s", "failed to stop llama-connect\n");
+    }
+}

tools/server/server-connect.h

@@ -0,0 +1,39 @@
+#pragma once
+
+#include "common.h"
+
+#include <atomic>
+#include <memory>
+#include <string>
+#include <thread>
+
+struct common_subproc;
+
+// spawns llama-connect, which exposes this server to a remote browser over WebRTC
+// core binary is in a separate project: https://github.com/ggml-org/llama-connect
+struct server_connect {
+    server_connect();
+    ~server_connect();
+
+    server_connect(const server_connect &) = delete;
+    server_connect & operator=(const server_connect &) = delete;
+
+    // path of the llama-connect binary, empty if not found
+    static std::string find_binary();
+
+    // why --connect cannot work here, empty if it can
+    static std::string unavailable_reason(const common_params & params);
+
+    static std::string resolve_code(const std::string & code);
+
+    // port is the bound one, it differs from params.port with --port 0 or the GCP override
+    bool start(const common_params & params, int port);
+
+    // idempotent, also called by the destructor
+    void stop();
+
+private:
+    std::unique_ptr<common_subproc> proc;
+    std::thread log_thread;
+    std::atomic<bool> stopping{false}; // tells the log thread the exit is expected
+};

tools/server/server-context.cpp

@@ -3781,6 +3781,9 @@ private:
 
                             llama_pos pos_next = slot.prompt.tokens.pos_next(n_past);
 
+                            // number of cached tokens shared with the new prompt, before any checkpoint is restored
+                            const int n_prefix = n_past;
+
                             // ref: https://github.com/ggml-org/llama.cpp/pull/24110
                             const bool has_new_tokens = (n_past < slot.task->n_tokens());
 
@@ -3891,10 +3894,10 @@ private:
                             }
 
                             {
-                                // erase any checkpoints with pos_max > pos_next
+                                // erase any checkpoints with pos_max > pos_next, or that cover tokens not shared with the new prompt
                                 for (auto it = slot.prompt.checkpoints.begin(); it != slot.prompt.checkpoints.end();) {
                                     const auto & cur = *it;
-                                    if (cur.pos_max > pos_next) {
+                                    if (cur.pos_max > pos_next || cur.n_tokens > n_prefix) {
                                         SLT_TRC(slot, "erased invalidated context checkpoint (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", n_swa = %d, pos_next = %d, size = %.3f MiB)\n", cur.pos_min, cur.pos_max, cur.n_tokens, n_swa, pos_next, (float) cur.size() / 1024 / 1024);
                                         it = slot.prompt.checkpoints.erase(it);
                                     } else {
@@ -5261,6 +5264,10 @@ static json get_res_props(const server_context_meta & meta, const common_params
         { "is_sleeping",                 is_sleeping },
         { "cors_proxy_enabled",          params.ui_mcp_proxy },
     };
+
+    if (!params.server_connect_code.empty()) {
+        props["connect_code"] = params.server_connect_code;
+    }
     if (params.use_jinja) {
         if (!tmpl_tools.empty()) {
             props["chat_template_tool_use"] = tmpl_tools;

tools/server/server-models.cpp

@@ -35,12 +35,6 @@
 extern char **environ;
 #endif
 
-#if defined(__APPLE__) && defined(__MACH__)
-// macOS: use _NSGetExecutablePath to get the executable path
-#include <mach-o/dyld.h>
-#include <limits.h>
-#endif
-
 #define DEFAULT_STOP_TIMEOUT 10 // seconds
 
 #define CMD_ROUTER_TO_CHILD_EXIT  "cmd_router_to_child:exit"
@@ -436,47 +430,6 @@ struct server_lru_sched {
 // delete). distinct from params.timeout_read/write which only applies to the generation proxy
 static constexpr int STREAM_LOOKUP_TIMEOUT_MS = 250;
 
-static std::filesystem::path get_server_exec_path() {
-#if defined(_WIN32)
-    wchar_t buf[32768] = { 0 };  // Large buffer to handle long paths
-    DWORD len = GetModuleFileNameW(nullptr, buf, _countof(buf));
-    if (len == 0 || len >= _countof(buf)) {
-        throw std::runtime_error("GetModuleFileNameW failed or path too long");
-    }
-    return std::filesystem::path(buf);
-#elif defined(__APPLE__) && defined(__MACH__)
-    char small_path[PATH_MAX];
-    uint32_t size = sizeof(small_path);
-
-    if (_NSGetExecutablePath(small_path, &size) == 0) {
-        // resolve any symlinks to get absolute path
-        try {
-            return std::filesystem::canonical(std::filesystem::path(small_path));
-        } catch (...) {
-            return std::filesystem::path(small_path);
-        }
-    } else {
-        // buffer was too small, allocate required size and call again
-        std::vector<char> buf(size);
-        if (_NSGetExecutablePath(buf.data(), &size) == 0) {
-            try {
-                return std::filesystem::canonical(std::filesystem::path(buf.data()));
-            } catch (...) {
-                return std::filesystem::path(buf.data());
-            }
-        }
-        throw std::runtime_error("_NSGetExecutablePath failed after buffer resize");
-    }
-#else
-    char path[FILENAME_MAX];
-    ssize_t count = readlink("/proc/self/exe", path, FILENAME_MAX);
-    if (count <= 0) {
-        throw std::runtime_error("failed to resolve /proc/self/exe");
-    }
-    return std::filesystem::path(std::string(path, count));
-#endif
-}
-
 static void unset_reserved_args(common_preset & preset, bool unset_model_args) {
     preset.unset_option("LLAMA_ARG_SSL_KEY_FILE");
     preset.unset_option("LLAMA_ARG_SSL_CERT_FILE");
@@ -486,6 +439,8 @@ static void unset_reserved_args(common_preset & preset, bool unset_model_args) {
     preset.unset_option("LLAMA_ARG_MODELS_MAX");
     preset.unset_option("LLAMA_ARG_MODELS_PRESET");
     preset.unset_option("LLAMA_ARG_MODELS_AUTOLOAD");
+    preset.unset_option("LLAMA_ARG_CONNECT");
+    preset.unset_option("LLAMA_ARG_CONNECT_CODE");
     if (unset_model_args) {
         preset.unset_option("LLAMA_ARG_MODEL");
         preset.unset_option("LLAMA_ARG_MMPROJ");
@@ -2282,7 +2237,7 @@ void server_models_routes::init_routes() {
         if (name.empty()) {
             // main instance
             auto res = std::make_unique<server_http_res>();
-            res_ok(res, {
+            json props = {
                 // TODO: add support for this on web UI
                 {"role",                 "router"},
                 {"max_instances",        params.models_max},
@@ -2298,7 +2253,11 @@ void server_models_routes::init_routes() {
                 {"ui_settings",          ui_settings},
                 {"build_info",           std::string(llama_build_info())},
                 {"cors_proxy_enabled",   params.ui_mcp_proxy},
-            });
+            };
+            if (!params.server_connect_code.empty()) {
+                props["connect_code"] = params.server_connect_code;
+            }
+            res_ok(res, props);
             return res;
         }
         return proxy_get(req);

tools/server/server.cpp

@@ -1,3 +1,4 @@
+#include "server-connect.h"
 #include "server-context.h"
 #include "server-http.h"
 #include "server-models.h"
@@ -191,6 +192,20 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         params.model_alias.insert(model_name);
     }
 
+    // check early, so a missing llama-connect fails-fast
+    if (params.server_connect) {
+        const std::string reason = server_connect::unavailable_reason(params);
+        if (!reason.empty()) {
+            SRV_ERR("--connect is not available: %s\n", reason.c_str());
+            return 1;
+        }
+        params.server_connect_code = server_connect::resolve_code(params.server_connect_code);
+    } else if (!params.server_connect_code.empty()) {
+        // the code is only meaningful while the tunnel runs, and /props keys off it being set
+        SRV_WRN("%s", "--connect-code is ignored without --connect\n");
+        params.server_connect_code.clear();
+    }
+
     // note: this is guaranteed to out-live ctx_http and tools
     server_mcp mcp_mgr;
 
@@ -347,6 +362,10 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         warn_names.push_back("router mode");
     }
 
+    if (params.server_connect) {
+        warn_names.push_back("peer-to-peer tunnel (llama-connect, experimental)");
+    }
+
     if (params.ui_mcp_proxy) {
         ctx_http.get ("/cors-proxy",      ex_wrapper(proxy_handler_get));
         ctx_http.post("/cors-proxy",      ex_wrapper(proxy_handler_post));
@@ -525,6 +544,16 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         SRV_INF("listening on %s\n", address.c_str());
     }
 
+    // spawn only once listening, so the port is bound and the child startup health check reaches us. the destructor stops it
+    server_connect connect_proc;
+    if (params.server_connect && !connect_proc.start(params, ctx_http.port)) {
+        SRV_ERR("%s", "exiting due to llama-connect error\n");
+        ctx_http.stop();
+        ctx_http.join();
+        clean_up();
+        return 1;
+    }
+
     if (is_router_server) {
         if (!params.models_preset_hf.empty()) {
             SRV_WRN(      "NOTE: using preset.ini from HF repo '%s'\n", params.models_preset_hf.c_str());

tools/server/tests/unit/test_completion.py

@@ -689,3 +689,81 @@ def test_completion_prompt_cache():
         assert "prompt_n" in timings and timings["prompt_n"] + timings["cache_n"] == n_prompt
         assert "predicted_n" in timings and timings["predicted_n"] == n_predict
         assert "tokens" in res.body and isinstance(res.body["tokens"], list)
+
+
+# a slot that switched to an unrelated prompt must not restore a context checkpoint of the previous one
+def test_completion_checkpoint_after_unrelated_prompt():
+    global server
+    server = ServerPreset.tinygemma3()
+    server.n_slots = 1
+    server.cache_ram = 0
+    server.temperature = 0.0
+    server.start()
+
+    def tokenize(content):
+        res = server.make_request("POST", "/tokenize", data={"content": content})
+        assert res.status_code == 200
+        return res.body["tokens"]
+
+    def top_logprobs(prompt, cache_prompt):
+        res = server.make_request("POST", "/completion", data={
+            "prompt": prompt,
+            "n_predict": 1,
+            "n_probs": 5,
+            "cache_prompt": cache_prompt,
+        })
+        assert res.status_code == 200
+        return [(p["id"], p["logprob"]) for p in res.body["completion_probabilities"][0]["top_logprobs"]]
+
+    tokens = tokenize("Purple elephants quietly compose sonnets about tax law")
+    expected = top_logprobs(tokens[:2], False)
+
+    # 5 tokens: the checkpoint 4 tokens before the end of the prompt holds only its first token
+    top_logprobs(tokenize("The quick brown fox jumps over the lazy dog")[:5], True)
+    top_logprobs(tokens[:1], True)
+    actual = top_logprobs(tokens[:2], True)
+
+    assert [tid for tid, _ in actual] == [tid for tid, _ in expected]
+    for (_, lp), (_, lp_expected) in zip(actual, expected):
+        assert lp == pytest.approx(lp_expected, abs=1e-3)
+
+
+# a restored context checkpoint must be kept for the next request that shares the same prefix
+def test_completion_checkpoint_kept_after_restore():
+    global server
+    server = ServerPreset.tinygemma3()
+    server.n_slots = 1
+    server.n_batch = 16
+    server.n_ubatch = 16
+    server.cache_ram = 0
+    server.temperature = 0.0
+    server.start()
+
+    def tokenize(content):
+        res = server.make_request("POST", "/tokenize", data={"content": content})
+        assert res.status_code == 200
+        return res.body["tokens"]
+
+    def cache_n(prompt):
+        res = server.make_request("POST", "/completion", data={
+            "prompt": prompt,
+            "n_predict": 1,
+            "cache_prompt": True,
+        })
+        assert res.status_code == 200
+        return res.body["timings"]["cache_n"]
+
+    prefix = tokenize("Purple elephants quietly compose sonnets about tax law in the moonlight")
+    suffixes = [tokenize(s) for s in [
+        "The quick brown fox jumps over the lazy dog while the farmer counts his sheep and the sun slowly sets behind the distant hills",
+        "Seven tiny robots repaired the broken lighthouse lens before the storm arrived, then danced on the rocks until the morning tide came in",
+        "Green apples rolled down the hill toward the river, where a curious otter collected them one by one and stacked them on a flat stone",
+    ]]
+
+    # 4 tokens: the checkpoint 4 tokens before the end of the prompt holds exactly the prefix
+    cache_n(prefix + suffixes[0][:4])
+
+    # suffixes longer than n_ubatch + 4: no new checkpoint within the prefix, only the restored one can be reused
+    for suffix in suffixes[1:]:
+        assert len(suffix) > server.n_ubatch + 4
+        assert cache_n(prefix + suffix) == len(prefix) - 1

tools/ui/src/lib/components/app/content/SyntaxHighlightedCode.svelte

@@ -147,12 +147,19 @@
 	.code-preview-wrapper pre {
 		background: transparent;
 		padding: 0;
+		/* Stretch to the content's intrinsic width so a single long line
+		   propagates its overflow to this scrolling wrapper instead of being
+		   clipped at the wrapper's width with no way to scroll to it. */
+		width: max-content;
+		min-width: 100%;
 	}
 
 	.code-preview-wrapper code {
 		background: transparent;
 		display: block;
 		padding: 0.5rem;
+		width: max-content;
+		min-width: 100%;
 	}
 
 	:global(.dark) .code-preview-wrapper {

tools/ui/src/lib/components/app/forms/SearchInput.svelte

@@ -61,7 +61,9 @@
 		bind:ref
 		bind:value
 		{autofocus}
-		class="max-md:h-10 pl-9 {showClearButton ? 'pr-9' : ''} {size === 'sm' ? 'h-8' : ''}"
+		class="max-md:h-10 pl-9 {showClearButton ? 'pr-9' : ''} {size === 'sm'
+			? 'h-8'
+			: ''} [&::-webkit-search-cancel-button]:hidden"
 		{id}
 		oninput={handleInput}
 		onkeydown={onKeyDown}

ty.toml

@@ -8,7 +8,6 @@ deprecated = "warn"
 [src]
 exclude = [
     "./tools/mtmd/legacy-models/**",
-    "./skills/**",
 ]
 
 [[overrides]]