ci : disable ubuntu-rocm (#26969)

* disable ubuntu-rocm

* link PR

.github/workflows/build-cache.yml

@@ -119,27 +119,27 @@ jobs:
           version_major: ${{ env.OPENVINO_VERSION_MAJOR }}
           version_full: ${{ env.OPENVINO_VERSION_FULL }}
 
-  windows-2022-rocm-cache:
-    runs-on: windows-2022
+  # windows-2022-rocm-cache:
+  #   runs-on: windows-2022
 
-    env:
-      # Make sure this is in sync with release.yml and build-cuda-windows.yml
-      ROCM_VERSION: "7.14.0"
+  #   env:
+  #     # Make sure this is in sync with release.yml and build-cuda-windows.yml
+  #     ROCM_VERSION: "7.14.0"
 
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
+  #   steps:
+  #     - name: Clone
+  #       id: checkout
+  #       uses: actions/checkout@v6
 
-      - name: Setup Cache
-        uses: actions/cache@v5
-        id: cache-rocm
-        with:
-          path: C:\TheRock\build
-          key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
+  #     - name: Setup Cache
+  #       uses: actions/cache@v5
+  #       id: cache-rocm
+  #       with:
+  #         path: C:\TheRock\build
+  #         key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
 
-      - name: Setup ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
-        uses: ./.github/actions/windows-setup-rocm
-        with:
-          version: ${{ env.ROCM_VERSION }}
+  #     - name: Setup ROCm
+  #       if: steps.cache-rocm.outputs.cache-hit != 'true'
+  #       uses: ./.github/actions/windows-setup-rocm
+  #       with:
+  #         version: ${{ env.ROCM_VERSION }}

.github/workflows/build-cmake-pkg.yml

@@ -21,6 +21,7 @@ jobs:
                 -DLLAMA_BUILD_TOOLS=OFF \
                 -DLLAMA_BUILD_EXAMPLES=OFF \
                 -DLLAMA_BUILD_APP=OFF \
+                -DLLAMA_BUILD_IS_DEV=OFF \
                 -DCMAKE_BUILD_TYPE=Release
           cmake --build build --config Release
           cmake --install build --prefix "$PREFIX" --config Release
@@ -29,7 +30,12 @@ jobs:
           tclsh <<'EOF'
           set build(commit)  [string trim [exec git rev-parse --short HEAD]]
           set build(number)  [string trim [exec git rev-list  --count HEAD]]
-          set build(version) "0.0.$build(number)"
+
+          set cmakelists [read [open "CMakeLists.txt" r]]
+          regexp {set\(LLAMA_VERSION_MAJOR\s+(\d+)\)} $cmakelists -> major
+          regexp {set\(LLAMA_VERSION_MINOR\s+(\d+)\)} $cmakelists -> minor
+          regexp {set\(LLAMA_VERSION_PATCH\s+(\d+)\)} $cmakelists -> patch
+          set build(version) "$major.$minor.$patch"
 
           set llamaconfig [read [open "$env(LLAMA_CONFIG)" r]]
           set checks [list "set\\(LLAMA_VERSION     \\s+$build(version)\\)" \

.github/workflows/build-cpu.yml

@@ -95,7 +95,8 @@ jobs:
         run: |
           cmake -B build \
             -DLLAMA_FATAL_WARNINGS=ON \
-            -DGGML_RPC=ON
+            -DGGML_RPC=ON \
+            -DGGML_NATIVE=OFF
           time cmake --build build --config Release -j $(nproc)
 
       - name: Test

.github/workflows/build-cuda-windows.yml

@@ -97,15 +97,15 @@ jobs:
         id: checkout
         uses: actions/checkout@v6
 
-      - name: Cache ROCm Installation
-        uses: actions/cache@v5
-        id: cache-rocm
-        with:
-          path: C:\TheRock\build
-          key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
+      # - name: Cache ROCm Installation
+      #   uses: actions/cache@v5
+      #   id: cache-rocm
+      #   with:
+      #     path: C:\TheRock\build
+      #     key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
 
       - name: Setup ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
+        # if: steps.cache-rocm.outputs.cache-hit != 'true'
         uses: ./.github/actions/windows-setup-rocm
         with:
           version: ${{ env.ROCM_VERSION }}

.github/workflows/make-release.yml

@@ -0,0 +1,46 @@
+name: Make Release
+
+on:
+  workflow_dispatch:
+    inputs:
+      dry_run:
+        description: 'Dry run - validate without creating the tag'
+        required: true
+        type: boolean
+        default: true
+
+env:
+  GH_TOKEN: ${{ github.token }}
+
+permissions:
+  contents: write
+
+jobs:
+  make-release:
+    runs-on: ubuntu-latest
+
+    steps:
+      - name: Checkout
+        uses: actions/checkout@v6
+
+      - name: Run release checks
+        id: checks
+        run: bash scripts/make-release-checks.sh ${{ github.event.inputs.dry_run == 'true' && '--dry-run' || '' }}
+        env:
+          GITHUB_REPOSITORY: ${{ github.repository }}
+
+      - name: Create release tag
+        if: ${{ github.event.inputs.dry_run == 'false' }}
+        run: |
+          VERSION="${{ steps.checks.outputs.version }}"
+          git config user.name "github-actions[bot]"
+          git config user.email "github-actions[bot]@users.noreply.github.com"
+          git tag -a "${VERSION}" -m "Release ${VERSION}"
+          git push origin "${VERSION}"
+          echo "Created and pushed tag ${VERSION}"
+
+      - name: Dry run summary
+        if: ${{ github.event.inputs.dry_run == 'true' }}
+        run: |
+          echo "Dry run complete - all checks passed."
+          echo "Would have created tag: ${{ steps.checks.outputs.version }}"

.github/workflows/release.yml

@@ -774,15 +774,15 @@ jobs:
           key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
           evict-old-files: 1d
 
-      - name: Cache ROCm Installation
-        id: cache-rocm
-        uses: actions/cache@v5
-        with:
-          path: C:\TheRock\build
-          key: rocm-wheels-${{ matrix.ROCM_VERSION }}-multi-arch-${{ runner.os }}
+      # - name: Cache ROCm Installation
+      #   id: cache-rocm
+      #   uses: actions/cache@v5
+      #   with:
+      #     path: C:\TheRock\build
+      #     key: rocm-wheels-${{ matrix.ROCM_VERSION }}-multi-arch-${{ runner.os }}
 
       - name: Setup ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
+        # if: steps.cache-rocm.outputs.cache-hit != 'true'
         uses: ./.github/actions/windows-setup-rocm
         with:
           version: ${{ matrix.ROCM_VERSION }}
@@ -1285,123 +1285,123 @@ jobs:
           path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz
           name: llama-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz
 
-  ubuntu-22-rocm:
-    needs: [check-release, get-version]
-    if: ${{ needs.check-release.outputs.should_release == 'true' }}
+#   ubuntu-22-rocm:
+#     needs: [check-release, get-version]
+#     if: ${{ needs.check-release.outputs.should_release == 'true' }}
 
-    runs-on: ubuntu-22.04
+#     runs-on: ubuntu-22.04
 
-    permissions:
-      actions: write
+#     permissions:
+#       actions: write
 
-    strategy:
-      matrix:
-        include:
-          - ROCM_VERSION: "7.14.0"
-            gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
-            build: 'x64'
+#     strategy:
+#       matrix:
+#         include:
+#           - ROCM_VERSION: "7.14.0"
+#             gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
+#             build: 'x64'
 
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-        with:
-          fetch-depth: 0
+#     steps:
+#       - name: Clone
+#         id: checkout
+#         uses: actions/checkout@v6
+#         with:
+#           fetch-depth: 0
 
-      - name: Setup Node.js
-        uses: actions/setup-node@v6
-        with:
-          node-version: "24"
-          cache: "npm"
-          cache-dependency-path: "tools/ui/package-lock.json"
+#       - name: Setup Node.js
+#         uses: actions/setup-node@v6
+#         with:
+#           node-version: "24"
+#           cache: "npm"
+#           cache-dependency-path: "tools/ui/package-lock.json"
 
-      - name: Free up disk space
-        uses: ggml-org/free-disk-space@v1.3.1
-        with:
-          tool-cache: true
+#       - name: Free up disk space
+#         uses: ggml-org/free-disk-space@v1.3.1
+#         with:
+#           tool-cache: true
 
-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.21
-        with:
-          key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+#       # - name: ccache
+#       #   uses: ggml-org/ccache-action@v1.2.21
+#       #   with:
+#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
 
-      - name: Dependencies
-        id: depends
-        run: |
-          sudo apt install -y build-essential git cmake wget
+#       - name: Dependencies
+#         id: depends
+#         run: |
+#           sudo apt install -y build-essential git cmake wget
 
-      - name: Setup TheRock with Wheels
-        id: therock_env
-        run: |
-          # Create Python virtual environment
-          python3 -m venv .venv
-          source .venv/bin/activate
+#       - name: Setup TheRock with Wheels
+#         id: therock_env
+#         run: |
+#           # Create Python virtual environment
+#           python3 -m venv .venv
+#           source .venv/bin/activate
 
-          # Install ROCm wheels for build
-          # libraries = HIP runtime and CMake configs needed for linking
-          # devel = compilers, headers, static libs
-          python -m pip install --upgrade pip
-          python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"
+#           # Install ROCm wheels for build
+#           # libraries = HIP runtime and CMake configs needed for linking
+#           # devel = compilers, headers, static libs
+#           python -m pip install --upgrade pip
+#           python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"
 
-          # Get ROCm installation paths using the rocm-sdk CLI tool
-          ROCM_PATH=$(rocm-sdk path --root)
-          CMAKE_PATH=$(rocm-sdk path --cmake)
-          BIN_PATH=$(rocm-sdk path --bin)
-          echo "ROCM_PATH=$ROCM_PATH"
-          echo "CMAKE_PATH=$CMAKE_PATH"
-          echo "BIN_PATH=$BIN_PATH"
+#           # Get ROCm installation paths using the rocm-sdk CLI tool
+#           ROCM_PATH=$(rocm-sdk path --root)
+#           CMAKE_PATH=$(rocm-sdk path --cmake)
+#           BIN_PATH=$(rocm-sdk path --bin)
+#           echo "ROCM_PATH=$ROCM_PATH"
+#           echo "CMAKE_PATH=$CMAKE_PATH"
+#           echo "BIN_PATH=$BIN_PATH"
 
-          # Set environment variables
-          echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
-          echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
-          echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
-          echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
-          echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
+#           # Set environment variables
+#           echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
+#           echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
+#           echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
+#           echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
+#           echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
 
-          # Keep venv activated for subsequent steps
-          echo "$(pwd)/.venv/bin" >> $GITHUB_PATH
+#           # Keep venv activated for subsequent steps
+#           echo "$(pwd)/.venv/bin" >> $GITHUB_PATH
 
-      - name: Build with native CMake HIP support
-        id: cmake_build
-        run: |
-          cmake -B build -S . \
-            -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
-            -DCMAKE_BUILD_TYPE=Release \
-            -DGGML_BACKEND_DL=ON \
-            -DGGML_NATIVE=OFF \
-            -DCMAKE_INSTALL_RPATH='$ORIGIN' \
-            -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
-            -DGGML_CPU_ALL_VARIANTS=ON \
-            -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
-            -DGGML_HIP=ON \
-            -DHIP_PLATFORM=amd \
-            -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
-            ${{ env.CMAKE_ARGS }}
-          cmake --build build --config Release -j $(nproc)
+#       - name: Build with native CMake HIP support
+#         id: cmake_build
+#         run: |
+#           cmake -B build -S . \
+#             -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
+#             -DCMAKE_BUILD_TYPE=Release \
+#             -DGGML_BACKEND_DL=ON \
+#             -DGGML_NATIVE=OFF \
+#             -DCMAKE_INSTALL_RPATH='$ORIGIN' \
+#             -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+#             -DGGML_CPU_ALL_VARIANTS=ON \
+#             -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
+#             -DGGML_HIP=ON \
+#             -DHIP_PLATFORM=amd \
+#             -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
+#             ${{ env.CMAKE_ARGS }}
+#           cmake --build build --config Release -j $(nproc)
 
-      - name: ccache-clear
-        uses: ./.github/actions/ccache-clear
-        with:
-          key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+#       # - name: ccache-clear
+#       #   uses: ./.github/actions/ccache-clear
+#       #   with:
+#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
 
-      - name: Determine tag name
-        id: tag
-        uses: ./.github/actions/get-tag-name
+#       - name: Determine tag name
+#         id: tag
+#         uses: ./.github/actions/get-tag-name
 
-      - name: Get ROCm short version
-        run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
+#       - name: Get ROCm short version
+#         run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
 
-      - name: Pack artifacts
-        id: pack_artifacts
-        run: |
-          cp LICENSE ./build/bin/
-          tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
+#       - name: Pack artifacts
+#         id: pack_artifacts
+#         run: |
+#           cp LICENSE ./build/bin/
+#           tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
 
-      - name: Upload artifacts
-        uses: actions/upload-artifact@v6
-        with:
-          path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
-          name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+#       - name: Upload artifacts
+#         uses: actions/upload-artifact@v6
+#         with:
+#           path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+#           name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
 
   ios-xcode:
     needs: [check-release, get-version]
@@ -1578,7 +1578,7 @@ jobs:
       #- windows-sycl
       - windows-rocm
       - windows-openvino
-      - ubuntu-22-rocm
+      #- ubuntu-22-rocm
       - ubuntu-cpu
       - ubuntu-vulkan
       - ubuntu-24-openvino
@@ -1688,7 +1688,7 @@ jobs:
             - [Ubuntu s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-s390x.tar.gz)
             - [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-x64.tar.gz)
             - [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-arm64.tar.gz)
-            - [Ubuntu x64 (ROCm 7.14)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-7.14-x64.tar.gz)
+            - Ubuntu x64 (ROCm 7.14)[DISABLED](https://github.com/ggml-org/llama.cpp/pull/26969)
             - [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-openvino-${{ needs.ubuntu-24-openvino.outputs.openvino_version }}-x64.tar.gz)
             - [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp32-x64.tar.gz)
             - [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp16-x64.tar.gz)

.github/workflows/winget.yml

@@ -19,6 +19,8 @@ jobs:
         run: |
           cargo binstall komac@2.16.0 -y
 
+      # TODO: This should later be updated to publish releases instead of
+      # development release builds.
       - name: Find latest release
         id: find_latest_release
         uses: actions/github-script@v8

CMakeLists.txt

@@ -2,6 +2,26 @@ cmake_minimum_required(VERSION 3.14...3.28) # for add_link_options and implicit
 project("llama.cpp" C CXX)
 include(CheckIncludeFileCXX)
 
+### llama.cpp version
+set(LLAMA_VERSION_MAJOR 0)
+set(LLAMA_VERSION_MINOR 1)
+set(LLAMA_VERSION_PATCH 0)
+set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
+
+# whether this is a development/nightly build
+# set this to OFF when making a release from a release tag (vX.Y.Z)
+# ref: https://github.com/ggml-org/ggml/discussions/1579
+option(LLAMA_BUILD_IS_DEV "llama: dev build" ON)
+
+if (LLAMA_BUILD_IS_DEV)
+    set(LLAMA_VERSION "${LLAMA_VERSION_BASE}-dev")
+else()
+    # TODO: check that the current commit is tagged correctly according to the version specified above
+    set(LLAMA_VERSION "${LLAMA_VERSION_BASE}")
+endif()
+
+message(STATUS "llama.cpp version: ${LLAMA_VERSION}")
+
 #set(CMAKE_WARN_DEPRECATED YES)
 set(CMAKE_WARN_UNUSED_CLI YES)
 
@@ -24,9 +44,6 @@ if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
     set(LLAMA_STANDALONE ON)
 
     include(git-vars)
-
-    # configure project version
-    # TODO
 else()
     set(LLAMA_STANDALONE OFF)
 endif()
@@ -139,7 +156,6 @@ endif()
 if (NOT DEFINED LLAMA_BUILD_COMMIT)
     set(LLAMA_BUILD_COMMIT        ${BUILD_COMMIT})
 endif()
-set(LLAMA_INSTALL_VERSION 0.0.${LLAMA_BUILD_NUMBER})
 
 # override ggml options
 set(GGML_ALL_WARNINGS   ${LLAMA_ALL_WARNINGS})
@@ -275,12 +291,12 @@ configure_package_config_file(
               LLAMA_BIN_INSTALL_DIR )
 
 write_basic_package_version_file(
-        ${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
-    VERSION ${LLAMA_INSTALL_VERSION}
+        ${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
+        VERSION ${LLAMA_VERSION}
     COMPATIBILITY SameMajorVersion)
 
 install(FILES ${CMAKE_CURRENT_BINARY_DIR}/llama-config.cmake
-              ${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
+              ${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
         DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/llama)
 
 configure_file(cmake/llama.pc.in

README.md

@@ -106,6 +106,7 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
 - [XCFramework](docs/xcframework.md)
 - [Completions](docs/completions.md)
 - [Models](docs/models.md)
+- [Release process](docs/release.md)
 
 ## Contributing
 

app/llama.cpp

@@ -1,5 +1,7 @@
 #include "build-info.h"
 
+#include "llama.h"
+
 #include <cstdio>
 #include <cstdlib>
 #include <string>
@@ -77,12 +79,12 @@ static const command cmds[] = {
 
 #undef UPDATE_HIDDEN
 
-static int version(int argc, char ** argv) {
-    printf("%s\n", llama_build_info());
+static int version(int /*argc*/, char ** /*argv*/) {
+    llama_print_build_info(llama_version());
     return 0;
 }
 
-static int licenses(int argc, char ** argv) {
+static int licenses(int /*argc*/, char ** /*argv*/) {
     for (int i = 0; LICENSES[i]; ++i) {
         printf("%s\n", LICENSES[i]);
     }

cmake/llama-config.cmake.in

@@ -1,4 +1,4 @@
-set(LLAMA_VERSION      @LLAMA_INSTALL_VERSION@)
+set(LLAMA_VERSION      @LLAMA_VERSION@)
 set(LLAMA_BUILD_COMMIT @LLAMA_BUILD_COMMIT@)
 set(LLAMA_BUILD_NUMBER @LLAMA_BUILD_NUMBER@)
 set(LLAMA_SHARED_LIB   @BUILD_SHARED_LIBS@)

cmake/llama.pc.in

@@ -5,6 +5,6 @@ includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@
 
 Name: llama
 Description: Port of Facebook's LLaMA model in C/C++
-Version: @LLAMA_INSTALL_VERSION@
+Version: @LLAMA_VERSION@
 Libs: -L${libdir} -lggml -lggml-base -lllama
 Cflags: -I${includedir}

common/CMakeLists.txt

@@ -121,8 +121,8 @@ add_library(${TARGET}
     )
 
 set_target_properties(${TARGET} PROPERTIES
-    VERSION ${LLAMA_INSTALL_VERSION}
-    SOVERSION 0
+    VERSION ${LLAMA_VERSION_BASE}
+    SOVERSION ${LLAMA_VERSION_MAJOR}
     MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
 )
 

common/arg.cpp

@@ -1390,8 +1390,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         {"--version"},
         "show version and build info",
         [](common_params &) {
-            fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit());
-            fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
+            llama_print_build_info(llama_version());
             exit(0);
         }
     ));

common/build-info.cpp.in

@@ -29,7 +29,7 @@ const char * llama_build_info(void) {
     return s.c_str();
 }
 
-void llama_print_build_info(void) {
-    fprintf(stderr, "%s: build = %d (%s)\n",      __func__, llama_build_number(), llama_commit());
-    fprintf(stderr, "%s: built with %s for %s\n", __func__, llama_compiler(), llama_build_target());
+void llama_print_build_info(const char * llama_version) {
+    fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version, llama_build_number(), llama_commit());
+    fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
 }

common/build-info.h

@@ -8,4 +8,4 @@ const char * llama_compiler(void);
 const char * llama_build_target(void);
 const char * llama_build_info(void);
 
-void llama_print_build_info(void);
+void llama_print_build_info(const char *);

conversion/gemma.py

@@ -665,7 +665,18 @@ class Gemma4Model(Gemma3Model):
         swa_layers = [t == "sliding_attention" for t in self.hparams["layer_types"]]
         self.gguf_writer.add_sliding_window_pattern(swa_layers)
 
-        head_dim_full = self.hparams["global_head_dim"]
+        per_layer_config = self.hparams.get("per_layer_config")
+        layer_types = self.hparams.get("layer_types", [])
+        if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
+                        head_dim_full = layer_config["head_dim"]
+                        break
+
+        assert head_dim_full is not None
+
         head_dim_swa = self.hparams["head_dim"]
         # correct the head dim for global/swa layers
         self.gguf_writer.add_key_length(head_dim_full)
@@ -685,8 +696,14 @@ class Gemma4Model(Gemma3Model):
             n_ff_arr = [n_ff if il < first_kv_shared_layer_idx else n_ff * 2 for il in range(self.block_count)]
             self.gguf_writer.add_feed_forward_length(n_ff_arr)
 
-        # handle num_global_key_value_heads
-        num_key_value_heads_full = self.hparams.get("num_global_key_value_heads")
+        if (num_key_value_heads_full := self.hparams.get("num_global_key_value_heads")) is None and per_layer_config is not None:
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "num_key_value_heads" in layer_config:
+                        num_key_value_heads_full = layer_config["num_key_value_heads"]
+                        break
+
         num_key_value_heads_swa = self.hparams.get("num_key_value_heads")
         if num_key_value_heads_full is not None and num_key_value_heads_swa is not None:
             value_arr = [num_key_value_heads_swa if is_swa else num_key_value_heads_full for is_swa in swa_layers]
@@ -708,7 +725,19 @@ class Gemma4Model(Gemma3Model):
         # IMPORTANT: this ROPE_FREQS tensor is ONLY used by the full_attention layers
         rope_params_full = self.hparams["rope_parameters"]["full_attention"]
         assert rope_params_full["rope_type"] == "proportional"
-        head_dim_full = (self.hparams["global_head_dim"])
+
+        per_layer_config = self.hparams.get("per_layer_config")
+        if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
+            layer_types = self.hparams.get("layer_types", [])
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
+                        head_dim_full = layer_config["head_dim"]
+                        break
+
+        assert head_dim_full is not None
+
         partial_rotary_factor_full = rope_params_full["partial_rotary_factor"]
         n_rot_full = int(head_dim_full * partial_rotary_factor_full / 2)
         n_unrot_full = int(head_dim_full / 2) - n_rot_full

docs/release.md

@@ -0,0 +1,49 @@
+# Release process
+
+llama.cpp uses [semantic versioning](https://semver.org) (`MAJOR.MINOR.PATCH`).
+
+## Version bump guidelines
+
+| Change type | Version component |
+|---|---|
+| Breaking change to the public C API (`include/llama.h`)         | `MAJOR` |
+| Backward-compatible features, model support, or API addition    | `MINOR` |
+| Bug fix with no API change                                      | `PATCH` |
+
+The version is set in the three variables at the top of the root `CMakeLists.txt`:
+
+```cmake
+set(LLAMA_VERSION_MAJOR 0)
+set(LLAMA_VERSION_MINOR 1)
+set(LLAMA_VERSION_PATCH 0)
+```
+
+_A version bump should be included in the PR that introduces the change, or in a
+dedicated bump commit merged before the release is cut._
+
+_TODO: add PR labels (`semver: patch`, `semver: minor`, `semver: major`) to help
+identify which PRs require a version bump before cutting a release._
+
+## Making a release
+
+Releases are created by running the [make-release](.github/workflows/make-release.yml)
+which is a manual workflow.
+
+The workflow creates an annotated git tag (e.g. `v0.1.0`) and pushes it to the
+remote. No GitHub Release object is created, the tag is the release artifact.
+
+## Building a release
+
+By default, `LLAMA_BUILD_IS_DEV=ON` which appends a `-dev` suffix to `LLAMA_VERSION`,
+marking the build as a nightly/development build. Distributors building from a
+release tag must pass `-DLLAMA_BUILD_IS_DEV=OFF` to produce a clean version string
+(e.g. `0.1.0` instead of `0.1.0-dev`).
+
+## How releases reach users
+Currently releases are not published to github releases, only nightly/development
+builds are available there. The way users can access releases are using the following
+channels:
+
+- **llama-install.sh**  — downloads pre-built binaries built from the release tag.
+- **Package managers**  — consume the git tag directly.
+- **Build from source** — users clone the repo and check out the tag.

examples/test-cmake/.gitignore

@@ -0,0 +1,3 @@
+llama-build-install
+install
+build

examples/test-cmake/CMakeLists.txt

@@ -0,0 +1,13 @@
+cmake_minimum_required(VERSION 3.14)
+project(llama-simple)
+
+set(CMAKE_CXX_STANDARD 17)
+
+find_package(llama 0.1.0 REQUIRED)
+
+add_executable(test-cmake test-cmake.cpp)
+target_link_libraries(test-cmake PRIVATE llama)
+target_compile_definitions(test-cmake PRIVATE
+    LLAMA_BUILD_NUMBER=${LLAMA_BUILD_NUMBER}
+    LLAMA_BUILD_COMMIT="${LLAMA_BUILD_COMMIT}"
+)

examples/test-cmake/README.md

@@ -0,0 +1,36 @@
+## cmake-test
+
+This is just for manually testing/developing of a llama.cpp installation to
+enable troubleshooting issues and exploration. The idea is that this can be used
+after making changes to llama.cpp installation cmake configuration and then
+verify it locally.
+
+### Usage
+The following will configure, build, and install llama.cpp
+
+Configuring/build/install:
+```console
+./build-install.sh
+```
+The above command will create a directory named `install` in the current directory
+which will have the follwing files in its lib directory:
+```console
+(venv) $ ls install/lib/
+cmake                   libggml.so          libllama-common.so.0      libllama.so.0.1.0  llama.cpp
+libggml-base.so         libggml.so.0        libllama-common.so.0.1.0  libmtmd.so         pkgconfig
+libggml-base.so.0       libggml.so.0.19.0   libllama.so               libmtmd.so.0
+libggml-base.so.0.19.0  libllama-common.so  libllama.so.0             libmtmd.so.0.1.0
+```
+
+Build/run this project using the installation created above:
+```console
+(venv) $ ./build.sh
+-- Configuring done (0.0s)
+-- Generating done (0.0s)
+-- Build files have been written to: /home/danbev/work/ai/llama.cpp/examples/test-cmake/build
+[100%] Built target test-cmake
+[test-cmake] Using llama.cpp version 0.1.0-dev-b10335
+[test-cmake] Initializing backend...
+load_backend: loaded CPU backend from /home/danbev/work/ai/llama.cpp/examples/test-cmake/install/lib/llama.cpp/libggml-cpu-alderlake.so
+[test-cmake] Backend initialized.
+```

examples/test-cmake/build-install.sh

@@ -0,0 +1,19 @@
+#!/bin/bash
+
+set -e
+
+rm -rf llama-build-install install
+
+cmake --fresh -S ../../. -B llama-build-install -DCMAKE_BUILD_TYPE=Release \
+  -DBUILD_SHARED_LIBS=ON \
+  -DGGML_BACKEND_DL=ON \
+  -DGGML_CPU_ALL_VARIANTS=ON \
+  -DLLAMA_TESTS_INSTALL=OFF \
+  -DCMAKE_INSTALL_PREFIX="${PWD}/install" \
+  -DGGML_BACKEND_DIR="${PWD}/install/lib/llama.cpp" \
+  -DGGML_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
+  -DLLAMA_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
+  -DLLAMA_TOOLS_INSTALL=OFF
+
+cmake --build llama-build-install --parallel 12
+cmake --install llama-build-install

examples/test-cmake/build.sh

@@ -0,0 +1,7 @@
+#!/bin/bash
+
+set -e
+
+cmake -S . -B build -DCMAKE_PREFIX_PATH="${PWD}/install"
+cmake --build build
+LD_LIBRARY_PATH="${PWD}/install/lib/llama.cpp:${PWD}/install/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" ./build/test-cmake

examples/test-cmake/test-cmake.cpp

@@ -0,0 +1,12 @@
+#include "llama.h"
+#include <cstdio>
+
+int main(void) {
+    printf("[test-cmake] version: %s, build: %d (%s)\n",
+           llama_version(), LLAMA_BUILD_NUMBER, LLAMA_BUILD_COMMIT);
+    printf("[test-cmake] Initializing backend...\n");
+    llama_backend_init();
+    printf("[test-cmake] Backend initialized.\n");
+    llama_backend_free();
+    return 0;
+}

ggml/CMakeLists.txt

@@ -402,7 +402,7 @@ configure_package_config_file(
               GGML_BIN_INSTALL_DIR)
 
 write_basic_package_version_file(
-        ${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
+        ${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
     VERSION ${GGML_INSTALL_VERSION}
     COMPATIBILITY SameMajorVersion)
 
@@ -414,7 +414,7 @@ message(STATUS "ggml version: ${GGML_INSTALL_VERSION}")
 message(STATUS "ggml commit:  ${GGML_BUILD_COMMIT}")
 
 install(FILES ${CMAKE_CURRENT_BINARY_DIR}/ggml-config.cmake
-              ${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
+              ${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
         DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/ggml)
 
 if (MSVC)

ggml/cmake/ggml-config.cmake.in

@@ -113,6 +113,7 @@ set_and_check(GGML_LIB_DIR "@PACKAGE_GGML_LIB_INSTALL_DIR@")
 if(NOT TARGET ggml::ggml)
     find_package(Threads REQUIRED)
 
+    unset(GGML_LIBRARY CACHE)
     find_library(GGML_LIBRARY ggml
         REQUIRED
         HINTS ${GGML_LIB_DIR}
@@ -121,8 +122,10 @@ if(NOT TARGET ggml::ggml)
     add_library(ggml::ggml UNKNOWN IMPORTED)
     set_target_properties(ggml::ggml
         PROPERTIES
-            IMPORTED_LOCATION "${GGML_LIBRARY}")
+            IMPORTED_LOCATION "${GGML_LIBRARY}"
+            INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}")
 
+    unset(GGML_BASE_LIBRARY CACHE)
     find_library(GGML_BASE_LIBRARY ggml-base
         REQUIRED
         HINTS ${GGML_LIB_DIR}
@@ -132,6 +135,7 @@ if(NOT TARGET ggml::ggml)
     set_target_properties(ggml::ggml-base
         PROPERTIES
             IMPORTED_LOCATION "${GGML_BASE_LIBRARY}"
+            INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}"
             INTERFACE_LINK_LIBRARIES "${GGML_BASE_INTERFACE_LINK_LIBRARIES}")
 
     set(_ggml_all_targets "")
@@ -140,6 +144,7 @@ if(NOT TARGET ggml::ggml)
             string(REPLACE "-" "_" _ggml_backend_pfx "${_ggml_backend}")
             string(TOUPPER "${_ggml_backend_pfx}" _ggml_backend_pfx)
 
+            unset(${_ggml_backend_pfx}_LIBRARY CACHE)
             find_library(${_ggml_backend_pfx}_LIBRARY ${_ggml_backend}
                 REQUIRED
                 HINTS ${GGML_LIB_DIR}

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

@@ -1,90 +1,11 @@
 #include "ggml-backend-impl.h"
+#include "ggml-feats.h"
 
-#if defined(__aarch64__)
-
-#if defined(__linux__)
-#include <sys/auxv.h>
-#elif defined(__APPLE__)
-#include <sys/sysctl.h>
-#endif
-
-#if !defined(HWCAP_FPHP)
-#define HWCAP_FPHP (1 << 9)
-#endif
-
-#if !defined(HWCAP_ASIMDHP)
-#define HWCAP_ASIMDHP (1 << 10)
-#endif
-
-#if !defined(HWCAP_ASIMDDP)
-#define HWCAP_ASIMDDP (1 << 20)
-#endif
-
-#if !defined(HWCAP_SVE)
-#define HWCAP_SVE (1 << 22)
-#endif
-
-#if !defined(HWCAP2_SVE2)
-#define HWCAP2_SVE2 (1 << 1)
-#endif
-
-#if !defined(HWCAP2_I8MM)
-#define HWCAP2_I8MM (1 << 13)
-#endif
-
-#if !defined(HWCAP2_SME)
-#define HWCAP2_SME (1 << 23)
-#endif
-
-struct aarch64_features {
-    // has_neon not needed, aarch64 has NEON guaranteed
-    bool has_dotprod     = false;
-    bool has_fp16        = false;
-    bool has_sve         = false;
-    bool has_sve2        = false;
-    bool has_i8mm        = false;
-    bool has_sme         = false;
-    bool has_sme2        = false;
-
-    aarch64_features() {
-#if defined(__linux__)
-        uint32_t hwcap = getauxval(AT_HWCAP);
-        uint32_t hwcap2 = getauxval(AT_HWCAP2);
-
-        has_dotprod = !!(hwcap & HWCAP_ASIMDDP);
-        has_fp16    = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP);
-        has_sve     = !!(hwcap & HWCAP_SVE);
-        has_sve2    = !!(hwcap2 & HWCAP2_SVE2);
-        has_i8mm    = !!(hwcap2 & HWCAP2_I8MM);
-        has_sme     = !!(hwcap2 & HWCAP2_SME);
-#elif defined(__APPLE__)
-        int oldp = 0;
-        size_t size = sizeof(oldp);
-
-        if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, NULL, 0) == 0) {
-            has_dotprod = static_cast<bool>(oldp);
-        }
-
-        if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, NULL, 0) == 0) {
-            has_i8mm = static_cast<bool>(oldp);
-        }
-
-        if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, NULL, 0) == 0) {
-            has_sme = static_cast<bool>(oldp);
-        }
-
-        if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, NULL, 0) == 0) {
-            has_sme2 = static_cast<bool>(oldp);
-        }
-
-        // Apple apparently does not implement SVE yet
-#endif
-    }
-};
+#if defined(__aarch64__) || defined(_M_ARM64)
 
 static int ggml_backend_cpu_aarch64_score() {
     int score = 1;
-    aarch64_features af;
+    ggml_feats_arch64_runtime_t af = ggml_get_aarch64_runtime_features();
 
 #ifdef GGML_USE_DOTPROD
     if (!af.has_dotprod) { return 0; }
@@ -116,4 +37,4 @@ static int ggml_backend_cpu_aarch64_score() {
 
 GGML_BACKEND_DL_SCORE_IMPL(ggml_backend_cpu_aarch64_score)
 
-# endif // defined(__aarch64__)
+# endif // defined(__aarch64__) || defined(_M_ARM64)

ggml/src/ggml-cpu/kleidiai/kleidiai.cpp

@@ -2,10 +2,12 @@
 // SPDX-License-Identifier: MIT
 //
 #include <arm_neon.h>
-#include <assert.h>
-#include <stdio.h>
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
 #include <atomic>
 #include <cfloat>
+#include <cctype>
 #include <algorithm>
 #include <cmath>
 #include <stdexcept>
@@ -17,25 +19,21 @@
 #include <cstddef>
 #include <cstdint>
 #include <fstream>
-#include <set>
+#include <map>
 #include <iostream>
 #include <climits>
+#include <charconv>
+#include <system_error>
 #if defined(__linux__)
 #include <asm/hwcap.h>
+#include <dirent.h>
 #include <sys/auxv.h>
 #include <sys/types.h>
 #include <sys/stat.h>
 #include <unistd.h>
-#ifndef HWCAP2_SME2
-#define HWCAP2_SME2 (1UL << 37)
-#endif
 #elif defined(__APPLE__)
-#include <string_view>
 #include <sys/sysctl.h>
 #include <sys/types.h>
-#elif defined(_WIN32)
-#include <windows.h>
-#include <excpt.h>
 #endif
 
 #include "kleidiai.h"
@@ -43,6 +41,7 @@
 #include "ggml-cpu.h"
 #include "ggml-cpu-impl.h"
 #include "ggml-impl.h"
+#include "ggml-feats.h"
 #include "ggml-backend-impl.h"
 #include "ggml-threading.h"
 #include "traits.h"
@@ -64,8 +63,8 @@ struct ggml_kleidiai_context {
     ggml_kleidiai_kernels * kernels_q4;
     ggml_kleidiai_kernels * kernels_q8;
     ggml_kleidiai_kernels * kernels_f32;
-    int sme_thread_cap; // <= 0 means “SME disabled/unknown”;
-    int thread_hint;    // <= 0 means “no hint”
+    int sme_thread_cap; // <= 0 means "SME disabled/unknown"
+    int thread_hint;    // <= 0 means "no hint"
     int chunk_multiplier;
 } static ctx = { CPU_FEATURE_NONE, nullptr, nullptr, nullptr, 0, -1, 4 };
 
@@ -93,24 +92,117 @@ static const char* cpu_feature_to_string(cpu_feature f) {
     }
 }
 
+#if defined(__linux__) && defined(__aarch64__)
+static bool parse_cpu_dir_name(const char* name, size_t* cpu) {
+    if (strncmp(name, "cpu", 3) != 0 ||
+        name[3] < '0' || name[3] > '9') {
+        return false;
+    }
+
+    const char* first = name + 3;
+    const char* last = name + strlen(name);
+
+    size_t value = 0;
+    const auto [end, ec] = std::from_chars(first, last, value, 10);
+
+    if (ec != std::errc{} || end != last) {
+        return false;
+    }
+
+    *cpu = value;
+    return true;
+}
+
+static std::vector<size_t> detect_cpu_ids() {
+    std::vector<size_t> cpus;
+
+    DIR * dir = opendir("/sys/devices/system/cpu");
+    if (dir == nullptr) {
+        return cpus;
+    }
+
+    while (dirent * entry = readdir(dir)) {
+        size_t cpu = 0;
+        if (parse_cpu_dir_name(entry->d_name, &cpu)) {
+            cpus.push_back(cpu);
+        }
+    }
+    closedir(dir);
+
+    std::sort(cpus.begin(), cpus.end());
+    cpus.erase(std::unique(cpus.begin(), cpus.end()), cpus.end());
+    return cpus;
+}
+#endif
+
+#if defined(__APPLE__) && defined(__aarch64__)
+static bool apple_sme_counted_perf_level(std::string name) {
+    for (std::string::size_type i = 0; i < name.size(); ++i) {
+        name[i] = (char) std::tolower((unsigned char) name[i]);
+    }
+
+    // Conservative ceiling: only count perf-level names observed to provide full SME throughput.
+    // Future names should be calibrated here before they raise the automatic SME thread cap.
+    return name.find("super") != std::string::npos ||
+           name.find("performance") != std::string::npos;
+}
+#endif
+
+static void add_smcus_from_smidr(uint64_t smidr, size_t & num_private, std::map<uint32_t, size_t> & shared_counts) {
+    // Arm ARM: SMIDR_EL1. SH==0 is implementation-defined; keep the existing
+    // conservative policy and only treat zero affinity as private.
+    const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
+    const uint32_t nsmc = (uint32_t)((smidr >> 56) & 0xF);
+    const size_t shared_count = nsmc == 0xF ? 1 : (size_t)nsmc + 1;
+    const uint32_t affinity = (uint32_t)(smidr & 0xFFFu);
+    const uint32_t affinity2 = (uint32_t)((smidr >> 32) & 0xFFFFFu);
+    const uint32_t id = (affinity2 << 12) | affinity;
+
+    if (nsmc == 0xF) {
+        GGML_LOG_WARN("kleidiai: NSMC detected as 0xF indicating reseved value, setting min safe shared SMCU count to 1");
+    }
+
+    switch (sh) {
+        case 2: // private SMCU
+            ++num_private;
+            break;
+        case 3: // shared SMCU
+            if (shared_counts[id] < shared_count) {
+                shared_counts[id] = shared_count;
+            }
+            break;
+        case 0:
+            if (id == 0) {
+                ++num_private;
+            } else if (shared_counts[id] < shared_count) {
+                shared_counts[id] = shared_count;
+            }
+            break;
+        default:
+            break;
+    }
+}
+
 static size_t detect_num_smcus() {
-    if (!ggml_cpu_has_sme()) {
+    auto runtime_feat = ggml_get_aarch64_runtime_features();
+    if (!runtime_feat.has_sme) {
         return 0;
     }
 
 #if defined(__linux__) && defined(__aarch64__)
     // Linux/aarch64: Best-effort count of Streaming Mode Compute Units (SMCUs) via SMIDR_EL1 sysfs.
     size_t num_private = 0;
-    std::set<uint32_t> shared_ids;
+    std::map<uint32_t, size_t> shared_counts;
 
-    for (size_t cpu = 0;; ++cpu) {
+    const std::vector<size_t> cpus = detect_cpu_ids();
+    for (const size_t cpu : cpus) {
         const std::string path =
             "/sys/devices/system/cpu/cpu" + std::to_string(cpu) +
             "/regs/identification/smidr_el1";
 
         std::ifstream file(path);
         if (!file.is_open()) {
-            break;
+            continue;
         }
 
         uint64_t smidr = 0;
@@ -118,54 +210,69 @@ static size_t detect_num_smcus() {
             continue;
         }
 
-        // Arm ARM: SMIDR_EL1
-        const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
-        // Build an "affinity-like" identifier for shared SMCUs.
-        // Keep the original packing logic, but isolate it here.
-        const uint32_t id = (uint32_t)((smidr & 0xFFFu) | ((smidr >> 20) & 0xFFFFF000u));
-
-        switch (sh) {
-            case 0b10: // private SMCU
-                ++num_private;
-                break;
-            case 0b11: // shared SMCU
-                shared_ids.emplace(id);
-                break;
-            case 0b00:
-                // Ambiguous / implementation-defined. Be conservative:
-                // treat id==0 as private, otherwise as shared.
-                if (id == 0) ++num_private;
-                else shared_ids.emplace(id);
-                break;
-            default:
-                break;
-        }
+        add_smcus_from_smidr(smidr, num_private, shared_counts);
     }
 
-    return num_private + shared_ids.size();
+    size_t total = num_private;
+    for (const auto & entry : shared_counts) {
+        total += entry.second;
+    }
+    return total;
 
 #elif defined(__APPLE__) && defined(__aarch64__)
-    // table for known M4 variants. Users can override via GGML_KLEIDIAI_SME=<n>.
-    char chip_name[256] = {};
-    size_t size = sizeof(chip_name);
+    int perf_levels = 0;
+    size_t size = sizeof(perf_levels);
+    if (sysctlbyname("hw.nperflevels", &perf_levels, &size, nullptr, 0) != 0 ||
+        size != sizeof(perf_levels) || perf_levels <= 0) {
+        return 0;
+    }
 
-    if (sysctlbyname("machdep.cpu.brand_string", chip_name, &size, nullptr, 0) == 0) {
-        const std::string brand(chip_name);
+    size_t units = 0;
+    for (int i = 0; i < perf_levels; ++i) {
+        char key[64] = {};
+        int physical_cpus = 0;
+        int cpus_per_l2 = 0;
 
-        struct ModelSMCU { const char *match; size_t smcus; };
-        static const ModelSMCU table[] = {
-            { "M4 Ultra", 2 },
-            { "M4 Max",   2 },
-            { "M4 Pro",   2 },
-            { "M4",       1 },
-        };
+        snprintf(key, sizeof(key), "hw.perflevel%d.physicalcpu", i);
+        size = sizeof(physical_cpus);
+        if (sysctlbyname(key, &physical_cpus, &size, nullptr, 0) != 0 ||
+            size != sizeof(physical_cpus) || physical_cpus <= 0) {
+            continue;
+        }
 
-        for (const auto &e : table) {
-            if (brand.find(e.match) != std::string::npos) {
-                return e.smcus;
-            }
+        snprintf(key, sizeof(key), "hw.perflevel%d.cpusperl2", i);
+        size = sizeof(cpus_per_l2);
+        if (sysctlbyname(key, &cpus_per_l2, &size, nullptr, 0) != 0 ||
+            size != sizeof(cpus_per_l2) || cpus_per_l2 <= 0) {
+            continue;
+        }
+
+        snprintf(key, sizeof(key), "hw.perflevel%d.name", i);
+        size = 0;
+        if (sysctlbyname(key, nullptr, &size, nullptr, 0) != 0 || size == 0) {
+            continue;
+        }
+
+        std::string name(size, '\0');
+        if (sysctlbyname(key, &name[0], &size, nullptr, 0) != 0) {
+            continue;
+        }
+        name.resize(size);
+        while (!name.empty() && name.back() == '\0') {
+            name.pop_back();
+        }
+
+        if (apple_sme_counted_perf_level(name)) {
+            units += (size_t) ((physical_cpus + cpus_per_l2 - 1) / cpus_per_l2);
         }
     }
+
+    return units;
+
+#elif defined(_WIN32) && (defined(_M_ARM64) || defined(__aarch64__))
+    // No verified Windows arm64 SMCU detection path yet. Return unknown and use
+    // GGML_KLEIDIAI_SME=N as a diagnostics/debug override for SME thread cap
+    // calibration until a detection mechanism is verified on real hardware.
     return 0;
 
 #else
@@ -198,15 +305,18 @@ static void init_kleidiai_context(void) {
     if (!initialized) {
         initialized = true;
 
+        // Optional diagnostics/debug overrides; production defaults come from runtime detection.
         const char *env_sme         = getenv("GGML_KLEIDIAI_SME");
         const char *env_threads     = getenv("GGML_TOTAL_THREADS");
         const char *env_chunk_mult  = getenv("GGML_KLEIDIAI_CHUNK_MULTIPLIER");
 
+        auto runtime_feat = ggml_get_aarch64_runtime_features();
+
         size_t detected_smcus = 0;
 
-        ctx.features  = (ggml_cpu_has_dotprod()     ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
-                        (ggml_cpu_has_matmul_int8() ? CPU_FEATURE_I8MM    : CPU_FEATURE_NONE) |
-                        ((ggml_cpu_has_sve() && ggml_cpu_get_sve_cnt() == QK8_0) ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);
+        ctx.features  = (runtime_feat.has_dotprod  ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
+                        (runtime_feat.has_i8mm     ? CPU_FEATURE_I8MM    : CPU_FEATURE_NONE) |
+                        (runtime_feat.sve_cnt == QK8_0 ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);
 
         if (env_threads) {
             bool ok = false;
@@ -224,54 +334,54 @@ static void init_kleidiai_context(void) {
             }
         }
 
-        // SME policy:
-        // - env unset => auto-detect SMCUs; enable SME only if detected > 0.
-        // - env=0     => force off.
-        // - env>0     => force N cores, if the binary was built with SME.
         int sme_cores = 0;
         bool sme_env_ok = false;
         bool sme_env_set = (env_sme != nullptr);
 
+        const bool has_supported_sme_family = runtime_feat.has_sme;
+        bool sme_cap_detected = false;
+
+        if (has_supported_sme_family) {
+            detected_smcus = detect_num_smcus();
+            sme_cap_detected = detected_smcus > 0;
+            // Some platforms expose SME without exposing a calibrated SMCU count.
+            // Use one SME thread as the conservative default; add platform SMCU detection to raise it.
+            sme_cores = sme_cap_detected ? (int)detected_smcus : 1;
+
+            if (!sme_env_set && !sme_cap_detected) {
+                GGML_LOG_INFO("kleidiai: SME detected; SMCU count unavailable, using conservative SME thread cap=1\n");
+            }
+        }
+
+        // Runtime-detect SME support and available SMCUs first. The detected SMCU
+        // count is used as the SME thread cap, and GGML_KLEIDIAI_SME can debug-override that:
+        //   - unset: use runtime detection.
+        //   - 0:     disable SME-family kernels.
+        //   - N > 0: use N as the SME thread cap, if an SME-family kernel is selectable.
         if (sme_env_set) {
             bool ok = false;
             int v = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
             sme_env_ok = ok;
 
-            if (!ok) {
-                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; falling back to runtime SME-core detection\n");
-                detected_smcus = detect_num_smcus();
-                sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
-            } else if (v == 0) {
-                sme_cores = 0;
-            } else if (!ggml_cpu_has_sme()) {
-                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but the binary was not built with SME; disabling SME\n", v);
-                sme_cores = 0;
+            if (ok) {
+                if (has_supported_sme_family) {
+                    sme_cores = v;
+                } else {
+                    if (v > 0) {
+                        GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but SME is not supported on this CPU; disabling SME-family kernels\n", v);
+                    }
+                    sme_cores = 0;
+                }
             } else {
-                sme_cores = v;
+                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; using automatic SME thread cap\n");
             }
-        } else {
-            detected_smcus = detect_num_smcus();
-            sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
         }
 
-        if (!sme_env_set && ggml_cpu_has_sme() && sme_cores == 0) {
-            GGML_LOG_WARN("kleidiai: runtime SME-core detection returned 0; falling back to NEON\n");
-        }
-
-        if (sme_cores > 0) {
+        if (sme_cores > 0 && has_supported_sme_family) {
             ctx.features |= CPU_FEATURE_SME;
-#if defined(__aarch64__) && defined(__linux__)
-            // ARM guarantees SME2 implies SME, so only check SME2 when SME is enabled.
-            if (getauxval(AT_HWCAP2) & HWCAP2_SME2) {
+            if (runtime_feat.has_sme2) {
                 ctx.features |= CPU_FEATURE_SME2;
             }
-#elif defined(__aarch64__) && defined(__APPLE__)
-            int feat_sme2 = 0;
-            size_t size = sizeof(feat_sme2);
-            if (sysctlbyname("hw.optional.arm.FEAT_SME2", &feat_sme2, &size, NULL, 0) == 0 && feat_sme2) {
-                ctx.features |= CPU_FEATURE_SME2;
-            }
-#endif
         }
 
         // Kernel selection
@@ -297,16 +407,19 @@ static void init_kleidiai_context(void) {
             GGML_LOG_INFO("kleidiai: primary f32 kernel feature %s\n", cpu_feature_to_string(ctx.kernels_f32->required_cpu));
         }
 
-        ctx.sme_thread_cap = (ctx.features & CPU_FEATURE_SME) ? sme_cores : 0;
+        const bool has_selected_sme_family_kernel =
+            (ctx.kernels_q4  && is_sme_family(ctx.kernels_q4->required_cpu)) ||
+            (ctx.kernels_q8  && is_sme_family(ctx.kernels_q8->required_cpu)) ||
+            (ctx.kernels_f32 && is_sme_family(ctx.kernels_f32->required_cpu));
+        ctx.sme_thread_cap = has_selected_sme_family_kernel ? sme_cores : 0;
 
-        if (ctx.features & CPU_FEATURE_SME) {
-            const bool has_sme2 = (ctx.features & CPU_FEATURE_SME2) != CPU_FEATURE_NONE;
+        if (has_selected_sme_family_kernel) {
             if (sme_env_set && sme_env_ok && sme_cores > 0) {
-                GGML_LOG_INFO("kleidiai: SME%s enabled (GGML_KLEIDIAI_SME=%d override)\n",
-                    has_sme2 ? "2" : "", sme_cores);
+                GGML_LOG_INFO("kleidiai: SME enabled (GGML_KLEIDIAI_SME=%d debug override)\n", sme_cores);
+            } else if (sme_cap_detected) {
+                GGML_LOG_INFO("kleidiai: SME enabled (runtime-detected SME thread cap=%d)\n", sme_cores);
             } else {
-                GGML_LOG_INFO("kleidiai: SME%s enabled (runtime-detected SME cores=%d)\n",
-                    has_sme2 ? "2" : "", sme_cores);
+                GGML_LOG_INFO("kleidiai: SME enabled (runtime SME detected, conservative thread cap=%d)\n", sme_cores);
             }
         } else {
             GGML_LOG_INFO("kleidiai: SME disabled\n");
@@ -467,7 +580,7 @@ static int kleidiai_collect_kernel_chain_common(
     }
 
     if (is_sme_family(primary->required_cpu)) {
-        const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~CPU_FEATURE_SME & ~CPU_FEATURE_SME2);
+        const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~(CPU_FEATURE_SME | CPU_FEATURE_SME2));
         if (fallback_mask != CPU_FEATURE_NONE) {
             ggml_kleidiai_kernels * fallback = select_fallback(fallback_mask);
             if (fallback && fallback != primary &&
@@ -1077,13 +1190,14 @@ class tensor_traits : public ggml::cpu::tensor_traits {
         const int ith_total = params->ith;
 
         int sme_slot = -1;
+        int non_sme_slot = -1;
         for (int i = 0; i < runtime_count; ++i) {
             if (is_sme_family(runtime[i].kernels->required_cpu)) {
                 sme_slot = i;
                 break;
             }
         }
-        int non_sme_slot = -1;
+
         for (int i = 0; i < runtime_count; ++i) {
             if (!is_sme_family(runtime[i].kernels->required_cpu)) {
                 non_sme_slot = i;

ggml/src/ggml-feats.h

@@ -0,0 +1,166 @@
+#pragma once
+
+#if defined(__aarch64__) || defined(_M_ARM64)
+
+#if defined(__linux__)
+#include <sys/auxv.h>
+#include <sys/prctl.h>
+
+#if !defined(HWCAP2_SVE2)
+#define HWCAP2_SVE2 (1ULL << 1)
+#endif
+
+#if !defined(HWCAP_FPHP)
+#define HWCAP_FPHP (1 << 9)
+#endif
+
+#if !defined(HWCAP_ASIMDHP)
+#define HWCAP_ASIMDHP (1 << 10)
+#endif
+
+#if !defined(HWCAP2_I8MM)
+#define HWCAP2_I8MM (1ULL << 13)
+#endif
+
+#if !defined(HWCAP_ASIMDDP)
+#define HWCAP_ASIMDDP (1 << 20)
+#endif
+
+#if !defined(HWCAP_SVE)
+#define HWCAP_SVE (1 << 22)
+#endif
+
+#if !defined(HWCAP2_SME)
+#define HWCAP2_SME (1ULL << 23)
+#endif
+
+#if !defined(HWCAP2_SME2)
+#define HWCAP2_SME2 (1ULL << 37)
+#endif
+
+#if !defined(PR_SVE_GET_VL)
+#define PR_SVE_GET_VL 51
+#endif
+
+#if !defined(PR_SVE_VL_LEN_MASK)
+#define PR_SVE_VL_LEN_MASK 0xffff
+#endif
+
+#elif defined(__APPLE__)
+#include <sys/sysctl.h>
+#elif defined(_WIN32)
+#include <windows.h>
+
+#if !defined(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE 43
+#endif
+
+#if !defined(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_SVE_INSTRUCTIONS_AVAILABLE 46
+#endif
+
+#if !defined(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE 47
+#endif
+
+#if !defined(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE 66
+#endif
+
+#if !defined(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE 67
+#endif
+
+#if !defined(PF_ARM_SME_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_SME_INSTRUCTIONS_AVAILABLE 70
+#endif
+
+#if !defined(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE)
+#define PF_ARM_SME2_INSTRUCTIONS_AVAILABLE 71
+#endif
+
+#endif
+
+typedef struct ggml_feats_arch64_runtime {
+    bool has_dotprod;
+    bool has_fp16;
+    bool has_sve;
+    bool has_sve2;
+    bool has_i8mm;
+    bool has_sme;
+    bool has_sme2;
+    int sve_cnt;
+} ggml_feats_arch64_runtime_t;
+
+static inline ggml_feats_arch64_runtime_t ggml_get_aarch64_runtime_features(void) {
+    ggml_feats_arch64_runtime_t runtime_feat = {};
+
+#if defined(__linux__)
+    const unsigned long hwcap  = getauxval(AT_HWCAP);
+    const unsigned long hwcap2 = getauxval(AT_HWCAP2);
+
+    runtime_feat.has_dotprod = !!(hwcap & HWCAP_ASIMDDP);
+    runtime_feat.has_fp16    = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP);;
+    runtime_feat.has_sve     = !!(hwcap & HWCAP_SVE);
+    runtime_feat.has_sve2    = !!(hwcap2 & HWCAP2_SVE2);
+    runtime_feat.has_i8mm    = !!(hwcap2 & HWCAP2_I8MM);
+    runtime_feat.has_sme     = !!(hwcap2 & HWCAP2_SME);
+    runtime_feat.has_sme2    = !!(hwcap2 & HWCAP2_SME2);
+
+    if (runtime_feat.has_sve) {
+        const int vl = prctl(PR_SVE_GET_VL);
+        if (vl >= 0) {
+            runtime_feat.sve_cnt = vl & PR_SVE_VL_LEN_MASK;
+        }
+    }
+#elif defined(__APPLE__)
+    int oldp = 0;
+    size_t size = sizeof(oldp);
+
+    if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_dotprod = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_FP16", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_fp16 = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_SVE", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_sve = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_SVE2", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_sve2 = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_i8mm = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_sme = static_cast<bool>(oldp);
+    }
+
+    if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, nullptr, 0) == 0) {
+        runtime_feat.has_sme2 = static_cast<bool>(oldp);
+    }
+
+    // Apple does not support userspace non-streaming SVE; keep SVE vector length unknown.
+    runtime_feat.sve_cnt = 0;
+#elif defined (_WIN32)
+    runtime_feat.has_dotprod = IsProcessorFeaturePresent(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_fp16    = IsProcessorFeaturePresent(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_sve     = IsProcessorFeaturePresent(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_sve2    = IsProcessorFeaturePresent(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_i8mm    = IsProcessorFeaturePresent(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_sme     = IsProcessorFeaturePresent(PF_ARM_SME_INSTRUCTIONS_AVAILABLE) != 0;
+    runtime_feat.has_sme2    = IsProcessorFeaturePresent(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE) != 0;
+
+    // Windows exposes SVE feature presence, but not the runtime SVE vector length here.
+    runtime_feat.sve_cnt = 0;
+#endif
+
+    return runtime_feat;
+}
+
+#endif // defined(__aarch64__) || defined(_M_ARM64)

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

@@ -4929,8 +4929,13 @@ static bool ggml_opencl_ensure_fa_variant(ggml_backend_opencl_context * backend_
         const int x = (e && e[0]) ? atoi(e) : 0;
         return (x == 8 || x == 16 || x == 32) ? x : 0;   // 0 = per-gen default
     }();
+    // X2E needs 16 to keep per-lane o_acc at 128B (the compiler spills the
+    // kernel-default width); X1E does not spill, but C=16 is still a measured
+    // +28-30% DK128-GQA4 decode win there (X1-85, kv 4096/8192), neutral on
+    // DK64 / GQA1 / quant-KV.
     const int fa_cl_c_gqa4 = fa_cl_c_env ? fa_cl_c_env
-        : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ? 16 : 0);
+        : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ||
+           backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E ? 16 : 0);
     const std::string opts_cl_c_gqa4 = fa_cl_c_gqa4
         ? " -D FA_CL_C=" + std::to_string(fa_cl_c_gqa4) : std::string();
     const std::string fa_cl_c_g8_val = std::to_string(fa_cl_c_gqa4 ? fa_cl_c_gqa4 * 2 : 16);
@@ -7076,6 +7081,19 @@ inline bool enable_adreno_trans_weight(const ggml_backend_opencl_context *backen
     return ((elem_num < 128 * 1024 * 1024) && adreno_kernel && shape_ok);  // max element num: 2**27
 }
 
+inline bool enable_adreno_trans_weight_q5_K(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
+    if (!use_adreno_kernels(backend_ctx, tensor)) {
+        return false;
+    }
+
+    const size_t elem_num = ggml_nelements(tensor);
+    const size_t q_img_width = elem_num / 8;
+    const size_t qh_img_width = elem_num / 16;
+
+    return q_img_width <= backend_ctx->image_max_buffer_size &&
+           qh_img_width <= backend_ctx->image_max_buffer_size;
+}
+
 static inline bool use_flat_gemv_for_large_m_q4_K(const ggml_tensor *tensor) {
     // gemv_noshuffle variant perf drops for large M, use flat variant for large M.
     // threshold is well above typical hidden/FFN dims, but below typical vocab sizes.
@@ -9255,7 +9273,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
 
 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
         cl_kernel kernel = backend_ctx->kernel_convert_block_q5_K;
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {
             kernel = backend_ctx->kernel_convert_block_q5_K_noshuffle;
         }
 #else
@@ -9290,7 +9308,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
 
         tensor->extra = extra;
 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {
 
             int M = tensor->ne[1];
             int K = tensor->ne[0];
@@ -10388,7 +10406,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
             CL_CHECK(clReleaseMemObject(data_device));
             return;
         }
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {
             int M = tensor->ne[1];
             int K = tensor->ne[0];
 
@@ -18928,7 +18946,8 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
         }
 
         // q5_K x fp32
-        if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32) {
+        if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32 &&
+            enable_adreno_trans_weight_q5_K(backend_ctx, src0)) {
             ggml_cl_mul_mat_q5_K_f32_adreno(backend, src0, src1, dst);
             return;
         }

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

@@ -4627,6 +4627,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q5_1], matmul_q5_1_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q8_0], matmul_q8_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q2_K], matmul_q2_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
+        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_TQ2_0], matmul_tq2_0_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q3_K], matmul_q3_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_K], matmul_q4_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q5_K], matmul_q5_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
@@ -4667,6 +4668,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
@@ -4739,6 +4741,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0], matmul_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
 
         CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_K], matmul_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_TQ2_0], matmul_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
         CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q3_K], matmul_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
         CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_K], matmul_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
         CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_K], matmul_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
@@ -4783,6 +4786,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
         CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
         CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
         CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
         CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
         CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
@@ -4873,6 +4877,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_1], matmul_q5_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0], matmul_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_K], matmul_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_TQ2_0], matmul_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q3_K], matmul_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_K], matmul_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_K], matmul_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
@@ -4921,6 +4926,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+            CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
@@ -4968,6 +4974,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_q5_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_q8_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_q2_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+            CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_tq2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_q3_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_q4_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_q5_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
@@ -5047,6 +5054,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         CREATE_MM(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0].f32acc, matmul_q8_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
 
         CREATE_MM(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_K].f32acc, matmul_q2_k_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+        CREATE_MM(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_TQ2_0].f32acc, matmul_tq2_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q3_K].f32acc, matmul_q3_k_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_K].f32acc, matmul_q4_k_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
         CREATE_MM(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_K].f32acc, matmul_q5_k_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
@@ -5094,6 +5102,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             CREATE_MM(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1].f32acc, matmul_id_subgroup_q5_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0].f32acc, matmul_id_subgroup_q8_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K].f32acc, matmul_id_subgroup_q2_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+            CREATE_MM(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0].f32acc, matmul_id_subgroup_tq2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K].f32acc, matmul_id_subgroup_q3_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K].f32acc, matmul_id_subgroup_q4_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
             CREATE_MM(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K].f32acc, matmul_id_subgroup_q5_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
@@ -5123,6 +5132,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             CREATE_MM(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1].f32acc, matmul_id_q5_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0].f32acc, matmul_id_q8_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K].f32acc, matmul_id_q2_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+            CREATE_MM(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0].f32acc, matmul_id_tq2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K].f32acc, matmul_id_q3_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K].f32acc, matmul_id_q4_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
             CREATE_MM(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K].f32acc, matmul_id_q5_k_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
@@ -5226,6 +5236,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q5_1][i], "mul_mat_vec_q5_1_f32_f32", arr_dmmv_q5_1_f32_f32_len[reduc], arr_dmmv_q5_1_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q8_0][i], "mul_mat_vec_q8_0_f32_f32", arr_dmmv_q8_0_f32_f32_len[reduc], arr_dmmv_q8_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_stdq, 1, 1}, {wg_size_subgroup, 1*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q2_K][i], "mul_mat_vec_q2_k_f32_f32", arr_dmmv_q2_k_f32_f32_len[reduc16], arr_dmmv_q2_k_f32_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_TQ2_0][i], "mul_mat_vec_tq2_0_f32_f32", arr_dmmv_tq2_0_f32_f32_len[reduc16], arr_dmmv_tq2_0_f32_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q3_K][i], "mul_mat_vec_q3_k_f32_f32", arr_dmmv_q3_k_f32_f32_len[reduc16], arr_dmmv_q3_k_f32_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_K][i], "mul_mat_vec_q4_k_f32_f32", arr_dmmv_q4_k_f32_f32_len[reduc16], arr_dmmv_q4_k_f32_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q5_K][i], "mul_mat_vec_q5_k_f32_f32", arr_dmmv_q5_k_f32_f32_len[reduc16], arr_dmmv_q5_k_f32_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
@@ -5253,6 +5264,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q5_1][i], "mul_mat_vec_q5_1_f16_f32", arr_dmmv_q5_1_f16_f32_len[reduc], arr_dmmv_q5_1_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q8_0][i], "mul_mat_vec_q8_0_f16_f32", arr_dmmv_q8_0_f16_f32_len[reduc], arr_dmmv_q8_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_stdq, 1, 1}, {wg_size_subgroup, 1*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q2_K][i], "mul_mat_vec_q2_k_f16_f32", arr_dmmv_q2_k_f16_f32_len[reduc16], arr_dmmv_q2_k_f16_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_TQ2_0][i], "mul_mat_vec_tq2_0_f16_f32", arr_dmmv_tq2_0_f16_f32_len[reduc16], arr_dmmv_tq2_0_f16_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q3_K][i], "mul_mat_vec_q3_k_f16_f32", arr_dmmv_q3_k_f16_f32_len[reduc16], arr_dmmv_q3_k_f16_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_K][i], "mul_mat_vec_q4_k_f16_f32", arr_dmmv_q4_k_f16_f32_len[reduc16], arr_dmmv_q4_k_f16_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q5_K][i], "mul_mat_vec_q5_k_f16_f32", arr_dmmv_q5_k_f16_f32_len[reduc16], arr_dmmv_q5_k_f16_f32_data[reduc16], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
@@ -5307,6 +5319,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q5_1], "mul_mat_vec_id_q5_1_f32",       arr_dmmv_id_q5_1_f32_f32_len[reduc],    arr_dmmv_id_q5_1_f32_f32_data[reduc],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size);
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q8_0], "mul_mat_vec_id_q8_0_f32",       arr_dmmv_id_q8_0_f32_f32_len[reduc],    arr_dmmv_id_q8_0_f32_f32_data[reduc],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {1*rm_stdq, 1, 1}, {wg_size_subgroup, 1*rm_stdq}, 1, true, use_subgroups, force_subgroup_size);
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q2_K], "mul_mat_vec_id_q2_k_f32",       arr_dmmv_id_q2_k_f32_f32_len[reduc16],    arr_dmmv_id_q2_k_f32_f32_data[reduc16],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq}, 1, true, use_subgroups16, force_subgroup_size16);
+        ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_TQ2_0], "mul_mat_vec_id_tq2_0_f32",     arr_dmmv_id_tq2_0_f32_f32_len[reduc16],   arr_dmmv_id_tq2_0_f32_f32_data[reduc16],   "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq}, 1, true, use_subgroups16, force_subgroup_size16);
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q3_K], "mul_mat_vec_id_q3_k_f32",       arr_dmmv_id_q3_k_f32_f32_len[reduc16],    arr_dmmv_id_q3_k_f32_f32_data[reduc16],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq}, 1, true, use_subgroups16, force_subgroup_size16);
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q4_K], "mul_mat_vec_id_q4_k_f32",       arr_dmmv_id_q4_k_f32_f32_len[reduc16],    arr_dmmv_id_q4_k_f32_f32_data[reduc16],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq}, 1, true, use_subgroups16, force_subgroup_size16);
         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q5_K], "mul_mat_vec_id_q5_k_f32",       arr_dmmv_id_q5_k_f32_f32_len[reduc16],    arr_dmmv_id_q5_k_f32_f32_data[reduc16],    "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_kq, 1, 1}, {wg_size_subgroup16, rm_kq}, 1, true, use_subgroups16, force_subgroup_size16);
@@ -5368,6 +5381,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q5_1], "dequant_q5_1", dequant_q5_1_len, dequant_q5_1_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q8_0], "dequant_q8_0", dequant_q8_0_len, dequant_q8_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q2_K], "dequant_q2_k", dequant_q2_k_len, dequant_q2_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_TQ2_0], "dequant_tq2_0", dequant_tq2_0_len, dequant_tq2_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q3_K], "dequant_q3_k", dequant_q3_k_len, dequant_q3_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_K], "dequant_q4_k", dequant_q4_k_len, dequant_q4_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 32, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q5_K], "dequant_q5_k", dequant_q5_k_len, dequant_q5_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
@@ -5396,6 +5410,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q5_1], "get_rows_q5_1", get_rows_q5_1_len, get_rows_q5_1_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q8_0], "get_rows_q8_0", get_rows_q8_0_len, get_rows_q8_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q2_K], "get_rows_q2_k", get_rows_q2_k_len, get_rows_q2_k_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_TQ2_0], "get_rows_tq2_0", get_rows_tq2_0_len, get_rows_tq2_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q3_K], "get_rows_q3_k", get_rows_q3_k_len, get_rows_q3_k_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_K], "get_rows_q4_k", get_rows_q4_k_len, get_rows_q4_k_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q5_K], "get_rows_q5_k", get_rows_q5_k_len, get_rows_q5_k_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
@@ -5424,6 +5439,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q5_1], "get_rows_q5_1_f32", get_rows_q5_1_f32_len, get_rows_q5_1_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q8_0], "get_rows_q8_0_f32", get_rows_q8_0_f32_len, get_rows_q8_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q2_K], "get_rows_q2_k_f32", get_rows_q2_k_f32_len, get_rows_q2_k_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_TQ2_0], "get_rows_tq2_0_f32", get_rows_tq2_0_f32_len, get_rows_tq2_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q3_K], "get_rows_q3_k_f32", get_rows_q3_k_f32_len, get_rows_q3_k_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_K], "get_rows_q4_k_f32", get_rows_q4_k_f32_len, get_rows_q4_k_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q5_K], "get_rows_q5_k_f32", get_rows_q5_k_f32_len, get_rows_q5_k_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
@@ -7638,6 +7654,7 @@ static vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_TQ2_0:
             break;
         default:
             return nullptr;
@@ -7712,6 +7729,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_TQ2_0:
             break;
         default:
             return nullptr;
@@ -7781,6 +7799,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context *
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_TQ2_0:
             break;
         default:
             return nullptr;
@@ -7874,6 +7893,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_TQ2_0:
             break;
         default:
             return nullptr;
@@ -7946,6 +7966,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec_id(ggml_backend_vk_context
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_TQ2_0:
             break;
         default:
             return nullptr;
@@ -18014,6 +18035,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                     case GGML_TYPE_IQ4_NL:
                     case GGML_TYPE_MXFP4:
                     case GGML_TYPE_NVFP4:
+                    case GGML_TYPE_TQ2_0:
                         break;
                     default:
                         return false;
@@ -18119,6 +18141,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                     case GGML_TYPE_IQ4_NL:
                     case GGML_TYPE_MXFP4:
                     case GGML_TYPE_NVFP4:
+                    case GGML_TYPE_TQ2_0:
                     case GGML_TYPE_I32:
                         return true;
                     default:

ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl

@@ -608,6 +608,20 @@ vec2 get_dm(uint ib, uint a_offset) {
 }
 #endif
 
+#if defined(DATA_A_TQ2_0)
+vec2 dequantize(uint ib, uint iqs, uint a_offset) {
+    // elem e -> byte qs[(e/128)*32 + e%32], bits 2*((e%128)/32); w = q - 1 (d applied via get_dm)
+    const uint qsi   = (iqs / 128) * 32 + (iqs % 32);  // iqs even -> qsi, qsi+1 in same group/level
+    const uint shift = 2 * ((iqs % 128) / 32);
+
+    const uvec2 qs = uvec2(data_a[a_offset + ib].qs[qsi], data_a[a_offset + ib].qs[qsi + 1]);
+    return vec2((qs >> shift) & 3) - 1.0;
+}
+vec2 get_dm(uint ib, uint a_offset) {
+    return vec2(float(data_a[a_offset + ib].d), 0);
+}
+#endif
+
 #if defined(DATA_A_Q3_K)
 vec2 dequantize(uint ib, uint iqs, uint a_offset) {
     iqs /= 2;

ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs_cm2.glsl

@@ -247,6 +247,44 @@ f16vec4 dequantFuncQ8_0_v(const in decodeBufQ8_0 bl, const in uint blockCoords[2
     return f16vec4(vec4(qi) * vec4(float(d)));
 }
 
+layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0 {
+   block_tq2_0 block;
+};
+
+layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0_packed16 {
+   block_tq2_0_packed16 block;
+};
+
+float16_t dequantFuncTQ2_0(const in decodeBufTQ2_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
+{
+    decodeBufTQ2_0_packed16 bl16 = decodeBufTQ2_0_packed16(bl);
+    const uint idx = coordInBlock[1];
+
+    const uint qsshift = (idx & 0x60) >> 4;             // 0,2,4,6
+
+    uint qs = uint32_t(bl16.block.qs[((idx & 0x80) >> 3) + ((idx & 0x1E) >> 1)]);
+    qs = (qs >> qsshift) & 0x0303;
+    qs = unpack8(qs)[idx & 1];
+
+    return bl.block.d * (float16_t(int(qs)) - float16_t(1.0));
+}
+
+f16vec4 dequantFuncTQ2_0_v(const in decodeBufTQ2_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
+{
+    const uint idx = coordInBlock[1];
+
+    const uint qsshift = (idx & 0x60) >> 4;             // 0,2,4,6
+    const uint qsi = ((idx & 0x80) >> 2) + (idx & 0x1C); // byte index of 4-aligned group
+
+    const uint qsw = (uint(bl.block.qs[qsi]))
+                   | (uint(bl.block.qs[qsi + 1]) << 8)
+                   | (uint(bl.block.qs[qsi + 2]) << 16)
+                   | (uint(bl.block.qs[qsi + 3]) << 24);
+    const u8vec4 q = unpack8((qsw >> qsshift) & 0x03030303);
+
+    return bl.block.d * (f16vec4(q) - f16vec4(1.0));
+}
+
 layout(buffer_reference, std430, buffer_reference_align = 4) buffer decodeBufQ2_K {
    block_q2_K block;
 };
@@ -1368,6 +1406,9 @@ f16vec4 dequantFuncNVFP4_v(const in decodeBufNVFP4 bl, const in uint blockCoords
 #elif defined(DATA_A_Q8_0)
 #define dequantFuncA dequantFuncQ8_0
 #define dequantFuncA_v dequantFuncQ8_0_v
+#elif defined(DATA_A_TQ2_0)
+#define dequantFuncA dequantFuncTQ2_0
+#define dequantFuncA_v dequantFuncTQ2_0_v
 #elif defined(DATA_A_Q2_K)
 #define dequantFuncA dequantFuncQ2_K
 #define dequantFuncA_v dequantFuncQ2_K_v

ggml/src/ggml-vulkan/vulkan-shaders/dequant_tq2_0.comp

@@ -0,0 +1,31 @@
+#version 450
+
+#include "dequant_head.glsl"
+
+layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
+
+layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
+layout (binding = 1) writeonly buffer D {D_TYPE data_b[];};
+
+void main() {
+    [[unroll]] for (uint wgy = 0; wgy < 256; wgy++) {
+        const uint i = gl_WorkGroupID.x * 256 + wgy;
+        if (i >= p.nel / QUANT_K) {
+            return;
+        }
+
+        const uint tid = gl_LocalInvocationID.x;
+        const uint ip = tid / 32;         // group 0,1 (128 elems each)
+        const uint il = tid - 32 * ip;    // byte in group 0..31
+
+        const uint y_idx = i * QUANT_K + 128 * ip + il;
+
+        const uint8_t qs = data_a[i].qs[32 * ip + il];
+
+        const FLOAT_TYPE d = FLOAT_TYPE(data_a[i].d);
+        data_b[y_idx +  0] = D_TYPE(d * FLOAT_TYPE(int((qs >> 0) & 3) - 1));
+        data_b[y_idx + 32] = D_TYPE(d * FLOAT_TYPE(int((qs >> 2) & 3) - 1));
+        data_b[y_idx + 64] = D_TYPE(d * FLOAT_TYPE(int((qs >> 4) & 3) - 1));
+        data_b[y_idx + 96] = D_TYPE(d * FLOAT_TYPE(int((qs >> 6) & 3) - 1));
+    }
+}

ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_tq2_0.comp

@@ -0,0 +1,102 @@
+#version 450
+#extension GL_EXT_shader_explicit_arithmetic_types_int32 : require
+
+#include "mul_mat_vec_base.glsl"
+
+layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
+
+FLOAT_TYPE temp[NUM_COLS][NUM_ROWS];
+
+// ternary TQ2_0: w = (q - 1) * d. Same qs group/level layout as q2_K, but a
+// single f16 scale per 256-block and no mins:
+//   sum_e b_e * (q_e - 1) * d  =  d * (sum_e b_e * q_e  -  sum_e b_e)
+void calc_superblock(const uint a_offset, const uint b_offset, const uint v_im, const uint q_offset, const uint y_offset, const uint i, const uint num_blocks_per_row, const uint first_row, const uint num_rows) {
+    const uint y_idx = i * QUANT_K + y_offset;
+
+    [[unroll]] for (uint n = 0; n < num_rows; ++n) {
+        const uint ib0 = a_offset + (first_row+n)*num_blocks_per_row;
+        if (i >= num_blocks_per_row) {
+            continue;
+        }
+
+        const uint32_t qs_u32 = uint32_t(data_a_packed16[ib0 + i].qs[q_offset / 2]) | (uint32_t(data_a_packed16[ib0 + i].qs[q_offset / 2 + 8]) << 16);
+        const vec4 qs_u32_0 = vec4(unpack8(qs_u32 & 0x03030303));
+        const vec4 qs_u32_2 = vec4(unpack8((qs_u32 >> 2) & 0x03030303));
+        const vec4 qs_u32_4 = vec4(unpack8((qs_u32 >> 4) & 0x03030303));
+        const vec4 qs_u32_6 = vec4(unpack8((qs_u32 >> 6) & 0x03030303));
+
+        const FLOAT_TYPE d = FLOAT_TYPE(data_a[ib0 + i].d);
+
+        [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
+            vec2 b0 =   vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 +  0]);
+            vec2 b16 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 +  8]);
+            vec2 b32 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 16]);
+            vec2 b48 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 24]);
+            vec2 b64 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 32]);
+            vec2 b80 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 40]);
+            vec2 b96 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 48]);
+            vec2 b112 = vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 56]);
+
+            FLOAT_TYPE sumq = FLOAT_TYPE(0.0);
+            FLOAT_TYPE sumb = FLOAT_TYPE(0.0);
+            [[unroll]] for (int l = 0; l < 2; ++l) {
+                sumq = fma(FLOAT_TYPE(b0[l]),   FLOAT_TYPE(qs_u32_0[l  ]),
+                       fma(FLOAT_TYPE(b16[l]),  FLOAT_TYPE(qs_u32_0[l+2]),
+                       fma(FLOAT_TYPE(b32[l]),  FLOAT_TYPE(qs_u32_2[l  ]),
+                       fma(FLOAT_TYPE(b48[l]),  FLOAT_TYPE(qs_u32_2[l+2]),
+                       fma(FLOAT_TYPE(b64[l]),  FLOAT_TYPE(qs_u32_4[l  ]),
+                       fma(FLOAT_TYPE(b80[l]),  FLOAT_TYPE(qs_u32_4[l+2]),
+                       fma(FLOAT_TYPE(b96[l]),  FLOAT_TYPE(qs_u32_6[l  ]),
+                       fma(FLOAT_TYPE(b112[l]), FLOAT_TYPE(qs_u32_6[l+2]), sumq))))))));
+                sumb += FLOAT_TYPE(b0[l]) + FLOAT_TYPE(b16[l]) + FLOAT_TYPE(b32[l]) + FLOAT_TYPE(b48[l])
+                      + FLOAT_TYPE(b64[l]) + FLOAT_TYPE(b80[l]) + FLOAT_TYPE(b96[l]) + FLOAT_TYPE(b112[l]);
+            }
+            temp[j][n] = fma(d, sumq - sumb, temp[j][n]);
+        }
+    }
+}
+
+void compute_outputs(const uint32_t first_row, const uint32_t num_rows) {
+    uint a_offset, b_offset, d_offset;
+    get_offsets(a_offset, b_offset, d_offset);
+
+    const uint num_blocks_per_row = p.ncols / QUANT_K;
+
+    // 16 threads are used to process each block
+    const uint it_size = gl_WorkGroupSize.x/16;
+    const uint tid = gl_LocalInvocationID.x;
+    const uint itid = tid%16;  // 0...15
+    const uint ix = tid/16;
+
+    const uint v_im = itid/8;                                // 0 or 1. 0 computes 0..., 1 computes 128...
+    const uint v_in = itid - 8*v_im;                         // 0...7
+
+    const uint l0 = 2*v_in;                                  // 0...15
+    const uint q_offset = 32*v_im + l0;
+    const uint y_offset = 128*v_im + l0;
+
+    [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
+        [[unroll]] for (uint i = 0; i < NUM_ROWS; ++i) {
+            temp[j][i] = FLOAT_TYPE(0);
+        }
+    }
+
+    for (uint i0 = 0; i0 < num_blocks_per_row; i0 += it_size)
+        calc_superblock(a_offset, b_offset, v_im, q_offset, y_offset, i0 + ix, num_blocks_per_row, first_row, num_rows);
+
+    reduce_result(temp, d_offset, first_row, num_rows, tid);
+}
+
+void main() {
+    const uint first_row = NUM_ROWS * (gl_WorkGroupID.x + gl_NumWorkGroups.x * gl_WorkGroupID.z);
+
+    // do NUM_ROWS at a time, unless there aren't enough remaining rows
+    if (first_row + NUM_ROWS <= p.stride_d) {
+        compute_outputs(first_row, NUM_ROWS);
+    } else {
+        if (first_row >= p.stride_d) {
+            return;
+        }
+        compute_outputs(first_row, p.stride_d - first_row);
+    }
+}

ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl

@@ -182,6 +182,22 @@ void load_a_to_shmem(const uint pos_a, const uint row, const uint col, const uin
 
             buf_a[buf_idx    ] = FLOAT_TYPEV2(v.xy);
             buf_a[buf_idx + 1] = FLOAT_TYPEV2(v.zw);
+#elif defined(DATA_A_TQ2_0)
+            const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row;
+            const uint buf_idx = col * SHMEM_STRIDE + row * LOAD_VEC_A / 2;
+
+            const uint ib = idx / 128;                         // 2 values per idx
+            const uint iqs = (idx % 128) * 2;                  // elem 0,2,4..254
+
+            const uint qsi   = (iqs / 128) * 32 + (iqs % 32);  // byte pair start
+            const uint shift = 2 * ((iqs % 128) / 32);         // 0,2,4,6
+
+            const uvec2 qs = uvec2(data_a[ib].qs[qsi], data_a[ib].qs[qsi + 1]);
+            const float d = float(data_a[ib].d);
+
+            const vec2 v = d * (vec2((qs >> shift) & 3) - 1.0);
+
+            buf_a[buf_idx] = FLOAT_TYPEV2(v.xy);
 #elif defined(DATA_A_Q3_K)
             const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row;
             const uint buf_idx = col * SHMEM_STRIDE + row * LOAD_VEC_A / 2;

ggml/src/ggml-vulkan/vulkan-shaders/types.glsl

@@ -303,6 +303,30 @@ struct block_q2_K_packed32
 #define DATA_A_QUANT_K
 #endif
 
+#define QUANT_K_TQ2_0 256
+
+// ternary (BitNet): 2-bit codes, w = (q - 1) * d; qs layout matches q2_K's
+// two 32-byte groups with four bit-levels per byte
+struct block_tq2_0
+{
+    uint8_t qs[QUANT_K_TQ2_0/4];
+    float16_t d;
+};
+
+struct block_tq2_0_packed16
+{
+    uint16_t qs[QUANT_K_TQ2_0/4/2];
+    float16_t d;
+};
+
+#if defined(DATA_A_TQ2_0)
+#define QUANT_K QUANT_K_TQ2_0
+#define QUANT_R 1
+#define A_TYPE block_tq2_0
+#define A_TYPE_PACKED16 block_tq2_0_packed16
+#define DATA_A_QUANT_K
+#endif
+
 #define QUANT_K_Q3_K 256
 
 struct block_q3_K

ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -72,6 +72,7 @@ const std::vector<std::string> type_names = {
     "iq4_nl",
     "mxfp4",
     "nvfp4",
+    "tq2_0",
     "bf16",
 };
 
@@ -733,7 +734,7 @@ void process_shaders() {
     for (const auto& tname : type_names) {
         // mul mat vec
         std::string data_a_key = "DATA_A_" + to_uppercase(tname);
-        std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_")) ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";
+        std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_") || tname == "tq2_0") ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";
 
         string_to_spv("mul_mat_vec_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}));
         string_to_spv("mul_mat_vec_" + tname + "_f16_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float16_t"}, {"B_TYPEV2", "f16vec2"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}));

ggml/src/gguf.cpp

@@ -611,6 +611,13 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr
         GGML_ASSERT(int64_t(ctx->kv.size()) == n_kv);
 
         const int alignment_idx = gguf_find_key(ctx, GGUF_KEY_GENERAL_ALIGNMENT);
+        if (alignment_idx != -1 && gguf_get_kv_type(ctx, alignment_idx) != GGUF_TYPE_UINT32) {
+            GGML_LOG_ERROR("%s: key '%s' must be of type %s but is %s\n",
+                __func__, GGUF_KEY_GENERAL_ALIGNMENT, gguf_type_name(GGUF_TYPE_UINT32),
+                gguf_type_name(gguf_get_kv_type(ctx, alignment_idx)));
+            gguf_free(ctx);
+            return nullptr;
+        }
         ctx->alignment = alignment_idx == -1 ? GGUF_DEFAULT_ALIGNMENT : gguf_get_val_u32(ctx, alignment_idx);
 
         if (ctx->alignment == 0 || (ctx->alignment & (ctx->alignment - 1)) != 0) {
@@ -682,9 +689,11 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr
             }
 
             // check that the total number of elements is representable
-            if (ok && ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) ||
-                       (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) ||
-                       (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) {
+            // (a zero-element tensor is trivially representable; the guard also avoids a division by zero below)
+            if (ok && ggml_nelements(&info.t) > 0 &&
+                ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) ||
+                 (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) ||
+                 (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) {
 
                 GGML_LOG_ERROR("%s: total number of elements in tensor '%s' with shape "
                     "(%" PRIi64 ", %" PRIi64 ", %" PRIi64 ", %" PRIi64 ") is >= %" PRIi64 "\n",

include/llama.h

@@ -457,6 +457,8 @@ extern "C" {
     // lora adapter
     struct llama_adapter_lora;
 
+    LLAMA_API const char * llama_version(void);
+
     // Helpers for getting default parameters
     // TODO: update API to start accepting pointers to params structs (https://github.com/ggml-org/llama.cpp/discussions/9172)
     LLAMA_API struct llama_model_params          llama_model_default_params(void);
@@ -883,6 +885,7 @@ extern "C" {
                const llama_token * tokens,
                           size_t   n_token_count);
 
+    // If tokens_out is NULL, only the token count is reported through n_token_count_out and no state is loaded
     LLAMA_API size_t llama_state_seq_load_file(
             struct llama_context * ctx,
                       const char * filepath,

scripts/make-release-checks.sh

@@ -0,0 +1,83 @@
+#!/bin/bash
+# Run all pre-release checks and determine the release version.
+#
+# Usage: make-release-checks.sh [--dry-run]
+#   --dry-run: warn on failures instead of aborting
+#
+# Env (when running in GitHub Actions): GH_TOKEN, GITHUB_REPOSITORY, GITHUB_OUTPUT
+set -euo pipefail
+
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)"
+
+DRY_RUN=false
+for arg in "$@"; do
+    case "$arg" in
+        --dry-run) DRY_RUN=true ;;
+        *) echo "Unknown argument: $arg"; exit 1 ;;
+    esac
+done
+
+MAJOR=$(grep "set(LLAMA_VERSION_MAJOR" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
+MINOR=$(grep "set(LLAMA_VERSION_MINOR" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
+PATCH=$(grep "set(LLAMA_VERSION_PATCH" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
+VERSION="v${MAJOR}.${MINOR}.${PATCH}"
+echo "Determined version: ${VERSION}"
+if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
+    echo "version=${VERSION}" >> "$GITHUB_OUTPUT"
+fi
+
+echo "Checking that tag ${VERSION} does not already exist..."
+if git ls-remote --tags origin "${VERSION}" | grep -q "${VERSION}"; then
+    echo "Error: tag ${VERSION} already exists on remote"
+    exit 1
+fi
+echo "Tag ${VERSION} does not exist on remote - OK"
+
+SHA=$(git rev-parse HEAD)
+echo "Checking release.yml status for commit ${SHA}..."
+if [[ -z "${GITHUB_REPOSITORY:-}" ]]; then
+    echo "Warning: GITHUB_REPOSITORY not set - skipping CI check (local run)"
+else
+    RUNS=$(gh api "repos/${GITHUB_REPOSITORY}/actions/workflows/release.yml/runs" \
+        --jq "[.workflow_runs[] | select(.head_sha == \"${SHA}\" and .conclusion == \"success\")] | length")
+    if [[ "$RUNS" -eq 0 ]]; then
+        if [[ "$DRY_RUN" == "true" ]]; then
+            echo "Warning: no successful release.yml run found for HEAD (${SHA}) (dry run, continuing)."
+        else
+            echo "Error: no successful release.yml run found for HEAD (${SHA})"
+            echo "The nightly build must complete successfully before making a release."
+            exit 1
+        fi
+    else
+        echo "Found successful release.yml run for HEAD."
+    fi
+fi
+
+MAJOR=$(grep "set(GGML_VERSION_MAJOR" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
+MINOR=$(grep "set(GGML_VERSION_MINOR" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
+PATCH=$(grep "set(GGML_VERSION_PATCH" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
+GGML_VERSION="v${MAJOR}.${MINOR}.${PATCH}"
+echo "Local ggml version: ${GGML_VERSION}"
+
+if ! git clone --depth 1 --branch "${GGML_VERSION}" https://github.com/ggml-org/ggml.git upstream-ggml 2>/dev/null; then
+    echo "Warning: tag ${GGML_VERSION} not found in upstream ggml - skipping comparison"
+else
+    echo "Comparing local ggml/ src and include with upstream ${GGML_VERSION}..."
+    DIFF=$(diff -rq "$REPO_ROOT/ggml/src"          upstream-ggml/src          2>&1 || true)
+    DIFF+=$(diff -rq "$REPO_ROOT/ggml/include"     upstream-ggml/include      2>&1 || true)
+    DIFF+=$(diff     "$REPO_ROOT/ggml/CMakeLists.txt" upstream-ggml/CMakeLists.txt 2>&1 || true)
+    rm -rf upstream-ggml
+    if [[ -n "$DIFF" ]]; then
+        echo "local ggml/ differs from upstream ${GGML_VERSION}:"
+        echo "$DIFF"
+        if [[ "$DRY_RUN" == "true" ]]; then
+            echo "Warning: would abort release due to ggml mismatch (dry run, continuing)."
+        else
+            echo "Error: ggml must match upstream before making a release."
+            exit 1
+        fi
+    else
+        echo "local ggml/ matches upstream ${GGML_VERSION}"
+    fi
+fi

scripts/sync-ggml.last

@@ -1 +1 @@
-30bf8685ed4eb0a47f2b06229543327749904150
+8846b79e66747bb9f68597420e95114c177315ce

src/CMakeLists.txt

@@ -45,11 +45,16 @@ add_library(llama
             )
 
 set_target_properties(llama PROPERTIES
-    VERSION ${LLAMA_INSTALL_VERSION}
-    SOVERSION 0
+    VERSION ${LLAMA_VERSION_BASE}
+    SOVERSION ${LLAMA_VERSION_MAJOR}
     MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
 )
 
+target_compile_definitions(llama PRIVATE
+    LLAMA_VERSION="${LLAMA_VERSION}"
+    LLAMA_COMMIT="${LLAMA_BUILD_COMMIT}"
+)
+
 target_include_directories(llama PRIVATE .)
 target_include_directories(llama PUBLIC ../include)
 target_compile_features   (llama PRIVATE cxx_std_17) # don't bump

src/llama-context.cpp

@@ -3110,6 +3110,17 @@ size_t llama_context::state_seq_load_file(llama_seq_id seq_id, const char * file
     {
         const uint32_t n_token_count = file.read_u32();
 
+        if (tokens_out == nullptr) {
+            const size_t n_token_max = (file.size() - file.tell()) / sizeof(llama_token);
+            if (n_token_count > n_token_max) {
+                LLAMA_LOG_ERROR("%s: token count in sequence state file exceeds the file size! %u > %zu\n", __func__, n_token_count, n_token_max);
+                return 0;
+            }
+
+            *n_token_count_out = n_token_count;
+            return file.tell();
+        }
+
         if (n_token_count > n_token_capacity) {
             LLAMA_LOG_ERROR("%s: token count in sequence state file exceeded capacity! %u > %zu\n", __func__, n_token_count, n_token_capacity);
             return 0;

src/llama-model.cpp

@@ -1122,6 +1122,9 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
 
         ml.get_key(LLM_KV_CONVNEXT_EMBEDDING_LENGTH, hparams.convnext.n_embd);
         ml.get_key(LLM_KV_CONVNEXT_BLOCK_COUNT,      hparams.convnext.n_layer);
+
+        GGML_ASSERT(hparams.posnet.n_layer   <= hparams.n_layer_all);
+        GGML_ASSERT(hparams.convnext.n_layer <= hparams.n_layer_all);
     }
 
     GGML_ASSERT(hparams.n_expert <= LLAMA_MAX_EXPERTS);

src/llama.cpp

@@ -114,6 +114,10 @@ bool llama_supports_rpc(void) {
     return ggml_backend_reg_by_name("RPC") != nullptr;
 }
 
+const char * llama_version(void) {
+    return LLAMA_VERSION;
+}
+
 void llama_backend_init(void) {
     ggml_time_init();
 

src/models/dflash.cpp

@@ -66,7 +66,7 @@ void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) {
     // DFlash has a single rope, so the SWA rope == main rope.
     if (ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false) && hparams.n_swa > 0) {
         hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
-        ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
+        ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
         hparams.rope_freq_base_train_swa  = hparams.rope_freq_base_train;
         hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train;
     }

tests/test-backend-ops.cpp

@@ -8000,7 +8000,8 @@ static const ggml_type all_types[] = {
     GGML_TYPE_Q2_K, GGML_TYPE_Q3_K,
     GGML_TYPE_Q4_K, GGML_TYPE_Q5_K,
     GGML_TYPE_Q6_K,
-    // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends
+    GGML_TYPE_TQ2_0,
+    // GGML_TYPE_TQ1_0, // TODO: implement for all backends
     GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S,
     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
     GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,
@@ -8027,7 +8028,8 @@ static const ggml_type other_types[] = {
     GGML_TYPE_Q2_K, GGML_TYPE_Q3_K,
     GGML_TYPE_Q5_K,
     GGML_TYPE_Q6_K,
-    // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends
+    GGML_TYPE_TQ2_0,
+    // GGML_TYPE_TQ1_0, // TODO: implement for all backends
     GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S,
     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
     GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,

tests/test-gguf.cpp

@@ -31,11 +31,13 @@ enum handcrafted_file_type {
     // HANDCRAFTED_KV_BAD_VALUE_SIZE          =  30 + offset_has_kv, // removed because it can result in allocations > 1 TB (default sanitizer limit)
     HANDCRAFTED_KV_DUPLICATE_KEY           =  40 + offset_has_kv,
     HANDCRAFTED_KV_BAD_ALIGN               =  50 + offset_has_kv,
+    HANDCRAFTED_KV_WRONG_TYPE_ALIGN        =  55 + offset_has_kv,
     HANDCRAFTED_KV_SUCCESS                 = 800 + offset_has_kv,
 
     HANDCRAFTED_TENSORS_BAD_NAME_SIZE      =  10 + offset_has_tensors,
     HANDCRAFTED_TENSORS_BAD_N_DIMS         =  20 + offset_has_tensors,
     HANDCRAFTED_TENSORS_BAD_SHAPE          =  30 + offset_has_tensors,
+    HANDCRAFTED_TENSORS_ZERO_DIM           =  35 + offset_has_tensors,
     HANDCRAFTED_TENSORS_NE_TOO_BIG         =  40 + offset_has_tensors,
     HANDCRAFTED_TENSORS_NBYTES_TOO_BIG     =  45 + offset_has_tensors,
     HANDCRAFTED_TENSORS_BAD_TYPE           =  50 + offset_has_tensors,
@@ -69,11 +71,13 @@ static std::string handcrafted_file_type_name(const enum handcrafted_file_type h
         case HANDCRAFTED_KV_BAD_TYPE:                return "KV_BAD_TYPE";
         case HANDCRAFTED_KV_DUPLICATE_KEY:           return "KV_DUPLICATE_KEY";
         case HANDCRAFTED_KV_BAD_ALIGN:               return "KV_BAD_ALIGN";
+        case HANDCRAFTED_KV_WRONG_TYPE_ALIGN:        return "KV_WRONG_TYPE_ALIGN";
         case HANDCRAFTED_KV_SUCCESS:                 return "KV_RANDOM_KV";
 
         case HANDCRAFTED_TENSORS_BAD_NAME_SIZE:      return "TENSORS_BAD_NAME_SIZE";
         case HANDCRAFTED_TENSORS_BAD_N_DIMS:         return "TENSORS_BAD_N_DIMS";
         case HANDCRAFTED_TENSORS_BAD_SHAPE:          return "TENSORS_BAD_SHAPE";
+        case HANDCRAFTED_TENSORS_ZERO_DIM:           return "TENSORS_ZERO_DIM";
         case HANDCRAFTED_TENSORS_NE_TOO_BIG:         return "TENSORS_NE_TOO_BIG";
         case HANDCRAFTED_TENSORS_NBYTES_TOO_BIG:     return "TENSORS_NBYTES_TOO_BIG";
         case HANDCRAFTED_TENSORS_BAD_TYPE:           return "TENSORS_BAD_TYPE";
@@ -95,6 +99,9 @@ static std::string handcrafted_file_type_name(const enum handcrafted_file_type h
 }
 
 static bool expect_context_not_null(const enum handcrafted_file_type hft) {
+    if (hft == HANDCRAFTED_TENSORS_ZERO_DIM) {
+        return true;
+    }
     if (hft < offset_has_kv) {
         return hft >= HANDCRAFTED_HEADER_EMPTY;
     }
@@ -257,9 +264,9 @@ static FILE * get_handcrafted_file(const unsigned int seed, const enum handcraft
     }
     {
         uint64_t n_kv = kv_types.size();
-        if (hft == HANDCRAFTED_KV_BAD_ALIGN      ||
-            hft == HANDCRAFTED_TENSORS_BAD_ALIGN || hft == HANDCRAFTED_TENSORS_CUSTOM_ALIGN ||
-            hft == HANDCRAFTED_DATA_BAD_ALIGN    || hft == HANDCRAFTED_DATA_CUSTOM_ALIGN) {
+        if (hft == HANDCRAFTED_KV_BAD_ALIGN        || hft == HANDCRAFTED_KV_WRONG_TYPE_ALIGN ||
+            hft == HANDCRAFTED_TENSORS_BAD_ALIGN   || hft == HANDCRAFTED_TENSORS_CUSTOM_ALIGN ||
+            hft == HANDCRAFTED_DATA_BAD_ALIGN      || hft == HANDCRAFTED_DATA_CUSTOM_ALIGN) {
 
             n_kv += 1;
         } else if (hft == HANDCRAFTED_HEADER_BAD_N_KV) {
@@ -344,15 +351,17 @@ static FILE * get_handcrafted_file(const unsigned int seed, const enum handcraft
         helper_write(file, data, hft == HANDCRAFTED_KV_BAD_TYPE ? 1 : gguf_type_size(type));
     }
 
-    if (hft == HANDCRAFTED_KV_BAD_ALIGN      ||
-        hft == HANDCRAFTED_TENSORS_BAD_ALIGN || hft == HANDCRAFTED_TENSORS_CUSTOM_ALIGN ||
-        hft == HANDCRAFTED_DATA_BAD_ALIGN    || hft == HANDCRAFTED_DATA_CUSTOM_ALIGN) {
+    if (hft == HANDCRAFTED_KV_BAD_ALIGN        || hft == HANDCRAFTED_KV_WRONG_TYPE_ALIGN ||
+        hft == HANDCRAFTED_TENSORS_BAD_ALIGN   || hft == HANDCRAFTED_TENSORS_CUSTOM_ALIGN ||
+        hft == HANDCRAFTED_DATA_BAD_ALIGN      || hft == HANDCRAFTED_DATA_CUSTOM_ALIGN) {
 
         const uint64_t n = strlen(GGUF_KEY_GENERAL_ALIGNMENT);
         helper_write(file, n);
         helper_write(file, GGUF_KEY_GENERAL_ALIGNMENT, n);
 
-        const int32_t type = gguf_type(GGUF_TYPE_UINT32);
+        // HANDCRAFTED_KV_WRONG_TYPE_ALIGN declares general.alignment with a non-UINT32 type,
+        // which the loader must reject cleanly instead of aborting on an assertion
+        const int32_t type = hft == HANDCRAFTED_KV_WRONG_TYPE_ALIGN ? int32_t(GGUF_TYPE_INT32) : int32_t(GGUF_TYPE_UINT32);
         helper_write(file, type);
 
         alignment = expect_context_not_null(hft) ? 1 : 13;
@@ -403,6 +412,9 @@ static FILE * get_handcrafted_file(const unsigned int seed, const enum handcraft
                 break;
             }
         }
+        if (hft == HANDCRAFTED_TENSORS_ZERO_DIM) {
+            n_dims = 2;
+        }
         if (hft == HANDCRAFTED_TENSORS_BAD_N_DIMS) {
             const uint32_t n_dims_bad = GGML_MAX_DIMS + 1;
             helper_write(file, n_dims_bad);
@@ -415,6 +427,9 @@ static FILE * get_handcrafted_file(const unsigned int seed, const enum handcraft
             for (uint32_t j = 0; j < n_dims; ++j) {
                 helper_write(file, bad_dim);
             }
+        } else if (hft == HANDCRAFTED_TENSORS_ZERO_DIM) {
+            const int64_t zero_shape[2] = { shape[0], 0 };
+            helper_write(file, zero_shape, 2*sizeof(int64_t));
         } else if (hft == HANDCRAFTED_TENSORS_NE_TOO_BIG){
             const int64_t big_dim = 4*int64_t(INT32_MAX);
             for (uint32_t j = 0; j < n_dims; ++j) {
@@ -446,6 +461,9 @@ static FILE * get_handcrafted_file(const unsigned int seed, const enum handcraft
         for (uint32_t i = 1; i < n_dims; ++i) {
             ne *= shape[i];
         }
+        if (hft == HANDCRAFTED_TENSORS_ZERO_DIM) {
+            ne = 0;
+        }
 
         offset += GGML_PAD(ggml_row_size(type, ne), (uint64_t) alignment);
     }
@@ -747,11 +765,13 @@ static std::pair<int, int> test_handcrafted_file(const unsigned int seed) {
         HANDCRAFTED_KV_BAD_TYPE,
         HANDCRAFTED_KV_DUPLICATE_KEY,
         HANDCRAFTED_KV_BAD_ALIGN,
+        HANDCRAFTED_KV_WRONG_TYPE_ALIGN,
         HANDCRAFTED_KV_SUCCESS,
 
         HANDCRAFTED_TENSORS_BAD_NAME_SIZE,
         HANDCRAFTED_TENSORS_BAD_N_DIMS,
         HANDCRAFTED_TENSORS_BAD_SHAPE,
+        HANDCRAFTED_TENSORS_ZERO_DIM,
         HANDCRAFTED_TENSORS_NE_TOO_BIG,
         HANDCRAFTED_TENSORS_NBYTES_TOO_BIG,
         HANDCRAFTED_TENSORS_BAD_TYPE,
@@ -840,7 +860,9 @@ static std::pair<int, int> test_handcrafted_file(const unsigned int seed) {
             ntest++;
         }
 
-        if (expect_context_not_null(hft) && hft >= offset_has_tensors) {
+        // HANDCRAFTED_TENSORS_ZERO_DIM deliberately mangles the tensor shapes to 0 elements,
+        // so only assert that it loads without crashing; skip the exact-geometry comparison.
+        if (expect_context_not_null(hft) && hft >= offset_has_tensors && hft != HANDCRAFTED_TENSORS_ZERO_DIM) {
             printf("%s:   - check_tensors: ", __func__);
             if (handcrafted_check_tensors(gguf_ctx, seed)) {
                 printf("\033[1;32mOK\033[0m\n");

tests/test-quantize-stats.cpp

@@ -301,7 +301,7 @@ int main(int argc, char ** argv) {
         return 1;
     }
 
-    llama_print_build_info();
+    llama_print_build_info(llama_version());
 
     // load the model
     fprintf(stderr, "Loading model\n");

tools/cvector-generator/cvector-generator.cpp

@@ -421,7 +421,7 @@ int main(int argc, char ** argv) {
     params.cb_eval_user_data = &cb_data;
     params.warmup = false;
 
-    llama_print_build_info();
+    llama_print_build_info(llama_version());
     llama_backend_init();
     llama_numa_init(params.numa);
 

tools/gguf-split/gguf-split.cpp

@@ -106,7 +106,7 @@ static void split_params_parse_ex(int argc, const char ** argv, split_params & p
             split_print_usage(argv[0]);
             exit(0);
         } else if (arg == "--version") {
-            fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit());
+            fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version(), llama_build_number(), llama_commit());
             fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
             exit(0);
         } else if (arg == "--dry-run") {

tools/mtmd/CMakeLists.txt

@@ -72,8 +72,8 @@ add_library(mtmd
             )
 
 set_target_properties(mtmd PROPERTIES
-    VERSION ${LLAMA_INSTALL_VERSION}
-    SOVERSION 0
+    VERSION ${LLAMA_VERSION_BASE}
+    SOVERSION ${LLAMA_VERSION_MAJOR}
     MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
 )
 

tools/quantize/quantize.cpp

@@ -611,7 +611,7 @@ int llama_quantize(int argc, char ** argv) {
         }
     }
 
-    llama_print_build_info();
+    llama_print_build_info(llama_version());
 
     if (params.dry_run) {
         fprintf(stderr, "%s: calculating quantization size for '%s' as %s", __func__, fname_inp.c_str(), ftype_str.c_str());

tools/server/server-common.cpp

@@ -13,6 +13,8 @@
 #include <sstream>
 #include <fstream>
 #include <limits>
+#include <cstring>
+#include <type_traits>
 
 json format_error_response(const std::string & message, const enum error_type type) {
     std::string type_str;
@@ -235,6 +237,102 @@ static inline raw_buffer base64_decode(const std::string & encoded_string) {
 // server_tokens implementation
 //
 
+namespace {
+
+constexpr uint32_t SERVER_TOKENS_STATE_VERSION = 1;
+
+uint32_t server_tokens_state_u32(size_t value) {
+    if (value > std::numeric_limits<uint32_t>::max()) {
+        throw std::runtime_error("Server tokens state is too large");
+    }
+    return value;
+}
+
+class server_tokens_state_writer {
+public:
+    template <typename T>
+    void write(T value) {
+        static_assert(std::is_trivially_copyable<T>::value, "T must be trivially copyable");
+        const auto * ptr = reinterpret_cast<const char *>(&value);
+        data.insert(data.end(), ptr, ptr + sizeof(value));
+    }
+
+    template <typename T>
+    void write(const std::vector<T> & values) {
+        static_assert(std::is_trivially_copyable<T>::value, "T must be trivially copyable");
+        write(server_tokens_state_u32(values.size()));
+        if (values.empty()) {
+            return;
+        }
+        const auto * ptr = reinterpret_cast<const char *>(values.data());
+        data.insert(data.end(), ptr, ptr + values.size() * sizeof(T));
+    }
+
+    void write_media_chunk(const mtmd_input_chunk * chunk) {
+        size_t chunk_size = 0;
+        if (mtmd_input_chunk_save(chunk, nullptr, 0, &chunk_size) != 0 || chunk_size == 0) {
+            throw std::runtime_error("Cannot serialize media chunk in server tokens");
+        }
+        std::vector<char> chunk_data(server_tokens_state_u32(chunk_size));
+        if (mtmd_input_chunk_save(chunk, chunk_data.data(), chunk_data.size(), nullptr) != 0) {
+            throw std::runtime_error("Cannot serialize media chunk in server tokens");
+        }
+        write(chunk_data);
+    }
+
+    std::vector<char> take() {
+        data.resize((data.size() + sizeof(llama_token) - 1) / sizeof(llama_token) * sizeof(llama_token), 0);
+        return std::move(data);
+    }
+
+private:
+    std::vector<char> data;
+};
+
+class server_tokens_state_reader {
+public:
+    server_tokens_state_reader(const char * data, size_t size) : data(data), size(size) {}
+
+    template <typename T>
+    T read() {
+        static_assert(std::is_trivially_copyable<T>::value, "T must be trivially copyable");
+        if (size - pos < sizeof(T)) {
+            throw std::runtime_error("Unexpected end of server tokens state");
+        }
+        T value;
+        std::memcpy(&value, data + pos, sizeof(value));
+        pos += sizeof(value);
+        return value;
+    }
+
+    template <typename T>
+    std::vector<T> read_vector() {
+        static_assert(std::is_trivially_copyable<T>::value, "T must be trivially copyable");
+        const uint32_t n_values = read<uint32_t>();
+        // reject before resizing, so that a small corrupted payload cannot request a huge allocation
+        if (n_values > remaining() / sizeof(T)) {
+            throw std::runtime_error("Unexpected end of server tokens state");
+        }
+        std::vector<T> values(n_values);
+        if (n_values > 0) {
+            std::memcpy(values.data(), data + pos, values.size() * sizeof(T));
+            pos += values.size() * sizeof(T);
+        }
+        return values;
+    }
+
+    size_t remaining() const {
+        return size - pos;
+    }
+
+private:
+    const char * data;
+    size_t size;
+    size_t pos = 0;
+};
+
+} // namespace
+
 server_tokens::server_tokens(mtmd::input_chunks & mtmd_chunks, bool has_mtmd) : has_mtmd(has_mtmd) {
     for (size_t i = 0; i < mtmd_chunks.size(); ++i) {
         push_back(mtmd_chunks[i]);
@@ -408,6 +506,73 @@ const llama_tokens & server_tokens::get_tokens() const {
     return tokens;
 }
 
+std::vector<char> server_tokens::serialize() const {
+    static_assert(sizeof(llama_token) == sizeof(uint32_t), "unexpected llama_token size");
+
+    server_tokens_state_writer writer;
+    writer.write((llama_token) LLAMA_TOKEN_NULL);
+    writer.write(SERVER_TOKENS_STATE_VERSION);
+    writer.write(tokens);
+
+    std::vector<uint32_t> media_keys;
+    media_keys.reserve(map_idx_to_media.size());
+    for (const auto & item : map_idx_to_media) {
+        media_keys.push_back(server_tokens_state_u32(item.first));
+    }
+    writer.write(media_keys);
+
+    for (const auto & item : map_idx_to_media) {
+        writer.write_media_chunk(item.second.get());
+    }
+
+    return writer.take();
+}
+
+server_tokens server_tokens::deserialize(const llama_tokens & packed, bool has_mtmd) {
+    static_assert(sizeof(llama_token) == sizeof(uint32_t), "unexpected llama_token size");
+
+    if (packed.empty() || packed[0] != LLAMA_TOKEN_NULL) {
+        // plain token list, as written by older versions
+        return server_tokens(packed, has_mtmd);
+    }
+
+    server_tokens_state_reader reader(reinterpret_cast<const char *>(packed.data()), packed.size() * sizeof(llama_token));
+    reader.read<llama_token>(); // format marker
+    if (reader.read<uint32_t>() != SERVER_TOKENS_STATE_VERSION) {
+        throw std::runtime_error("Unsupported server tokens state version");
+    }
+
+    const llama_tokens tokens = reader.read_vector<llama_token>();
+
+    // the media start indices, followed by the media chunks in the same order
+    const std::vector<uint32_t> media_keys = reader.read_vector<uint32_t>();
+    if (!media_keys.empty() && !has_mtmd) {
+        throw std::runtime_error("Cannot restore media tokens without an mmproj");
+    }
+
+    server_tokens result(tokens, has_mtmd);
+
+    for (const uint32_t key : media_keys) {
+        const size_t start_idx = key;
+        const std::vector<char> chunk_data = reader.read_vector<char>();
+        if (chunk_data.empty()) {
+            throw std::runtime_error("Cannot load media chunk from server tokens state");
+        }
+
+        mtmd::input_chunk_ptr chunk(mtmd_input_chunk_load(chunk_data.data(), chunk_data.size()));
+        if (!chunk) {
+            throw std::runtime_error("Cannot load media chunk from server tokens state");
+        }
+        result.map_idx_to_media[start_idx] = std::move(chunk);
+    }
+
+    if (reader.remaining() >= sizeof(llama_token)) {
+        throw std::runtime_error("Trailing data in server tokens state");
+    }
+
+    return result;
+}
+
 llama_tokens server_tokens::get_text_tokens() const {
     llama_tokens res;
     res.reserve(tokens.size());
@@ -530,14 +695,28 @@ bool server_tokens::validate(const struct llama_context * ctx) const {
     const llama_model * model = llama_get_model(ctx);
     const llama_vocab * vocab = llama_model_get_vocab(model);
     const int32_t n_vocab = llama_vocab_n_tokens(vocab);
+    size_t n_media = 0;
 
     for (size_t i = 0; i < tokens.size(); ++i) {
         const auto & t = tokens[i];
         if (t == LLAMA_TOKEN_NULL) {
             try {
                 const auto & chunk = find_chunk(i);
-                size_t n_tokens = mtmd_input_chunk_get_n_tokens(chunk.get());
-                i += n_tokens - 1; // will be +1 by the for loop
+                if (mtmd_input_chunk_get_type(chunk.get()) == MTMD_INPUT_CHUNK_TYPE_TEXT) {
+                    return false;
+                }
+                const size_t n_tokens = mtmd_input_chunk_get_n_tokens(chunk.get());
+                const llama_pos n_pos = mtmd_input_chunk_get_n_pos(chunk.get());
+                if (n_tokens == 0 || n_pos <= 0 || n_tokens > tokens.size() - i) {
+                    return false;
+                }
+                for (size_t j = i; j < i + n_tokens; ++j) {
+                    if (tokens[j] != LLAMA_TOKEN_NULL) {
+                        return false;
+                    }
+                }
+                ++n_media;
+                i += n_tokens - 1;
             } catch (const std::exception & e) {
                 return false;
             }
@@ -545,7 +724,7 @@ bool server_tokens::validate(const struct llama_context * ctx) const {
             return false;
         }
     }
-    return true;
+    return n_media == map_idx_to_media.size();
 }
 
 server_tokens server_tokens::clone() const {

tools/server/server-common.h

@@ -201,11 +201,14 @@ public:
     // for compatibility with context shift and prompt truncation
     void insert(const llama_tokens & inp_tokens);
 
-    // for compatibility with speculative decoding, ctx shift, slot save/load
+    // for compatibility with speculative decoding, ctx shift
     const llama_tokens & get_tokens() const;
 
     llama_tokens get_text_tokens() const;
 
+    std::vector<char> serialize() const;
+    static server_tokens deserialize(const llama_tokens & packed, bool has_mtmd);
+
     // for compatibility with speculative decoding
     void set_token(llama_pos pos, llama_token id);
 
@@ -213,9 +216,6 @@ public:
 
     bool empty() const { return tokens.empty(); }
 
-    // true if the sequence actually contains image/audio chunks.
-    bool has_media() const { return !map_idx_to_media.empty(); }
-
     void clear() {
         map_idx_to_media.clear();
         tokens.clear();
@@ -230,7 +230,7 @@ public:
     // split the tokens into message spans, skipping over media chunks
     common_chat_msg_spans find_message_spans(const common_chat_msg_delimiters & delims) const;
 
-    // make sure all text tokens are within the vocab range
+    // check text token IDs and the mapping between media chunks and token ranges
     bool validate(const struct llama_context * ctx) const;
 
     server_tokens clone() const;

tools/server/server-context.cpp

@@ -2072,18 +2072,6 @@ private:
         queue_results.send(std::move(res));
     }
 
-    // Gate slot save/restore/erase on slot content (does it hold media),
-    // not model capability: a multimodal model may hold a pure-text slot.
-    bool check_slot_no_media(const server_slot & slot, const int id_task) {
-        if (slot.prompt.tokens.has_media()) {
-            send_error(id_task,
-                "This operation is not supported while the slot holds image/audio tokens (a pure-text prefix is supported)",
-                ERROR_TYPE_NOT_SUPPORTED);
-            return false;
-        }
-        return true;
-    }
-
     void send_partial_response(server_slot & slot, const completion_token_output & tkn, bool is_progress, bool is_begin = false) {
         auto res = std::make_unique<server_task_result_cmpl_partial>();
 
@@ -2577,9 +2565,6 @@ private:
                         send_error(task, "Invalid slot ID", ERROR_TYPE_INVALID_REQUEST);
                         break;
                     }
-                    if (!check_slot_no_media(*slot, task.id)) {
-                        break;
-                    }
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);
@@ -2592,9 +2577,22 @@ private:
                     std::string filename = task.slot_action.filename;
                     std::string filepath = task.slot_action.filepath;
 
-                    const llama_tokens tokens = slot->prompt.tokens.get_text_tokens();
-                    const size_t token_count = tokens.size();
-                    const size_t nwrite = llama_state_seq_save_file(ctx_tgt, filepath.c_str(), slot->id, tokens.data(), token_count);
+                    std::vector<char> packed;
+                    try {
+                        packed = slot->prompt.tokens.serialize();
+                    } catch (const std::exception & err) {
+                        send_error(task, err.what(), ERROR_TYPE_NOT_SUPPORTED);
+                        break;
+                    }
+
+                    GGML_ASSERT(packed.size() % sizeof(llama_token) == 0);
+                    const size_t nwrite = llama_state_seq_save_file(
+                        ctx_tgt, filepath.c_str(), slot->id,
+                        reinterpret_cast<const llama_token *>(packed.data()), packed.size() / sizeof(llama_token));
+                    if (nwrite == 0) {
+                        send_error(task, "Unable to save slot", ERROR_TYPE_SERVER);
+                        break;
+                    }
 
                     const int64_t t_end = ggml_time_us();
                     const double t_save_ms = (t_end - t_start) / 1000.0;
@@ -2604,7 +2602,7 @@ private:
                     res->id_slot  = id_slot;
                     res->filename = filename;
                     res->is_save  = true;
-                    res->n_tokens = token_count;
+                    res->n_tokens = slot->prompt.tokens.size();
                     res->n_bytes  = nwrite;
                     res->t_ms     = t_save_ms;
                     queue_results.send(std::move(res));
@@ -2629,18 +2627,37 @@ private:
                     std::string filename = task.slot_action.filename;
                     std::string filepath = task.slot_action.filepath;
 
-                    llama_tokens tokens;
-                    tokens.resize(slot->n_ctx);
-                    size_t token_count = 0;
-                    size_t nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, tokens.data(), tokens.size(), &token_count);
-                    if (nread == 0) {
-                        slot->prompt.clear(); // KV may already been invalidated?
-                        send_error(task, "Unable to restore slot, no available space in KV cache or invalid slot save file", ERROR_TYPE_INVALID_REQUEST);
+                    size_t nread = 0;
+                    try {
+                        size_t n_packed = 0;
+                        llama_tokens packed;
+                        nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, nullptr, 0, &n_packed);
+                        if (nread != 0) {
+                            packed.resize(std::max<size_t>(1, n_packed));
+                            nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, packed.data(), packed.size(), &n_packed);
+                        }
+                        if (nread == 0) {
+                            throw std::runtime_error("No available space in KV cache or invalid slot save file");
+                        }
+                        packed.resize(n_packed);
+
+                        server_tokens restored = server_tokens::deserialize(packed, mctx != nullptr);
+
+                        if (restored.size() > (size_t) slot->n_ctx) {
+                            throw std::runtime_error("Restored prompt does not fit in the slot context");
+                        }
+
+                        if (!restored.validate(ctx_tgt)) {
+                            throw std::runtime_error("Invalid tokens in slot save file");
+                        }
+
+                        slot->prompt.clear();
+                        slot->prompt.tokens = std::move(restored);
+                    } catch (const std::exception & err) {
+                        slot->prompt_clear();
+                        send_error(task, std::string("Unable to restore slot: ") + err.what(), ERROR_TYPE_INVALID_REQUEST);
                         break;
                     }
-                    tokens.resize(token_count);
-                    slot->prompt.clear();
-                    slot->prompt.tokens.insert(tokens);
 
                     const int64_t t_end = ggml_time_us();
                     const double t_restore_ms = (t_end - t_start) / 1000.0;
@@ -2650,7 +2667,7 @@ private:
                     res->id_slot  = id_slot;
                     res->filename = filename;
                     res->is_save  = false;
-                    res->n_tokens = token_count;
+                    res->n_tokens = slot->prompt.tokens.size();
                     res->n_bytes  = nread;
                     res->t_ms     = t_restore_ms;
                     queue_results.send(std::move(res));
@@ -2663,10 +2680,6 @@ private:
                         send_error(task, "Invalid slot ID", ERROR_TYPE_INVALID_REQUEST);
                         break;
                     }
-                    // Gate on slot content, consistent with save/restore.
-                    if (!check_slot_no_media(*slot, task.id)) {
-                        break;
-                    }
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);

tools/server/server-tools.cpp

@@ -1,6 +1,7 @@
 #include "server-tools.h"
 
 #include "subproc.h"
+#include "base64.hpp"
 
 #include <filesystem>
 #include <fstream>
@@ -864,6 +865,7 @@ static bool path_glob_match(const std::string & pattern, const std::string & rel
 //
 
 static constexpr size_t SERVER_TOOL_READ_FILE_MAX_SIZE = 16 * 1024; // 16 KB
+static constexpr size_t SERVER_TOOL_READ_FILE_MAX_SIZE_BASE64 = 32 * 1024 * 1024; // 32 MB
 
 struct server_tool_read_file : server_tool {
     server_tool_read_file() {
@@ -899,6 +901,8 @@ struct server_tool_read_file : server_tool {
         int  start_line   = json_value(params, "start_line", 1);
         int  end_line     = json_value(params, "end_line",  -1); // -1 = no limit
         bool append_loc   = json_value(params, "append_loc", false);
+        // comes from the x-resp-type header, the model cannot ask for it
+        bool as_base64    = json_value(params, "resp_type", std::string()) == "base64";
 
         auto io = make_tools_io(params);
 
@@ -906,6 +910,23 @@ struct server_tool_read_file : server_tool {
         if (!io->file_size(path, file_size)) {
             return {{"error", "cannot stat file: " + path}};
         }
+
+        if (as_base64) {
+            if (file_size > SERVER_TOOL_READ_FILE_MAX_SIZE_BASE64) {
+                return {{"error", string_format(
+                    "file too large (%zu bytes, max %zu)",
+                    (size_t)file_size, SERVER_TOOL_READ_FILE_MAX_SIZE_BASE64)}};
+            }
+            std::string content;
+            if (!io->read_file(path, content)) {
+                return {{"error", "failed to open file: " + path}};
+            }
+            return {
+                {"base64",     base64::encode(content.data(), content.size())},
+                {"size_bytes", (size_t) content.size()},
+            };
+        }
+
         if (file_size > SERVER_TOOL_READ_FILE_MAX_SIZE && end_line == -1) {
             return {{"error", string_format(
                 "file too large (%zu bytes, max %zu). Use start_line/end_line to read a portion.",
@@ -2135,6 +2156,15 @@ void server_tools::setup(const std::vector<std::string> & enabled_tools,
                 params["runtime"] = runtime->spec();
             }
 
+            // x-resp-type header is only used by read_file for now
+            if (params.contains("resp_type")) {
+                params.erase("resp_type");
+            }
+            auto resp_type = get_header(req.headers, "x-resp-type");
+            if (!resp_type.empty()) {
+                params["resp_type"] = resp_type;
+            }
+
             server_tool & tool = find_tool(tools, tool_name, stream);
 
             if (stream) {

tools/server/tests/unit/test_slot_save.py

@@ -2,6 +2,10 @@ import pytest
 from utils import *
 import base64
 import requests
+import struct
+
+# sequence state file: magic(4) version(4) payload_size(4), then payload_size llama_token words
+STATE_FILE_HEADER_SIZE = 12
 
 server = ServerPreset.tinyllama2()
 
@@ -72,6 +76,60 @@ def test_slot_save_restore():
     assert res.body["timings"]["prompt_n"] == 1
 
 
+def test_slot_restore_legacy_token_list():
+    global server
+    server.start()
+
+    res = server.make_request("POST", "/completion", data={
+        "prompt": "What is the capital of France?",
+        "id_slot": 1,
+        "cache_prompt": True,
+    })
+    assert res.status_code == 200
+
+    res = server.make_request("POST", "/slots/1?action=save", data={
+        "filename": "slot_legacy.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_saved"] == 84
+
+    # rewrite the token payload into a plain token list, as written by servers that predate the packed server_tokens format
+    path = os.path.join("tmp", "slot_legacy.bin")
+    with open(path, "rb") as f:
+        data = bytearray(f.read())
+
+    # the payload written by this server starts with a packed header: LLAMA_TOKEN_NULL(4) version(4) n_tokens(4)
+    packed_header_size = 12
+
+    payload_size = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE - 4)[0]
+    payload_end = STATE_FILE_HEADER_SIZE + payload_size * 4
+    n_tokens = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE + 8)[0]
+    assert n_tokens == 84
+
+    tokens_start = STATE_FILE_HEADER_SIZE + packed_header_size
+    data = data[:STATE_FILE_HEADER_SIZE] + data[tokens_start:tokens_start + n_tokens * 4] + data[payload_end:]
+    struct.pack_into("=I", data, STATE_FILE_HEADER_SIZE - 4, n_tokens)
+
+    with open(path, "wb") as f:
+        f.write(data)
+
+    # the plain token list must restore, and the restored KV must be reusable
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "slot_legacy.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == 84
+
+    res = server.make_request("POST", "/completion", data={
+        "prompt": "What is the capital of Germany?",
+        "id_slot": 0,
+        "cache_prompt": True,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["prompt_n"] == 6  # only the different part is processed
+
+
+
 def test_slot_erase():
     global server
     server.start()
@@ -103,14 +161,12 @@ def test_slot_erase():
 #
 # Multimodal server (mmproj loaded) slot save/restore.
 #
-# Regression coverage for issue #21133: slot save/restore/erase must be gated on
-# the slot's CONTENT (does it actually hold image/audio tokens) rather than the
-# model's CAPABILITY (is an mmproj loaded). A pure-text slot on a multimodal
-# server must save/restore/erase normally; a slot that actually holds an image
-# must be rejected with ERROR_TYPE_NOT_SUPPORTED (HTTP 501).
+# A pure-text slot on a multimodal server and a slot containing images must both support save/restore.
+# Erase remains gated on the slot's content.
 #
 
 IMG_URL_CAT = "https://huggingface.co/ggml-org/tinygemma3-GGUF/resolve/main/test/91_cat.png"
+IMG_URL_TRUCK = "https://huggingface.co/ggml-org/tinygemma3-GGUF/resolve/main/test/11_truck.png"
 
 
 def _get_img_base64(url: str) -> str:
@@ -121,8 +177,7 @@ def _get_img_base64(url: str) -> str:
 
 @pytest.fixture
 def mmproj_server():
-    # tinygemma3 is a small multimodal model: the mmproj is provided by the HF
-    # registry API and auto-downloaded on first run.
+    # tinygemma3 is a small multimodal model: the mmproj is provided by the HF registry API and auto-downloaded on first run.
     os.environ['LLAMA_MEDIA_MARKER'] = '<__media__>'
     mm_server = ServerPreset.tinygemma3()
     mm_server.slot_save_path = "./tmp"
@@ -159,10 +214,7 @@ def test_slot_save_restore_text_only_on_multimodal(mmproj_server):
     assert res.status_code == 200
     assert res.body["n_restored"] == n_saved
 
-    # The restored slot is usable for a follow-up completion. We do NOT assert
-    # prefix reuse here: tinygemma3 is a SWA model, which forces full prompt
-    # re-processing after a restore (a model property, not the save/restore gate
-    # under test).
+    # Prefix reuse is not checked with the default SWA cache.
     res = server.make_request("POST", "/completion", data={
         "prompt": "The quick brown fox jumps over the lazy dog.",
         "id_slot": 0,
@@ -171,54 +223,307 @@ def test_slot_save_restore_text_only_on_multimodal(mmproj_server):
     assert res.status_code == 200
 
 
-def test_slot_save_rejected_when_slot_holds_image(mmproj_server):
+def test_slot_save_restore_with_image(mmproj_server):
     server = mmproj_server
+    # Use the full SWA cache so the restored image prefix can be reused.
+    server.swa_full = True
     server.start()
 
-    # Process a prompt that actually contains an image on slot 1.
+    prompt_cat = {
+        "prompt_string": "What is this: <__media__>\n",
+        "multimodal_data": [_get_img_base64(IMG_URL_CAT)],
+    }
     res = server.make_request("POST", "/completions", data={
         "temperature": 0.0,
         "top_k": 1,
         "id_slot": 1,
         "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    content_cat = res.body["content"]
+    prompt_n_full = res.body["timings"]["prompt_n"]
+    assert res.body["timings"]["cache_n"] == 0
+    assert prompt_n_full > 32  # text plus image tokens are all processed
+
+    res = server.make_request("POST", "/slots/1?action=save", data={
+        "filename": "mm_slot_image.bin",
+    })
+    assert res.status_code == 200
+    n_saved = res.body["n_saved"]
+    n_written = res.body["n_written"]
+    assert n_saved > 0
+    assert n_written > 0
+
+    res = server.make_request("POST", "/slots/1?action=erase")
+    assert res.status_code == 200
+
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_image.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == n_saved
+    assert res.body["n_read"] == n_written
+
+    # a different image must not reuse the restored image tokens; only the text prefix before the image is common
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
         "prompt": {
             "prompt_string": "What is this: <__media__>\n",
-            "multimodal_data": [ _get_img_base64(IMG_URL_CAT) ],
+            "multimodal_data": [_get_img_base64(IMG_URL_TRUCK)],
+        },
+    })
+    assert res.status_code == 200
+    cache_n = res.body["timings"]["cache_n"]
+    assert cache_n < 16
+    assert res.body["timings"]["prompt_n"] == prompt_n_full - cache_n
+
+    # restore again and resend the same image: the image tokens must be reused and greedy sampling must reproduce the original content
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_image.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == n_saved
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+    assert res.body["timings"]["prompt_n"] == 1
+    assert res.body["content"] == content_cat
+
+
+def test_slot_save_restore_with_two_images(mmproj_server):
+    server = mmproj_server
+    server.swa_full = True
+    server.n_ctx = 2048  # two images need more than the default 512 per slot
+    server.start()
+
+    prompt = {
+        "prompt_string": "A: <__media__> B: <__media__>\n",
+        "multimodal_data": [_get_img_base64(IMG_URL_CAT), _get_img_base64(IMG_URL_TRUCK)],
+    }
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 1,
+        "cache_prompt": True,
+        "prompt": prompt,
+    })
+    assert res.status_code == 200
+    content = res.body["content"]
+    prompt_n_full = res.body["timings"]["prompt_n"]
+    assert prompt_n_full > 64
+
+    res = server.make_request("POST", "/slots/1?action=save", data={
+        "filename": "mm_slot_two_images.bin",
+    })
+    assert res.status_code == 200
+    n_saved = res.body["n_saved"]
+
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_two_images.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == n_saved
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+    assert res.body["timings"]["prompt_n"] == 1
+    assert res.body["content"] == content
+
+
+def test_slot_save_restore_with_image_across_restart(mmproj_server):
+    server = mmproj_server
+    server.swa_full = True
+    server.start()
+
+    prompt_cat = {
+        "prompt_string": "What is this: <__media__>\n",
+        "multimodal_data": [_get_img_base64(IMG_URL_CAT)],
+    }
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    content = res.body["content"]
+    prompt_n_full = res.body["timings"]["prompt_n"]
+
+    res = server.make_request("POST", "/slots/0?action=save", data={
+        "filename": "mm_slot_restart.bin",
+    })
+    assert res.status_code == 200
+    n_saved = res.body["n_saved"]
+
+    # restart the server with the same model and mmproj: the saved file must restore in the new process and the image KV must be reused
+    server.stop()
+    server.start()
+
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_restart.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == n_saved
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+    assert res.body["timings"]["prompt_n"] == 1
+    assert res.body["content"] == content
+
+
+def test_slot_save_restore_image_payload_larger_than_context(mmproj_server):
+    server = mmproj_server
+    server.swa_full = True
+    server.start()
+
+    # the slot context, as the server computed it (n_ctx split across the slots)
+    res = server.make_request("GET", "/props")
+    assert res.status_code == 200
+    n_ctx_slot = res.body["default_generation_settings"]["n_ctx"]
+
+    # a filler token, used to grow the prompt up to the slot context
+    res = server.make_request("POST", "/tokenize", data={"content": " hello" * 8})
+    assert res.status_code == 200
+    assert len(res.body["tokens"]) == 8
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": {
+            "prompt_string": "What is this: <__media__>\n",
+            "multimodal_data": [_get_img_base64(IMG_URL_CAT)],
         },
     })
     assert res.status_code == 200
 
-    # Saving a slot that holds image tokens must be rejected (HTTP 501,
-    # not_supported_error).
-    res = server.make_request("POST", "/slots/1?action=save", data={
-        "filename": "mm_slot_image.bin",
+    prompt_cat = {
+        "prompt_string": "What is this: <__media__>\n" + " hello" * (n_ctx_slot - res.body["timings"]["prompt_n"] - 8),
+        "multimodal_data": [_get_img_base64(IMG_URL_CAT)],
+    }
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
     })
-    assert res.status_code != 200
-    assert res.body["error"]["type"] == "not_supported_error"
+    assert res.status_code == 200
+    prompt_n_full = res.body["timings"]["cache_n"] + res.body["timings"]["prompt_n"]
+
+    res = server.make_request("POST", "/slots/0?action=save", data={
+        "filename": "mm_slot_large_payload.bin",
+    })
+    assert res.status_code == 200
+
+    path = os.path.join("tmp", "mm_slot_large_payload.bin")
+    with open(path, "rb") as f:
+        data = bytearray(f.read())
+    payload_size = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE - 4)[0]
+    assert payload_size > n_ctx_slot  # the scenario under test: the payload does not fit in n_ctx
+
+    # drop the image from the slot, then restore it from the file
+    res = server.make_request("POST", "/completion", data={
+        "prompt": "The quick brown fox",
+        "id_slot": 0,
+        "cache_prompt": True,
+    })
+    assert res.status_code == 200
+
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_large_payload.bin",
+    })
+    assert res.status_code == 200
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+    assert res.body["timings"]["prompt_n"] == 1
 
 
-def test_slot_erase_text_only_on_multimodal(mmproj_server):
+def test_slot_restore_media_file_without_mmproj(mmproj_server):
     server = mmproj_server
     server.start()
 
-    res = server.make_request("POST", "/completion", data={
-        "prompt": "The quick brown fox jumps over the lazy dog.",
-        "id_slot": 1,
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
         "cache_prompt": True,
+        "prompt": {
+            "prompt_string": "What is this: <__media__>\n",
+            "multimodal_data": [_get_img_base64(IMG_URL_CAT)],
+        },
     })
     assert res.status_code == 200
-    prompt_n = res.body["timings"]["prompt_n"]
-    assert prompt_n > 0  # all tokens are processed
 
-    # Erasing a pure-text slot must succeed even though an mmproj is loaded.
-    res = server.make_request("POST", "/slots/1?action=erase")
-    assert res.status_code == 200
-
-    # Re-running the same prompt should process all tokens again.
-    res = server.make_request("POST", "/completion", data={
-        "prompt": "The quick brown fox jumps over the lazy dog.",
-        "id_slot": 1,
-        "cache_prompt": True,
+    res = server.make_request("POST", "/slots/0?action=save", data={
+        "filename": "mm_slot_no_mmproj.bin",
     })
     assert res.status_code == 200
-    assert res.body["timings"]["prompt_n"] == prompt_n  # all tokens are processed again
+
+    # restart the same model without the mmproj: restoring the media file must fail gracefully and leave the slot usable
+    server.stop()
+    server.no_mmproj = True
+    server.start()
+
+    res = server.make_request("POST", "/slots/0?action=restore", data={
+        "filename": "mm_slot_no_mmproj.bin",
+    })
+    assert res.status_code == 400
+    assert "Cannot restore media tokens without an mmproj" in res.body["error"]["message"]
+
+    # A failed restore must leave the slot empty and usable.
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 1,
+        "cache_prompt": True,
+        "prompt": "The quick brown fox",
+    })
+    assert res.status_code == 200
+    content = res.body["content"]
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": "The quick brown fox",
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == 0
+    assert res.body["content"] == content

tools/server/tests/unit/test_speculative.py

@@ -27,8 +27,8 @@ def test_with_and_without_draft():
     global server
     request = {
         "prompt": "I believe the meaning of life is",
-        "temperature": 0.8,
-        "top_k": 40,
+        "temperature": 0.2,
+        "top_k": 5,
         "seed": 4242,
         "n_predict": 16,
         "return_tokens": True,
@@ -36,7 +36,6 @@ def test_with_and_without_draft():
 
     server.model_draft = None  # disable draft model
     server.spec_type = None
-    server.backend_sampling = True
     server.start()
     res = server.make_request("POST", "/completion", data=request)
     assert res.status_code == 200
@@ -45,7 +44,6 @@ def test_with_and_without_draft():
 
     # create new server with draft model
     create_server()
-    server.backend_sampling = True
     server.start()
     res = server.make_request("POST", "/completion", data=request)
     assert res.status_code == 200

tools/server/tests/utils.py

@@ -86,6 +86,7 @@ class ServerProcess:
     server_reranking: bool | None = False
     server_metrics: bool | None = False
     kv_unified: bool | None = False
+    swa_full: bool | None = False
     server_slots: bool | None = False
     pooling: str | None = None
     api_key: str | None = None
@@ -106,6 +107,7 @@ class ServerProcess:
     chat_template_file: str | None = None
     server_path: str | None = None
     mmproj_url: str | None = None
+    no_mmproj: bool | None = None
     media_path: str | None = None
     sleep_idle_seconds: int | None = None
     cache_ram: int | None = None
@@ -198,6 +200,8 @@ class ServerProcess:
             server_args.append("--metrics")
         if self.kv_unified:
             server_args.append("--kv-unified")
+        if self.swa_full:
+            server_args.append("--swa-full")
         if self.server_slots:
             server_args.append("--slots")
         else:
@@ -259,6 +263,8 @@ class ServerProcess:
             server_args.extend(["--chat-template-file", self.chat_template_file])
         if self.mmproj_url:
             server_args.extend(["--mmproj-url", self.mmproj_url])
+        if self.no_mmproj:
+            server_args.append("--no-mmproj")
         if self.media_path:
             server_args.extend(["--media-path", self.media_path])
         if self.sleep_idle_seconds is not None:

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte

@@ -7,6 +7,7 @@
 	import ChatMessageToolCallBlockGetInfo from './ChatMessageToolCallBlockGetInfo.svelte';
 	import ChatMessageToolCallBlockGrepSearch from './ChatMessageToolCallBlockGrepSearch.svelte';
 	import ChatMessageToolCallBlockReadFile from './ChatMessageToolCallBlockReadFile.svelte';
+	import ChatMessageToolCallBlockReadMedia from './ChatMessageToolCallBlockReadMedia.svelte';
 	import ChatMessageToolCallBlockRunJavascript from './ChatMessageToolCallBlockRunJavascript.svelte';
 	import ChatMessageToolCallBlockSearchResults from './ChatMessageToolCallBlockSearchResults.svelte';
 	import ChatMessageToolCallBlockWriteFile from './ChatMessageToolCallBlockWriteFile.svelte';
@@ -45,6 +46,8 @@
 	<ChatMessageToolCallBlockGetInfo {section} {isStreaming} />
 {:else if section.toolName === BuiltInTool.READ_FILE}
 	<ChatMessageToolCallBlockReadFile {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.READ_MEDIA}
+	<ChatMessageToolCallBlockReadMedia {section} {open} {isStreaming} {onToggle} />
 {:else if section.toolName === BuiltInTool.EDIT_FILE}
 	<ChatMessageToolCallBlockEditFile {section} {open} {isStreaming} {onToggle} />
 {:else if section.toolName === BuiltInTool.WRITE_FILE}

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte

@@ -8,14 +8,16 @@
 	import { MarkdownContent, SyntaxHighlightedCode } from '$lib/components/app';
 	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
 	import { getBuiltinToolUi } from '$lib/constants/built-in-tools';
-	import { FileTypeText, ToolResultKind } from '$lib/enums';
+	import { AttachmentType, FileTypeText, MimeTypeAudio, ToolResultKind } from '$lib/enums';
 	import type { DatabaseMessageExtra } from '$lib/types';
 	import {
 		type AgenticSection,
 		classifyToolResult,
 		formatJsonPretty,
-		parseToolResultWithImages
+		parseToolResultWithMedia,
+		type ToolResultLine
 	} from '$lib/utils';
+	import { createBase64DataUrl } from '$lib/utils/data-url';
 
 	interface Props {
 		section: AgenticSection;
@@ -29,8 +31,8 @@
 
 	const title = $derived(getBuiltinToolUi(section.toolName)?.label ?? section.toolName ?? '');
 	const outputKind = $derived(classifyToolResult(section.toolResult));
-	const parsedLines = $derived(
-		section.toolResult ? parseToolResultWithImages(section.toolResult, attachments) : []
+	const parsedLines: ToolResultLine[] = $derived(
+		section.toolResult ? parseToolResultWithMedia(section.toolResult, attachments) : []
 	);
 </script>
 
@@ -103,13 +105,26 @@
 							<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">
 								{line.text}
 							</div>
-							{#if line.image}
-								<img
-									src={line.image.base64Url}
-									alt={line.image.name}
-									class="mt-2 mb-2 h-auto max-w-full rounded-lg"
-									loading="lazy"
-								/>
+							{#if line.media}
+								{#if line.media.type === AttachmentType.AUDIO}
+									{@const audioMimeType = line.media.mimeType ?? MimeTypeAudio.MP3_MPEG}
+									<div class="mt-2 mb-2">
+										<audio controls class="w-full rounded-lg">
+											<source
+												src={createBase64DataUrl(audioMimeType, line.media.base64Data)}
+												type={audioMimeType}
+											/>
+											Your browser does not support the audio element.
+										</audio>
+									</div>
+								{:else}
+									<img
+										src={line.media.base64Url}
+										alt={line.media.name}
+										class="mt-2 mb-2 h-auto max-w-full rounded-lg"
+										loading="lazy"
+									/>
+								{/if}
 							{/if}
 						{/each}
 					</div>

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte

@@ -23,7 +23,7 @@
 		isExitCodeSummaryLine,
 		parseExecShellCommandError,
 		parseExecShellCommandExitStatus,
-		parseToolResultWithImages,
+		parseToolResultWithMedia,
 		type ToolResultLine
 	} from '$lib/utils';
 
@@ -53,7 +53,7 @@
 	);
 
 	const parsedLines: ToolResultLine[] = $derived(
-		section.toolResult ? parseToolResultWithImages(section.toolResult, attachments) : []
+		section.toolResult ? parseToolResultWithMedia(section.toolResult, attachments) : []
 	);
 
 	// Drop the trailing "[exit code: N]" line - rendered as a colored
@@ -223,10 +223,10 @@
 			>
 				{#each outputLines as line, i (i)}
 					<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">{line.text}</div>
-					{#if line.image}
+					{#if line.media}
 						<img
-							src={line.image.base64Url}
-							alt={line.image.name}
+							src={line.media.base64Url}
+							alt={line.media.name}
 							class="mt-2 mb-2 h-auto max-w-full rounded-lg"
 							loading="lazy"
 						/>

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadMedia.svelte

@@ -0,0 +1,99 @@
+<script lang="ts">
+	import { parseReadMediaMeta } from './parsers/read-media';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+	import { ATTACHMENT_SAVED_REGEX } from '$lib/constants/agentic';
+	import { AttachmentType, MimeTypeAudio } from '$lib/enums';
+	import type { DatabaseMessageExtraAudioFile, DatabaseMessageExtraImageFile } from '$lib/types';
+	import { type AgenticSection } from '$lib/utils';
+	import { createBase64DataUrl } from '$lib/utils/data-url';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { isStreaming, onToggle, open, section }: Props = $props();
+
+	const readMediaMeta = $derived(parseReadMediaMeta(section));
+
+	// extractBase64Attachments swapped the data URI line for [Attachment saved: name]
+	// and moved the bytes to the message extras, so the name is the only link back
+	const mediaAttachment = $derived.by(() => {
+		const extras = section.toolResultExtras;
+
+		if (!extras || extras.length === 0) return null;
+
+		const match = section.toolResult?.match(ATTACHMENT_SAVED_REGEX);
+
+		if (!match) return null;
+
+		const attachmentName = match[1];
+
+		return (
+			extras.find(
+				(e): e is DatabaseMessageExtraImageFile | DatabaseMessageExtraAudioFile =>
+					(e.type === AttachmentType.IMAGE || e.type === AttachmentType.AUDIO) &&
+					e.name === attachmentName
+			) ?? null
+		);
+	});
+
+	const audioMimeType = $derived(readMediaMeta?.mimeType ?? MimeTypeAudio.MP3_MPEG);
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={readMediaMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		<span class="text-muted-foreground">Read media </span>
+		<span class="font-mono">{readMediaMeta?.fileName}</span>
+	{/snippet}
+
+	{#snippet children(_meta, _ctx)}
+		{#if section.toolResult}
+			{#if !mediaAttachment}
+				<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+					Media attachment not found in message extras
+				</div>
+			{:else if mediaAttachment.type === AttachmentType.AUDIO}
+				<div class="mt-2">
+					<audio controls class="w-full rounded-lg">
+						<source
+							src={createBase64DataUrl(audioMimeType, mediaAttachment.base64Data)}
+							type={audioMimeType}
+						/>
+						Your browser does not support the audio element.
+					</audio>
+				</div>
+			{:else}
+				<div class="mt-2">
+					<img
+						src={mediaAttachment.base64Url}
+						alt={readMediaMeta?.fileName ?? 'media'}
+						class="max-h-[60vh] max-w-full rounded-lg object-contain shadow-lg"
+						loading="lazy"
+					/>
+				</div>
+			{/if}
+
+			{#if readMediaMeta?.sizeBytes || readMediaMeta?.mimeType}
+				<div class="mt-2 flex gap-4 text-xs text-muted-foreground">
+					{#if readMediaMeta?.sizeBytes}
+						<span>Size: {readMediaMeta.sizeBytes} bytes</span>
+					{/if}
+					{#if readMediaMeta?.mimeType}
+						<span>MIME: {readMediaMeta.mimeType}</span>
+					{/if}
+				</div>
+			{/if}
+
+			{#if readMediaMeta?.path}
+				<div class="mt-1 font-mono text-xs text-muted-foreground/60">{readMediaMeta.path}</div>
+			{/if}
+		{:else}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+				Waiting for media data...
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-media.ts

@@ -0,0 +1,56 @@
+import { FILE_PATH_SEPARATOR_REGEX, NEWLINE } from '$lib/constants/code';
+import {
+	PREFIX_FILE,
+	PREFIX_MIME,
+	PREFIX_SIZE,
+	READ_MEDIA_SIZE_REGEX
+} from '$lib/constants/read-media';
+import type { AgenticSection } from '$lib/utils';
+
+export interface ReadMediaMeta {
+	fileName: string;
+	path: string;
+	sizeBytes?: number;
+	mimeType?: string;
+}
+
+/**
+ * Parse read_media tool result to extract metadata.
+ * Expected format (after extractBase64Attachments processing):
+ *   File: /path/to/file.png
+ *   Size: 12345 bytes
+ *   MIME: image/png
+ *   [Attachment saved: mcp-attachment-xxx.png]
+ *
+ * The data URI line is replaced by the attachment marker by
+ * agenticStore.extractBase64Attachments before storage.
+ */
+export function parseReadMediaMeta(section: AgenticSection): ReadMediaMeta | null {
+	if (!section.toolResult) return null;
+
+	const lines = section.toolResult.split(NEWLINE);
+
+	let fileName = '';
+	let path = '';
+	let sizeBytes: number | undefined;
+	let mimeType: string | undefined;
+
+	for (const line of lines) {
+		const trimmed = line.trim();
+
+		if (trimmed.startsWith(PREFIX_FILE)) {
+			path = trimmed.slice(PREFIX_FILE.length).trim();
+			fileName = path.split(FILE_PATH_SEPARATOR_REGEX).pop() ?? path;
+		} else if (trimmed.startsWith(PREFIX_SIZE)) {
+			const match = trimmed.match(READ_MEDIA_SIZE_REGEX);
+
+			if (match) sizeBytes = Number(match[1]);
+		} else if (trimmed.startsWith(PREFIX_MIME)) {
+			mimeType = trimmed.slice(PREFIX_MIME.length).trim();
+		}
+	}
+
+	if (!path) return null;
+
+	return { fileName, mimeType, path, sizeBytes };
+}

tools/ui/src/lib/constants/built-in-tools.ts

@@ -10,6 +10,7 @@
 import {
 	Braces,
 	Clock,
+	Eye,
 	FilePen,
 	FilePlus,
 	FileSearch,
@@ -47,6 +48,7 @@ export const BUILTIN_TOOL_UI: Readonly<Record<BuiltInTool, BuiltinToolUiEntry>>
 		source: ToolSource.BUILTIN
 	},
 	[BuiltInTool.READ_FILE]: { icon: FileText, label: 'Read file', source: ToolSource.BUILTIN },
+	[BuiltInTool.READ_MEDIA]: { icon: Eye, label: 'Read media', source: ToolSource.FRONTEND },
 	[BuiltInTool.RUN_JAVASCRIPT]: {
 		icon: Braces,
 		label: 'Run JavaScript',

tools/ui/src/lib/constants/code.ts

@@ -18,6 +18,9 @@ export const TRIM_TRAILING_PADDING_REGEX = /(?:\n[ \t]*)+$/;
 // `C:\foo\bar.txt`. Used wherever a parameter accepts a user-supplied path.
 export const FILE_PATH_SEPARATOR_REGEX = /[\\/]/;
 
+// Separates a file name from its extension, e.g. the '.' in `cover.png`.
+export const FILE_EXTENSION_SEPARATOR = '.';
+
 // Matches the `text:` prefix that file-type identifiers use to denote a
 // plain-text language (e.g. `text:typescript`). Used by tool-call renderers
 // to recover the underlying highlight.js language.

tools/ui/src/lib/constants/index.ts

@@ -49,6 +49,7 @@ export * from './sse';
 export * from './precision';
 export * from './processing-info';
 export * from './pwa';
+export * from './read-media';
 export * from './routes';
 export * from './sandbox';
 export * from './settings-keys';

tools/ui/src/lib/constants/mcp-resource.ts

@@ -1,4 +1,4 @@
-import { MimeTypeImage } from '$lib/enums';
+import { MimeTypeAudio, MimeTypeImage } from '$lib/enums';
 
 // File extension patterns for resource type detection
 export const IMAGE_FILE_EXTENSION_REGEX = /\.(png|jpg|jpeg|gif|svg|webp)$/i;
@@ -27,6 +27,9 @@ export const MCP_RESOURCE_ATTACHMENT_ID_PREFIX = 'res';
 // Default file extension for unknown image types
 export const DEFAULT_IMAGE_EXTENSION = 'img';
 
+// Default file extension for unknown audio types
+export const DEFAULT_AUDIO_EXTENSION = 'mp3';
+
 // Default filename for resource content downloads
 export const DEFAULT_RESOURCE_FILENAME = 'resource.txt';
 
@@ -53,3 +56,18 @@ export const IMAGE_MIME_TO_EXTENSION: Record<string, string> = {
 	[MimeTypeImage.PNG]: 'png',
 	[MimeTypeImage.WEBP]: 'webp'
 } as const;
+
+/**
+ * Mapping from audio MIME types to file extensions.
+ * Used for generating attachment filenames from MIME types.
+ */
+export const AUDIO_MIME_TO_EXTENSION: Record<string, string> = {
+	[MimeTypeAudio.MP3]: 'mp3',
+	[MimeTypeAudio.MP3_MPEG]: 'mp3',
+	[MimeTypeAudio.VND_WAVE]: 'wav',
+	[MimeTypeAudio.WAV]: 'wav',
+	[MimeTypeAudio.WAVE]: 'wav',
+	[MimeTypeAudio.X_PN_WAV]: 'wav',
+	[MimeTypeAudio.X_WAV]: 'wav',
+	[MimeTypeAudio.X_WAVE]: 'wav'
+} as const;

tools/ui/src/lib/constants/read-media.ts

@@ -0,0 +1,66 @@
+import {
+	BuiltInTool,
+	JsonSchemaType,
+	MimeTypeAudio,
+	MimeTypeImage,
+	ToolCallType
+} from '$lib/enums';
+import type { OpenAIToolDefinition } from '$lib/types';
+
+export const READ_MEDIA_TOOL_NAME = BuiltInTool.READ_MEDIA;
+
+// header lines of the tool result, parsed back by the read_media renderer
+export const PREFIX_FILE = 'File: ';
+export const PREFIX_SIZE = 'Size: ';
+export const PREFIX_MIME = 'MIME: ';
+
+/** Byte count of the `Size: ` header line, e.g. `Size: 12345 bytes` -> capture group 1 is `12345`. */
+export const READ_MEDIA_SIZE_REGEX = new RegExp(`^${PREFIX_SIZE}\\s*(\\d+)\\s*bytes`);
+
+/** Image extensions the tool accepts. The server decodes images with stb_image, which has no webp or tiff. */
+export const READ_MEDIA_IMAGE_MIME: Record<string, string> = {
+	gif: MimeTypeImage.GIF,
+	jpeg: MimeTypeImage.JPEG,
+	jpg: MimeTypeImage.JPEG,
+	png: MimeTypeImage.PNG
+} as const;
+
+/** Audio extensions the tool accepts. The `input_audio` API only takes wav and mp3. */
+export const READ_MEDIA_AUDIO_MIME: Record<string, string> = {
+	mp3: MimeTypeAudio.MP3_MPEG,
+	wav: MimeTypeAudio.WAV
+} as const;
+
+/**
+ * Build the read_media tool definition for the modalities the active model has.
+ * At least one of the two flags must be true, otherwise the tool is not offered
+ * at all - a model that cannot see or hear has nothing to do with the bytes.
+ */
+export function buildReadMediaToolDefinition(
+	supportsVision: boolean,
+	supportsAudio: boolean
+): OpenAIToolDefinition {
+	const kinds: string[] = [];
+
+	if (supportsVision) kinds.push(`images (${Object.keys(READ_MEDIA_IMAGE_MIME).join(', ')})`);
+
+	if (supportsAudio) kinds.push(`audio (${Object.keys(READ_MEDIA_AUDIO_MIME).join(', ')})`);
+
+	return {
+		function: {
+			description: `Read a media file and attach it to the conversation so it can be perceived directly. Supports ${kinds.join(' and ')}.`,
+			name: READ_MEDIA_TOOL_NAME,
+			parameters: {
+				properties: {
+					path: {
+						description: 'Path to the media file',
+						type: JsonSchemaType.STRING
+					}
+				},
+				required: ['path'],
+				type: JsonSchemaType.OBJECT
+			}
+		},
+		type: ToolCallType.FUNCTION
+	};
+}

tools/ui/src/lib/constants/tools.ts

@@ -3,6 +3,12 @@ import { ToolSource } from '$lib/enums/tools.enums';
 /** HTTP header carrying the working directory a tool call runs in. The server resolves relative paths against it; the model cannot override it. */
 export const X_TOOL_CWD_HEADER = 'x-tool-cwd';
 
+/** HTTP header asking the server to encode a tool's output differently, e.g. read_file returning base64. Not a tool parameter, so it stays out of the definition the model sees. */
+export const X_RESP_TYPE_HEADER = 'x-resp-type';
+
+/** `X_RESP_TYPE_HEADER` value that makes read_file return the raw bytes as base64 instead of text. */
+export const RESP_TYPE_BASE64 = 'base64';
+
 export const TOOL_GROUP_LABELS = {
 	[ToolSource.BUILTIN]: 'Built-in',
 	[ToolSource.CUSTOM]: 'JSON Schema',

tools/ui/src/lib/enums/files.enums.ts

@@ -163,6 +163,7 @@ export enum FileExtensionText {
 // MIME type prefixes and includes for content detection
 export enum MimeTypePrefix {
 	IMAGE = 'image/',
+	AUDIO = 'audio/',
 	TEXT = 'text'
 }
 

tools/ui/src/lib/enums/tools.enums.ts

@@ -37,6 +37,7 @@ export enum GlobSearchType {
  */
 export enum BuiltInTool {
 	READ_FILE = 'read_file',
+	READ_MEDIA = 'read_media',
 	EDIT_FILE = 'edit_file',
 	WRITE_FILE = 'write_file',
 	GET_DATETIME = 'get_datetime',

tools/ui/src/lib/services/chat.service.ts

@@ -1,4 +1,5 @@
 import { settingsStore } from '../stores/settings.svelte';
+import { getAudioInputFormat } from '../utils/audio-format';
 import { capImageDataURLSize } from '../utils/cap-img-size';
 import {
 	API_CHAT,
@@ -20,18 +21,12 @@ import {
 import {
 	AttachmentType,
 	ContentPartType,
-	FileTypeAudio,
 	MessageRole,
-	MimeTypeAudio,
 	ReasoningFormat,
 	StreamConnectionState
 } from '$lib/enums';
 import { modelsStore } from '$lib/stores/models.svelte';
-import type {
-	AudioInputFormat,
-	DatabaseMessageExtraMcpPrompt,
-	DatabaseMessageExtraMcpResource
-} from '$lib/types';
+import type { DatabaseMessageExtraMcpPrompt, DatabaseMessageExtraMcpResource } from '$lib/types';
 import type {
 	ApiChatCompletionToolCall,
 	ApiChatMessageContentPart,
@@ -43,23 +38,6 @@ import { getAuthHeaders, getJsonHeaders } from '$lib/utils/api-headers';
 import { formatAttachmentText } from '$lib/utils/formatters';
 import { streamIdentity } from '$lib/utils/stream-identity';
 
-function getAudioInputFormat(mimeType: string): AudioInputFormat {
-	const normalizedMimeType = mimeType.trim().toLowerCase();
-
-	if (
-		normalizedMimeType === MimeTypeAudio.WAV ||
-		normalizedMimeType === MimeTypeAudio.WAVE ||
-		normalizedMimeType === MimeTypeAudio.X_WAV ||
-		normalizedMimeType === MimeTypeAudio.X_WAVE ||
-		normalizedMimeType === MimeTypeAudio.VND_WAVE ||
-		normalizedMimeType === MimeTypeAudio.X_PN_WAV
-	) {
-		return FileTypeAudio.WAV;
-	}
-
-	return FileTypeAudio.MP3;
-}
-
 interface ResumableStreamState {
 	bytesReceived: number;
 	updatedAt: number;

tools/ui/src/lib/services/read-media.service.ts

@@ -0,0 +1,112 @@
+import { ToolsService } from './tools.service';
+import {
+	FILE_EXTENSION_SEPARATOR,
+	FILE_PATH_SEPARATOR_REGEX,
+	NEWLINE,
+	PREFIX_FILE,
+	PREFIX_MIME,
+	PREFIX_SIZE,
+	READ_MEDIA_AUDIO_MIME,
+	READ_MEDIA_IMAGE_MIME,
+	RESP_TYPE_BASE64
+} from '$lib/constants';
+import { BuiltInTool, ToolResponseField } from '$lib/enums';
+import type { ToolExecutionResult } from '$lib/types';
+
+/** Modalities of the model the tool call runs for. */
+export interface ReadMediaCapabilities {
+	audio: boolean;
+	vision: boolean;
+}
+
+/** Lowercase extension of a path, without the dot. Empty when the file name has none. */
+function fileExtension(path: string): string {
+	const name = path.split(FILE_PATH_SEPARATOR_REGEX).pop() ?? '';
+	const dot = name.lastIndexOf(FILE_EXTENSION_SEPARATOR);
+
+	return dot > 0 ? name.slice(dot + 1).toLowerCase() : '';
+}
+
+/**
+ * **ReadMediaService** - frontend executor for the `read_media` tool
+ *
+ * The tool is synthetic: no such tool exists on the server. It reads the file
+ * through the built-in `read_file` tool with the `base64` response type, then
+ * turns the bytes into a data URI line. The agentic store lifts that line into
+ * an image or audio attachment on the tool result message, which is what makes
+ * the model perceive the file instead of reading a wall of base64.
+ *
+ * Living in the frontend is what lets it exist only for models that can
+ * actually use the result - the server has no idea which model is selected.
+ *
+ * @see buildReadMediaToolDefinition in constants/read-media.ts - tool schema sent to the LLM
+ * @see agenticStore in stores/agentic.svelte.ts - tool dispatch and attachment extraction
+ */
+export class ReadMediaService {
+	static async executeTool(
+		params: Record<string, unknown>,
+		capabilities: ReadMediaCapabilities,
+		signal?: AbortSignal,
+		cwd?: string
+	): Promise<ToolExecutionResult> {
+		const path = typeof params.path === 'string' ? params.path : '';
+
+		if (!path) {
+			return { content: 'Error: missing "path" argument.', isError: true };
+		}
+
+		const extension = fileExtension(path);
+		const imageMime = READ_MEDIA_IMAGE_MIME[extension];
+		const audioMime = READ_MEDIA_AUDIO_MIME[extension];
+
+		let resolvedMime: string | undefined;
+
+		if (imageMime && capabilities.vision) resolvedMime = imageMime;
+		else if (audioMime && capabilities.audio) resolvedMime = audioMime;
+
+		if (!resolvedMime) {
+			const supported = [
+				...(capabilities.vision ? Object.keys(READ_MEDIA_IMAGE_MIME) : []),
+				...(capabilities.audio ? Object.keys(READ_MEDIA_AUDIO_MIME) : [])
+			];
+			// an unreadable-by-this-model file is a dead end, so say why instead of failing silently
+			const reason =
+				imageMime || audioMime
+					? `the current model cannot perceive ".${extension}" files`
+					: `".${extension}" is not a supported media type`;
+
+			return {
+				content: `Error: ${reason}. Supported: ${supported.join(', ')}.`,
+				isError: true
+			};
+		}
+
+		const raw = await ToolsService.executeToolRaw(
+			BuiltInTool.READ_FILE,
+			{ path },
+			signal,
+			cwd,
+			RESP_TYPE_BASE64
+		);
+
+		if (ToolResponseField.ERROR in raw) {
+			return { content: String(raw[ToolResponseField.ERROR]), isError: true };
+		}
+
+		const base64 = typeof raw.base64 === 'string' ? raw.base64 : '';
+
+		if (!base64) {
+			return { content: `Error: no data returned for ${path}.`, isError: true };
+		}
+
+		const sizeBytes = typeof raw.size_bytes === 'number' ? raw.size_bytes : 0;
+		const content = [
+			`${PREFIX_FILE}${path}`,
+			`${PREFIX_SIZE}${sizeBytes} bytes`,
+			`${PREFIX_MIME}${resolvedMime}`,
+			`data:${resolvedMime};base64,${base64}`
+		].join(NEWLINE);
+
+		return { content, isError: false };
+	}
+}

tools/ui/src/lib/services/tools.service.ts

@@ -1,5 +1,5 @@
 import { base } from '$app/paths';
-import { API_TOOLS, X_TOOL_CWD_HEADER } from '$lib/constants';
+import { API_TOOLS, X_RESP_TYPE_HEADER, X_TOOL_CWD_HEADER } from '$lib/constants';
 import { ToolResponseField } from '$lib/enums';
 import type { ServerBuiltinToolInfo, ToolExecutionResult } from '$lib/types';
 import { apiFetch } from '$lib/utils';
@@ -51,16 +51,26 @@ export class ToolsService {
 	 * Execute a built-in tool and return the raw JSON response. Unlike
 	 * executeTool, this preserves structured fields (e.g. file_glob_search's
 	 * `entries` and `base`) that the flattened ToolExecutionResult drops.
+	 *
+	 * @param respType - sent as the x-resp-type request header. Only read_file
+	 * honors it, with `base64` to get the raw bytes instead of decoded text.
 	 */
 	static async executeToolRaw(
 		toolName: string,
 		params: Record<string, unknown>,
 		signal?: AbortSignal,
-		cwd?: string
+		cwd?: string,
+		respType?: string
 	): Promise<Record<string, unknown>> {
+		const headers: Record<string, string> = {};
+
+		if (cwd) headers[X_TOOL_CWD_HEADER] = cwd;
+
+		if (respType) headers[X_RESP_TYPE_HEADER] = respType;
+
 		return apiFetch<Record<string, unknown>>(API_TOOLS.EXECUTE, {
 			body: JSON.stringify({ params, tool: toolName }),
-			headers: cwd ? { [X_TOOL_CWD_HEADER]: cwd } : undefined,
+			headers: Object.keys(headers).length > 0 ? headers : undefined,
 			method: 'POST',
 			signal
 		});

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

@@ -22,7 +22,9 @@
 
 import { DEFAULT_AGENTIC_CONFIG, NEWLINE } from '$lib/constants';
 import {
+	AUDIO_MIME_TO_EXTENSION,
 	DATA_URI_BASE64_REGEX,
+	DEFAULT_AUDIO_EXTENSION,
 	DEFAULT_IMAGE_EXTENSION,
 	IMAGE_MIME_TO_EXTENSION,
 	MCP_ATTACHMENT_NAME_PREFIX
@@ -36,6 +38,7 @@ import {
 	ToolCallType
 } from '$lib/enums';
 import { ChatService } from '$lib/services';
+import { ReadMediaService } from '$lib/services/read-media.service';
 import { SandboxService } from '$lib/services/sandbox.service';
 import { ToolsService } from '$lib/services/tools.service';
 import { conversationsStore } from '$lib/stores/conversations.svelte';
@@ -75,9 +78,10 @@ import type {
 import type {
 	DatabaseMessage,
 	DatabaseMessageExtra,
+	DatabaseMessageExtraAudioFile,
 	DatabaseMessageExtraImageFile
 } from '$lib/types/database';
-import { isAbortError } from '$lib/utils';
+import { getAudioInputFormat, isAbortError } from '$lib/utils';
 import { SvelteMap } from 'svelte/reactivity';
 
 function createDefaultSession(): AgenticSession {
@@ -900,7 +904,18 @@ class AgenticStore {
 							if (executionResult.isError) toolSuccess = false;
 						} else if (toolSource === ToolSource.FRONTEND) {
 							const args = this.parseToolArguments(toolCall.function.arguments);
-							const executionResult = await SandboxService.executeTool(toolName, args, signal);
+							const executionResult =
+								toolName === BuiltInTool.READ_MEDIA
+									? await ReadMediaService.executeTool(
+											args,
+											{
+												audio: modelsStore.modelSupportsAudio(effectiveModel),
+												vision: modelsStore.modelSupportsVision(effectiveModel)
+											},
+											signal,
+											conversationsStore.activeConversation?.cwd
+										)
+									: await SandboxService.executeTool(toolName, args, signal);
 
 							result = executionResult.content;
 
@@ -990,7 +1005,19 @@ class AgenticStore {
 				];
 
 				for (const attachment of attachments) {
-					if (attachment.type === AttachmentType.IMAGE) {
+					if (attachment.type === AttachmentType.AUDIO) {
+						if (modelsStore.modelSupportsAudio(effectiveModel)) {
+							contentParts.push({
+								input_audio: {
+									data: (attachment as DatabaseMessageExtraAudioFile).base64Data,
+									format: getAudioInputFormat(
+										(attachment as DatabaseMessageExtraAudioFile).mimeType
+									)
+								},
+								type: ContentPartType.INPUT_AUDIO
+							});
+						}
+					} else if (attachment.type === AttachmentType.IMAGE) {
 						if (modelsStore.modelSupportsVision(effectiveModel)) {
 							contentParts.push({
 								image_url: {
@@ -1101,6 +1128,18 @@ class AgenticStore {
 				return `[Attachment saved: ${name}]`;
 			}
 
+			if (mimeType.startsWith(MimeTypePrefix.AUDIO)) {
+				// audio extras hold the bare base64, the input_audio part has no room for a data URI
+				attachments.push({
+					base64Data,
+					mimeType,
+					name,
+					type: AttachmentType.AUDIO
+				});
+
+				return `[Attachment saved: ${name}]`;
+			}
+
 			return line;
 		});
 
@@ -1108,7 +1147,9 @@ class AgenticStore {
 	}
 
 	private buildAttachmentName(mimeType: string, index: number): string {
-		const extension = IMAGE_MIME_TO_EXTENSION[mimeType] ?? DEFAULT_IMAGE_EXTENSION;
+		const extension = mimeType.startsWith(MimeTypePrefix.AUDIO)
+			? (AUDIO_MIME_TO_EXTENSION[mimeType] ?? DEFAULT_AUDIO_EXTENSION)
+			: (IMAGE_MIME_TO_EXTENSION[mimeType] ?? DEFAULT_IMAGE_EXTENSION);
 
 		return `${MCP_ATTACHMENT_NAME_PREFIX}-${Date.now()}-${index}.${extension}`;
 	}

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

@@ -18,7 +18,10 @@ import { ModelsService } from '$lib/services/models.service';
 import { PropsService } from '$lib/services/props.service';
 import { conversationsStore } from '$lib/stores/conversations.svelte';
 import { isRouterMode, serverStore } from '$lib/stores/server.svelte';
-import { getAuthHeaders, TTLCache } from '$lib/utils';
+// deep imports, not the '$lib/utils' barrel: it re-exports modules that reach back
+// into the stores, and going through it here would read a half-built module
+import { getAuthHeaders } from '$lib/utils/api-headers';
+import { TTLCache } from '$lib/utils/cache-ttl';
 import {
 	detectThinkingSupport,
 	detectThinkingSupportWithReason

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

@@ -1,4 +1,5 @@
 import {
+	buildReadMediaToolDefinition,
 	buildSandboxToolDefinition,
 	DISABLED_TOOL_KEYS_LOCALSTORAGE_KEY,
 	HOME_TILDE,
@@ -15,6 +16,7 @@ import {
 } from '$lib/enums';
 import { ToolsService } from '$lib/services/tools.service';
 import { mcpStore } from '$lib/stores/mcp.svelte';
+import { modelsStore, selectedModelName } from '$lib/stores/models.svelte';
 import { config } from '$lib/stores/settings.svelte';
 import type { OpenAIToolDefinition, ToolEntry, ToolGroup } from '$lib/types';
 import { SvelteMap, SvelteSet } from 'svelte/reactivity';
@@ -168,9 +170,42 @@ class ToolsStore {
 	}
 
 	get frontendTools(): OpenAIToolDefinition[] {
-		return config().jsSandboxEnabled
-			? [buildSandboxToolDefinition(!!config().symbolicMathEnabled)]
-			: [];
+		const tools: OpenAIToolDefinition[] = [];
+
+		if (config().jsSandboxEnabled) {
+			tools.push(buildSandboxToolDefinition(!!config().symbolicMathEnabled));
+		}
+
+		const readMedia = this.readMediaTool();
+
+		if (readMedia) tools.push(readMedia);
+
+		return tools;
+	}
+
+	/**
+	 * `read_media` runs in the frontend on top of the server's `read_file`, so it
+	 * exists only when that tool is served and the active model can perceive the
+	 * bytes. The server cannot make this call - it does not know which model the
+	 * conversation uses.
+	 */
+	private readMediaTool(): OpenAIToolDefinition | null {
+		const hasReadFile = this._builtinTools.some(
+			(def) => def.function.name === BuiltInTool.READ_FILE
+		);
+
+		if (!hasReadFile) return null;
+
+		const model = selectedModelName() ?? modelsStore.models[0]?.model ?? '';
+
+		if (!model) return null;
+
+		const vision = modelsStore.modelSupportsVision(model);
+		const audio = modelsStore.modelSupportsAudio(model);
+
+		if (!vision && !audio) return null;
+
+		return buildReadMediaToolDefinition(vision, audio);
 	}
 
 	get customTools(): OpenAIToolDefinition[] {

tools/ui/src/lib/utils/agentic.ts

@@ -50,11 +50,11 @@ export interface AgenticSection {
 }
 
 /**
- * Represents a tool result line that may reference an image attachment
+ * Represents a tool result line that may reference a media attachment (image or audio)
  */
 export type ToolResultLine = {
 	text: string;
-	image?: DatabaseMessageExtraImageFile;
+	media?: DatabaseMessageExtraImageFile | DatabaseMessageExtraAudioFile;
 };
 
 /**
@@ -301,16 +301,16 @@ export function splitSearchSummaryList(
 	return { lines };
 }
 
-/** Bounded cache for parseToolResultWithImages results. */
+/** Bounded cache for parseToolResultWithMedia results. */
 const TOOL_RESULT_LINES_CACHE_MAX_SIZE = 32;
 const toolResultLinesCache = new Map<string, ToolResultLine[]>();
 
 /**
- * Parse tool result text into lines, matching image attachments by name.
+ * Parse tool result text into lines, matching media attachments (images and audio) by name.
  * Memoized: called per render during streaming on unchanged tool result
  * strings with unchanged extras.
  */
-export function parseToolResultWithImages(
+export function parseToolResultWithMedia(
 	toolResult: string,
 	extras?: DatabaseMessageExtra[]
 ): ToolResultLine[] {
@@ -332,12 +332,13 @@ export function parseToolResultWithImages(
 		if (!match || !extras) return { text: line };
 
 		const attachmentName = match[1];
-		const image = extras.find(
-			(e): e is DatabaseMessageExtraImageFile =>
-				e.type === AttachmentType.IMAGE && e.name === attachmentName
+		const media = extras.find(
+			(e): e is DatabaseMessageExtraImageFile | DatabaseMessageExtraAudioFile =>
+				(e.type === AttachmentType.IMAGE || e.type === AttachmentType.AUDIO) &&
+				e.name === attachmentName
 		);
 
-		return { image, text: line };
+		return { media, text: line };
 	});
 
 	if (toolResultLinesCache.size >= TOOL_RESULT_LINES_CACHE_MAX_SIZE) {

tools/ui/src/lib/utils/audio-format.ts

@@ -0,0 +1,22 @@
+import { FileTypeAudio, MimeTypeAudio } from '$lib/enums';
+import type { AudioInputFormat } from '$lib/types/api';
+
+/**
+ * Map a MIME type to the AudioInputFormat expected by the API.
+ */
+export function getAudioInputFormat(mimeType: string): AudioInputFormat {
+	const normalizedMimeType = mimeType.trim().toLowerCase();
+
+	if (
+		normalizedMimeType === MimeTypeAudio.WAV ||
+		normalizedMimeType === MimeTypeAudio.WAVE ||
+		normalizedMimeType === MimeTypeAudio.X_WAV ||
+		normalizedMimeType === MimeTypeAudio.X_WAVE ||
+		normalizedMimeType === MimeTypeAudio.VND_WAVE ||
+		normalizedMimeType === MimeTypeAudio.X_PN_WAV
+	) {
+		return FileTypeAudio.WAV;
+	}
+
+	return FileTypeAudio.MP3;
+}

tools/ui/src/lib/utils/index.ts

@@ -248,7 +248,7 @@ export {
 export {
 	deriveAgenticSections,
 	buildAssistantRawOutput,
-	parseToolResultWithImages,
+	parseToolResultWithMedia,
 	splitSearchSummaryList,
 	hasAgenticContent,
 	classifyToolResult,
@@ -325,3 +325,6 @@ export { uuid } from './uuid';
 
 // CSS utilities
 export { remToPx } from './css';
+
+// Audio format helper (used by agentic store and chat service)
+export { getAudioInputFormat } from './audio-format';

tools/ui/tests/client/README-perf.md

@@ -33,12 +33,12 @@ npx vitest --project=client --run tests/client/agentic-stream.perf.svelte.test.t
 
 The point of the harness is the _scaling curve_, not any single number.
 
-| Knob                        | Reads on                                                                                             |
-| --------------------------- | ---------------------------------------------------------------------------------------------------- |
-| `priorToolCalls` (0/1/5/20) | the reactive fan-out. Flat => no fan-out. Linear => confirmed.                                       |
-| `toolResultBytes`           | whole-blob string scans (`extractSearchResults`, `parseToolResultWithImages`, `classifyToolResult`). |
-| `editFileEdits`             | `computeLineDiff`, the O(m\*n) LCS.                                                                  |
-| `openCodeFence`             | `hljs.highlightAuto` on partial code.                                                                |
+| Knob                        | Reads on                                                                                            |
+| --------------------------- | --------------------------------------------------------------------------------------------------- |
+| `priorToolCalls` (0/1/5/20) | the reactive fan-out. Flat => no fan-out. Linear => confirmed.                                      |
+| `toolResultBytes`           | whole-blob string scans (`extractSearchResults`, `parseToolResultWithMedia`, `classifyToolResult`). |
+| `editFileEdits`             | `computeLineDiff`, the O(m\*n) LCS.                                                                 |
+| `openCodeFence`             | `hljs.highlightAuto` on partial code.                                                               |
 
 Deliberately no hard assertions: CI timing is noisy and the value here is the
 before/after delta, not a gate.

tools/ui/tests/unit/agentic-hotpath.bench.ts

@@ -6,7 +6,7 @@
 //
 // Run: npx vitest bench --project=unit tests/unit/agentic-hotpath.bench.ts
 
-import { classifyToolResult, parseToolResultWithImages } from '$lib/utils/agentic';
+import { classifyToolResult, parseToolResultWithMedia } from '$lib/utils/agentic';
 import { detectIncompleteCodeBlock, highlightCode } from '$lib/utils/code';
 import { computeLineDiff } from '$lib/utils/compute-line-diff';
 import { preprocessLaTeX } from '$lib/utils/latex-protection';
@@ -200,17 +200,17 @@ describe('exit-code regex', () => {
 
 // --- per-line result parsers ----------------------------------------------
 
-describe('parseToolResultWithImages', () => {
+describe('parseToolResultWithMedia', () => {
 	bench('1KB', () => {
-		parseToolResultWithImages(SHELL_OUTPUT_1KB, []);
+		parseToolResultWithMedia(SHELL_OUTPUT_1KB, []);
 	});
 
 	bench('200KB', () => {
-		parseToolResultWithImages(SHELL_OUTPUT_200KB, []);
+		parseToolResultWithMedia(SHELL_OUTPUT_200KB, []);
 	});
 
 	bench('2MB', () => {
-		parseToolResultWithImages(SHELL_OUTPUT_2MB, []);
+		parseToolResultWithMedia(SHELL_OUTPUT_2MB, []);
 	});
 });