demo, wip

2026-08-19 10:46:37 +02:00 · 2026-08-18 00:43:50 +02:00
173 changed files with 2791 additions and 2561 deletions
@@ -1,18 +1,18 @@
-ARG OPENVINO_VERSION_MAJOR=2026.3
-ARG OPENVINO_VERSION_FULL=2026.3.0.22451.bd8d6542e3c
+ARG OPENVINO_VERSION_MAJOR=2026.2.1
+ARG OPENVINO_VERSION_FULL=2026.2.1.21919.ede283a88e3
 ARG UBUNTU_VERSION=24.04

 # Intel GPU driver versions. https://github.com/intel/compute-runtime/releases
-ARG IGC_VERSION=v2.38.2
-ARG IGC_VERSION_FULL=2_2.38.2+22051
-ARG COMPUTE_RUNTIME_VERSION=26.27.39122.11
-ARG COMPUTE_RUNTIME_VERSION_FULL=26.27.39122.11-0
+ARG IGC_VERSION=v2.36.3
+ARG IGC_VERSION_FULL=2_2.36.3+21719
+ARG COMPUTE_RUNTIME_VERSION=26.22.38646.4
+ARG COMPUTE_RUNTIME_VERSION_FULL=26.22.38646.4-0
 ARG IGDGMM_VERSION=22.10.0

 # Intel NPU driver versions. https://github.com/intel/linux-npu-driver/releases
-ARG NPU_DRIVER_VERSION=v1.35.0
-ARG NPU_DRIVER_FULL=v1.35.0.20260722-29947505341
-ARG LIBZE1_VERSION=1.28.2-1~24.04~ppa1
+ARG NPU_DRIVER_VERSION=v1.33.0
+ARG NPU_DRIVER_FULL=v1.33.0.20260529-26625960453
+ARG LIBZE1_VERSION=1.27.0-1~24.04~ppa2

 # Optional proxy build arguments
 ARG http_proxy=
@@ -170,7 +170,7 @@ RUN --mount=type=cache,target=/var/cache/intel-npu,sharing=locked \
    fi; \
    DEB=/var/cache/intel-npu/libze1_${LIBZE1_VERSION}_amd64.deb; \
    if [ ! -f "$DEB" ]; then \
-        wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260606T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
+        wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260324T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
    fi; \
    mkdir /tmp/npu/ && cd /tmp/npu/ && tar -xf "$TGZ" && cp "$DEB" .; \
    apt-get update; \
@@ -0,0 +1,20 @@
+name: "Linux - Setup Vulkan SDK"
+description: "Setup Vulkan SDK for Linux"
+inputs:
+  path:
+    description: "Installation path"
+    required: true
+  version:
+    description: "Vulkan SDK version"
+    required: true
+
+runs:
+  using: "composite"
+  steps:
+    - name: Setup Vulkan SDK
+      id: setup
+      uses: ./.github/actions/unarchive-tar
+      with:
+        url: https://sdk.lunarg.com/sdk/download/${{ inputs.version }}/linux/vulkan_sdk.tar.xz
+        path: ${{ inputs.path }}
+        strip: 1
@@ -6,7 +6,8 @@ inputs:
    required: true
  cuda_arch:
    description: "CUDA target architecture"
-    required: true
+    required: false
+    default: "x64"

 runs:
  using: "composite"
@@ -10,6 +10,33 @@ concurrency:
  cancel-in-progress: true

 jobs:
+  ubuntu-24-vulkan-cache:
+    runs-on: ubuntu-24.04
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+
+      - name: Get latest Vulkan SDK version
+        id: vulkan_sdk_version
+        run: |
+          echo "VULKAN_SDK_VERSION=$(curl https://vulkan.lunarg.com/sdk/latest/linux.txt)" >> "$GITHUB_ENV"
+
+      - name: Setup Cache
+        uses: actions/cache@v5
+        id: cache-sdk
+        with:
+          path: ./vulkan_sdk
+          key: cache-gha-vulkan-sdk-${{ env.VULKAN_SDK_VERSION }}-${{ runner.os }}
+
+      - name: Setup Vulkan SDK
+        if: steps.cache-sdk.outputs.cache-hit != 'true'
+        uses: ./.github/actions/linux-setup-vulkan
+        with:
+          path: ./vulkan_sdk
+          version: ${{ env.VULKAN_SDK_VERSION }}
+
  #ubuntu-24-spacemit-cache:
  #  runs-on: ubuntu-24.04

@@ -40,9 +67,9 @@ jobs:
    runs-on: ubuntu-24.04

    env:
-      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Clone
@@ -69,8 +96,8 @@ jobs:

    env:
      # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Clone
@@ -21,7 +21,6 @@ on:
    paths: [
      '.github/workflows/build-cpu.yml',
      '.github/workflows/build-cmake-pkg.yml',
-      'ggml/src/ggml-rpc/**',
      '**/CMakeLists.txt',
      '**/.cmake',
      '**/*.h',
@@ -124,6 +123,7 @@ jobs:
    env:
      OPENBLAS_VERSION: 0.3.23
      SDE_VERSION: 9.33.0-2024-01-07
+      VULKAN_VERSION: 1.4.357.0

    strategy:
      matrix:
@@ -134,6 +134,9 @@ jobs:
          - build: 'x64-openblas'
            arch: 'x64'
            defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_OPENMP=OFF -DGGML_BLAS=ON -DGGML_BLAS_VENDOR=OpenBLAS -DBLAS_INCLUDE_DIRS="$env:RUNNER_TEMP/openblas/include" -DBLAS_LIBRARIES="$env:RUNNER_TEMP/openblas/lib/openblas.lib"'
+          - build: 'x64-vulkan'
+            arch: 'x64'
+            defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DCMAKE_BUILD_TYPE=Release -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_VULKAN=ON'
          - build: 'arm64'
            arch: 'arm64'
            defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON'
@@ -164,6 +167,15 @@ jobs:
          $lib =  $(join-path $msvc 'bin\Hostx64\x64\lib.exe')
          & $lib /machine:x64 "/def:${env:RUNNER_TEMP}/openblas/lib/libopenblas.def" "/out:${env:RUNNER_TEMP}/openblas/lib/openblas.lib" /name:openblas.dll

+      - name: Install Vulkan SDK
+        id: get_vulkan
+        if: ${{ matrix.build == 'x64-vulkan' }}
+        run: |
+          curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe"
+          & "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install
+          Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}"
+          Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin"
+
      - name: Install Ninja
        id: install_ninja
        run: |
@@ -22,7 +22,6 @@ env:

 jobs:
  cuda:
-    name: windows-cuda (${{ matrix.cuda }}, ${{ matrix.arch }})
    runs-on: windows-2022

    permissions:
@@ -30,16 +29,7 @@ jobs:

    strategy:
      matrix:
-        include:
-          - cuda: '12.4'
-            arch: x64
-            defines: '-DGGML_CUDA_CUB_3DOT2=ON'
-          - cuda: '13.3'
-            arch: x64
-            defines: ''
-          - cuda: '13.4'
-            arch: arm64
-            defines: '-DCMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-msvc-cuda.cmake'
+        cuda: ['12.4', '13.3']

    steps:
      - name: Clone
@@ -49,13 +39,12 @@ jobs:
      - name: ccache
        uses: ggml-org/ccache-action@v1.2.21
        with:
-          key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
+          key: release-windows-2022-x64-cuda-${{ matrix.cuda }}

      - name: Install Cuda Toolkit
        uses: ./.github/actions/windows-setup-cuda
        with:
          cuda_version: ${{ matrix.cuda }}
-          cuda_arch: ${{ matrix.arch }}

      - name: Install Ninja
        id: install_ninja
@@ -65,21 +54,26 @@ jobs:
      - name: Build
        id: cmake_build
        shell: cmd
+        # TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project
        run: |
-          call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" ${{ matrix.arch == 'x64' && 'x64' || 'amd64_arm64' }}
+          call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64
          cmake -S . -B build -G "Ninja Multi-Config" ^
-            -DGGML_BACKEND_DL=ON ^
+            -DLLAMA_BUILD_SERVER=ON ^
+            -DLLAMA_BUILD_BORINGSSL=ON ^
            -DGGML_NATIVE=OFF ^
-            -DGGML_CPU=OFF ^
+            -DGGML_BACKEND_DL=ON ^
+            -DGGML_CPU_ALL_VARIANTS=ON ^
            -DGGML_CUDA=ON ^
-            -DLLAMA_BUILD_BORINGSSL=ON ${{ matrix.defines }}
+            -DGGML_RPC=ON ^
+            -DGGML_CUDA_CUB_3DOT2=ON
          set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1
-          cmake --build build --config Release -j %NINJA_JOBS% --target ggml-cuda
+          cmake --build build --config Release -j %NINJA_JOBS% -t ggml
+          cmake --build build --config Release

      - name: ccache-clear
        uses: ./.github/actions/ccache-clear
        with:
-          key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
+          key: release-windows-2022-x64-cuda-${{ matrix.cuda }}

  hip:
    runs-on: windows-2022
@@ -39,8 +39,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Clone
@@ -81,7 +81,7 @@ jobs:
        # TODO: fix and re-enable the `test-llama-archs` test below
        run: |
          cd ${{ github.workspace }}
-          ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" --verbose --timeout 2000
+          ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 2000

      - name: Test (GPU)
        id: cmake_test_gpu
@@ -89,15 +89,15 @@ jobs:
        run: |
          cd ${{ github.workspace }}
          export GGML_OPENVINO_DEVICE=GPU
-          ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" --verbose --timeout 3000
+          ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 3000

  openvino-windows-2022:
    runs-on: windows-2022

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Clone
@@ -166,4 +166,4 @@ jobs:
          call "%OPENVINO_ROOT%\setupvars.bat"

          cd build
-          ctest --test-dir ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" -C Release --verbose --timeout 3000
+          ctest --test-dir ReleaseOV -L main -E "test-llama-archs" -C Release --verbose --timeout 3000
@@ -0,0 +1,66 @@
+name: CI (rpc)
+
+on:
+  workflow_dispatch: # allows manual triggering
+  push:
+    branches:
+      - master
+    paths: [
+      '.github/workflows/build-rpc.yml',
+      '**/CMakeLists.txt',
+      '**/.cmake',
+      '**/*.h',
+      '**/*.hpp',
+      '**/*.c',
+      '**/*.cpp'
+    ]
+
+  pull_request:
+    types: [opened, synchronize, reopened]
+    paths: [
+      '.github/workflows/build-rpc.yml',
+      'ggml/src/ggml-rpc/**'
+    ]
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
+  cancel-in-progress: true
+
+env:
+  GGML_NLOOP: 3
+  GGML_N_THREADS: 1
+  LLAMA_ARG_LOG_COLORS: 1
+  LLAMA_ARG_LOG_PREFIX: 1
+  LLAMA_ARG_LOG_TIMESTAMPS: 1
+
+jobs:
+  ubuntu-24-rpc:
+    runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }}
+
+    continue-on-error: true
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+
+      - name: Dependencies
+        id: depends
+        run: |
+          sudo apt-get update
+          sudo apt-get install build-essential libssl-dev ninja-build
+
+      - name: Build
+        id: cmake_build
+        run: |
+          cmake -B build \
+            -G "Ninja" \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DGGML_RPC=ON
+          time cmake --build build --config Release -j $(nproc)
+
+      - name: Test
+        id: cmake_test
+        run: |
+          cd build
+          ctest -L main --verbose
@@ -288,8 +288,8 @@ jobs:

    env:
      # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Clone
@@ -93,13 +93,19 @@ jobs:
        run: |
          echo "VULKAN_SDK_VERSION=$(curl https://vulkan.lunarg.com/sdk/latest/linux.txt)" >> "$GITHUB_ENV"

-      - name: Setup Vulkan SDK
-        id: setup
-        uses: ./.github/actions/unarchive-tar
+      - name: Use Vulkan SDK Cache
+        uses: actions/cache@v5
+        id: cache-sdk
        with:
-          url: https://sdk.lunarg.com/sdk/download/${{ env.VULKAN_SDK_VERSION }}/linux/vulkan_sdk.tar.xz
          path: ./vulkan_sdk
-          strip: 1
+          key: cache-gha-vulkan-sdk-${{ env.VULKAN_SDK_VERSION }}-${{ runner.os }}
+
+      - name: Setup Vulkan SDK
+        if: steps.cache-sdk.outputs.cache-hit != 'true'
+        uses: ./.github/actions/linux-setup-vulkan
+        with:
+          path: ./vulkan_sdk
+          version: ${{ env.VULKAN_SDK_VERSION }}

      - name: ccache
        uses: ggml-org/ccache-action@v1.2.21
@@ -127,56 +133,3 @@ jobs:
          # This is using llvmpipe and runs slower than other backends
          # test-backend-ops is too slow on llvmpipe, skip it
          ctest -L main -E test-backend-ops --verbose --timeout 900
-
-  windows:
-    runs-on: windows-2025
-
-    env:
-      VULKAN_VERSION: 1.4.357.0
-
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-
-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.21
-        with:
-          key: cpu-windows-2025-x64-vulkan
-          variant: ccache
-          evict-old-files: 1d
-          save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
-
-      - name: Install Vulkan SDK
-        id: get_vulkan
-        run: |
-          curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe"
-          & "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install
-          Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}"
-          Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin"
-
-      - name: Install Ninja
-        id: install_ninja
-        run: |
-          choco install ninja
-
-      - name: Build
-        id: cmake_build
-        run: |
-          cmake -S . -B build -G "Ninja Multi-Config" `
-            -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake `
-            -DCMAKE_BUILD_TYPE=Release `
-            -DGGML_NATIVE=OFF `
-            -DLLAMA_BUILD_SERVER=ON `
-            -DGGML_RPC=ON `
-            -DGGML_BACKEND_DL=ON `
-            -DGGML_CPU_ALL_VARIANTS=ON `
-            -DGGML_VULKAN=ON `
-            -DLLAMA_BUILD_BORINGSSL=ON
-          cmake --build build --config Release -j ${env:NUMBER_OF_PROCESSORS}
-
-      - name: Test
-        id: cmake_test
-        run: |
-          cd build
-          ctest -L main -C Release --verbose --timeout 900
@@ -49,33 +49,6 @@ jobs:
          git push origin "${VERSION}"
          echo "Created and pushed tag ${VERSION}"

-      - name: Generate release description
-        id: desc
-        run: bash scripts/make-release-desc.sh "${{ steps.checks.outputs.version }}"
-        env:
-          GITHUB_REPOSITORY: ${{ github.repository }}
-
-      - name: Create release
-        if: ${{ github.event.inputs.dry_run == 'false' }}
-        uses: ggml-org/action-create-release@v1
-        env:
-          GITHUB_TOKEN: ${{ github.token }}
-        with:
-          tag_name: ${{ steps.checks.outputs.version }}
-          # TODO: remove the prerelease flag once the semantic versioning workflow is ready
-          # ref: https://github.com/ggml-org/ggml/discussions/1579
-          prerelease: true
-          body: |
-            > [!NOTE]
-            > Semantic versioning is still work in progress.
-            > More info can be found in https://github.com/ggml-org/ggml/discussions/1579
-
-            ${{ steps.desc.outputs.nightly }}
-
-            ## ${{ steps.desc.outputs.changelog_title }}
-
-            ${{ steps.desc.outputs.changelog }}
-
      - name: Dry run summary
        if: ${{ github.event.inputs.dry_run == 'true' }}
        run: |
@@ -446,8 +446,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Set OpenVINO version output
@@ -562,8 +562,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3"
-      OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
+      OPENVINO_VERSION_MAJOR: "2026.2.1"
+      OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"

    steps:
      - name: Set OpenVINO version output
@@ -5,7 +5,7 @@ include(CheckIncludeFileCXX)
 ### llama.cpp version
 set(LLAMA_VERSION_MAJOR 0)
 set(LLAMA_VERSION_MINOR 1)
-set(LLAMA_VERSION_PATCH 2)
+set(LLAMA_VERSION_PATCH 1)
 set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")

 # whether this is a development/nightly build
@@ -224,10 +224,12 @@ add_subdirectory(src)
 # utils, programs, examples and tests
 #

-add_subdirectory(vendor)
+# mtmd needs this even when common is not built
+add_subdirectory(vendor/hash)

 if (LLAMA_BUILD_COMMON)
    add_subdirectory(common)
+    add_subdirectory(vendor/cpp-httplib)
 endif()

 if (LLAMA_BUILD_COMMON AND LLAMA_BUILD_TESTS AND NOT CMAKE_JS_VERSION)
@@ -7,11 +7,10 @@
 <b>LLM inference in C/C++</b>

 [![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT)
-[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp?filter=v*)](https://github.com/ggml-org/llama.cpp/releases?q=tag:v0)
-[![Nightly](https://img.shields.io/github/v/release/ggml-org/llama.cpp?label=nightly)](https://github.com/ggml-org/llama.cpp/releases)
+[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp)](https://github.com/ggml-org/llama.cpp/releases)
 [![Server](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml)
-[![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
-[![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
+[![Docker](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
+[![Winget](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)

 [manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3A0cc4m%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [compile times](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-compile-times.md) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)

@@ -120,7 +119,7 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
 ## Acknowledgements

 - [yhirose/cpp-httplib](https://github.com/yhirose/cpp-httplib) - Single-header HTTP server, used by `llama-server` - MIT license
- [nothings/stb](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
+- [stb-image](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
 - [nlohmann/json](https://github.com/nlohmann/json) - Single-header JSON library, used by various tools/examples - MIT License
- [mackron/miniaudio](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
- [sheredom/subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
+- [miniaudio.h](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
+- [subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
@@ -290,7 +290,6 @@ combine_static_libraries() {
        "${base_dir}/${build_dir}/ggml/src/ggml-metal/${release_dir}/libggml-metal.a"
        "${base_dir}/${build_dir}/ggml/src/ggml-blas/${release_dir}/libggml-blas.a"
        "${base_dir}/${build_dir}/tools/mtmd/${release_dir}/libmtmd.a"
-        "${base_dir}/${build_dir}/vendor/hash/${release_dir}/libvendor-hash.a"
    )

    # Create temporary directory for processing
@@ -190,7 +190,7 @@ if [ ! -z ${GG_BUILD_OPENVINO} ]; then
    CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_OPENVINO=ON"

    # TODO: fix and re-enable the `test-llama-archs` test below
-    CTEST_EXTRA="-E test-llama-archs|test-recurrent-state-rollback-nemotron-h"
+    CTEST_EXTRA="-E test-llama-archs"
 fi

 ## helpers
@@ -126,8 +126,7 @@ set_target_properties(${TARGET} PROPERTIES
    MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
 )

-target_include_directories(${TARGET} PUBLIC .)
-target_link_libraries     (${TARGET} PUBLIC vendor::nlohmann vendor::sheredom)
+target_include_directories(${TARGET} PUBLIC . ../vendor)
 target_compile_features   (${TARGET} PUBLIC cxx_std_17)

 if (LLAMA_SUBPROCESS)
@@ -1710,6 +1710,38 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
            params.cache_ram_mib = value;
        }
    ).set_env("LLAMA_ARG_CACHE_RAM").set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}));
+    add_opt(common_arg(
+        {"-cdisk", "--cache-disk"}, "PATH",
+        "directory for the disk prompt cache; prompts evicted from the RAM cache are saved here and restored on later requests, including across restarts (default: disabled, requires cache-ram)",
+        [](common_params & params, const std::string & value) {
+            params.cache_disk_path = value;
+            if (!fs_is_directory(params.cache_disk_path)) {
+                throw std::invalid_argument("not a directory: " + value);
+            }
+            // if doesn't end with DIRECTORY_SEPARATOR, add it
+            if (params.cache_disk_path[params.cache_disk_path.size() - 1] != DIRECTORY_SEPARATOR) {
+                params.cache_disk_path += DIRECTORY_SEPARATOR;
+            }
+        }
+    ).set_env("LLAMA_ARG_CACHE_DISK").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--cache-disk-limit"}, "N",
+        string_format("total size budget of the disk prompt cache directory in MiB; oldest entries are deleted when exceeded (default: %d, -1 - no limit)", params.cache_disk_limit_mib),
+        [](common_params & params, int value) {
+            if (value == 0 || value < -1) {
+                throw std::invalid_argument("cache-disk-limit must be positive or -1 (no limit)");
+            }
+            params.cache_disk_limit_mib = value;
+        }
+    ).set_env("LLAMA_ARG_CACHE_DISK_LIMIT").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--cache-disk-write-through"},
+        {"--no-cache-disk-write-through"},
+        "write prompts to the disk cache every time they are saved to the RAM cache, instead of only when evicted from it (default: disabled)",
+        [](common_params & params, bool value) {
+            params.cache_disk_write_through = value;
+        }
+    ).set_env("LLAMA_ARG_CACHE_DISK_WRITE_THROUGH").set_examples({LLAMA_EXAMPLE_SERVER}));
    add_opt(common_arg(
        {"-kvu", "--kv-unified"},
        {"-no-kvu", "--no-kv-unified"},
@@ -1750,18 +1750,6 @@ struct ggml_threadpool_params ggml_threadpool_params_from_cpu_params(const commo
    return tpp;
 }

-namespace {
-
-bool can_share_threadpool(const ggml_threadpool_params & tpp1, const ggml_threadpool_params & tpp2) {
-    // n_threads does not matter -> we'll use what's larger
-    ggml_threadpool_params tpp_comparison = tpp1;
-    tpp_comparison.n_threads = tpp2.n_threads;
-
-    return ggml_threadpool_params_match(&tpp_comparison, &tpp2);
-}
-
-}  // namespace
-
 common_threadpools::~common_threadpools() {
    if (!free_fn) {
        return;
@@ -1790,9 +1778,7 @@ void common_threadpools::init(llama_context * ctx, const common_params & params)
    struct ggml_threadpool_params tpp =
            ggml_threadpool_params_from_cpu_params(params.cpuparams);

-    if (can_share_threadpool(tpp, tpp_batch)) {
-        tpp.n_threads = std::max(tpp.n_threads, tpp_batch.n_threads);
-    } else {
+    if (!ggml_threadpool_params_match(&tpp, &tpp_batch)) {
        threadpool_batch = ggml_threadpool_new_fn(&tpp_batch);
        if (!threadpool_batch) {
            COM_WRN("batch threadpool create failed : n_threads %d\n", tpp_batch.n_threads);
@@ -614,6 +614,10 @@ struct common_params {
    int32_t checkpoint_min_step = 8192;  // minimum spacing between context checkpoints
    int32_t cache_ram_mib       = 8192;  // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.

+    std::string cache_disk_path;                 // disk prompt cache directory, empty = disabled
+    int32_t     cache_disk_limit_mib     = -1;   // total size budget for the disk prompt cache dir, -1 = no limit
+    bool        cache_disk_write_through = false; // also write to disk whenever a prompt is saved to the RAM cache
+
    std::string hostname      = "127.0.0.1";
    std::string public_path   = "";                                                                         // NOLINT
    std::string api_prefix    = "";                                                                         // NOLINT
@@ -237,8 +237,8 @@ chmod +x ubuntu-llamacpp-ov-install.sh
 # ============================================
 set -euo pipefail

-OPENVINO_VERSION_MAJOR="2026.3"
-OPENVINO_VERSION_FULL="2026.3.0.22451.bd8d6542e3c"
+OPENVINO_VERSION_MAJOR="2026.2.1"
+OPENVINO_VERSION_FULL="2026.2.1.21919.ede283a88e3"

 SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
 OPENVINO_INSTALL_DIR="/opt/intel/openvino_${OPENVINO_VERSION_MAJOR}"
@@ -334,7 +334,7 @@ echo "  ./build/ReleaseOV/bin/llama-cli -m model.gguf"
 ```

 > [!NOTE]
-> The script pins OpenVINO `2026.3` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release.
+> The script pins OpenVINO `2026.2.1` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release.

 </details>

@@ -364,8 +364,8 @@ REM ============================================
 REM llama.cpp OpenVINO Build Script (Ninja)
 REM ============================================

-set "OPENVINO_VERSION_MAJOR=2026.3"
-set "OPENVINO_VERSION_FULL=2026.3.0.22451.bd8d6542e3c"
+set "OPENVINO_VERSION_MAJOR=2026.2.1"
+set "OPENVINO_VERSION_FULL=2026.2.1.21919.ede283a88e3"

 set "SCRIPT_DIR=%~dp0"
 set "VCPKG_DIR=C:\vcpkg"
@@ -547,7 +547,7 @@ endlocal
 ```

 > [!NOTE]
-> The script pins OpenVINO `2026.3` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release. From any new shell, source the matching `setupvars` script via the junction — `call "C:\Intel\openvino\setupvars.bat"` from `cmd`, or `& "C:\Intel\openvino\setupvars.ps1"` from PowerShell. If `winget` cannot register Visual Studio Build Tools on first run, install them once manually and re-run the script from an elevated **Developer Command Prompt for VS 2022**.
+> The script pins OpenVINO `2026.2.1` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release. From any new shell, source the matching `setupvars` script via the junction — `call "C:\Intel\openvino\setupvars.bat"` from `cmd`, or `& "C:\Intel\openvino\setupvars.ps1"` from PowerShell. If `winget` cannot register Visual Studio Build Tools on first run, install them once manually and re-run the script from an elevated **Developer Command Prompt for VS 2022**.

 </details>

@@ -2,5 +2,5 @@ set(TARGET llama-gguf-hash)
 add_executable(${TARGET} gguf-hash.cpp)
 install(TARGETS ${TARGET} RUNTIME)

-target_link_libraries(${TARGET} PRIVATE vendor::hash ggml ${CMAKE_THREAD_LIBS_INIT})
+target_link_libraries(${TARGET} PRIVATE vendor-hash ggml ${CMAKE_THREAD_LIBS_INIT})
 target_compile_features(${TARGET} PRIVATE cxx_std_17)
@@ -17,15 +17,15 @@
 extern "C" {
 #endif

-#include "hash/xxhash/xxhash.h"
-#include "hash/sha256/sha256.h"
+#include "xxhash/xxhash.h"
+#include "sha256/sha256.h"

 #ifdef __cplusplus
 }
 #endif

 // sha1 is compiled as C++ and lives in a namespace, see scripts/sync_vendor.py
-#include "hash/sha1/sha1.h"
+#include "sha1/sha1.h"
 using namespace vendor_hash;


@@ -5,7 +5,7 @@ project("ggml" C CXX ASM)
 ### GGML Version
 set(GGML_VERSION_MAJOR 0)
 set(GGML_VERSION_MINOR 20)
-set(GGML_VERSION_PATCH 2)
+set(GGML_VERSION_PATCH 1)
 set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")

 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
@@ -7,7 +7,7 @@ extern "C" {
 #endif

 #define RPC_PROTO_MAJOR_VERSION    5
-#define RPC_PROTO_MINOR_VERSION    1
+#define RPC_PROTO_MINOR_VERSION    0
 #define RPC_PROTO_PATCH_VERSION    0

 #ifdef  __cplusplus
@@ -4,7 +4,6 @@
 #include "vecdotq.cuh"

 #include <cstdint>
-#include <type_traits>

 typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);

@@ -70,8 +69,7 @@ enum mmvq_parameter_table_id {
    MMVQ_PARAMETERS_GCN,
    MMVQ_PARAMETERS_RDNA2,
    MMVQ_PARAMETERS_RDNA3_0,
-    MMVQ_PARAMETERS_RDNA4,
-    MMVQ_PARAMETERS_GB10
+    MMVQ_PARAMETERS_RDNA4
 };

 static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
@@ -85,8 +83,6 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
    return MMVQ_PARAMETERS_GCN;
 #elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE
    return MMVQ_PARAMETERS_TURING;
-#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK
-    return MMVQ_PARAMETERS_GB10;
 #else
    return MMVQ_PARAMETERS_GENERIC;
 #endif
@@ -108,9 +104,6 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
    if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) {
        return MMVQ_PARAMETERS_TURING;
    }
-    if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) {
-        return MMVQ_PARAMETERS_GB10;
-    }
    return MMVQ_PARAMETERS_GENERIC;
 }

@@ -358,7 +351,7 @@ static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() {
 #endif
 }

-static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id, bool small_k = false, bool halve_iters = false) {
+static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id) {
    if (table_id == MMVQ_PARAMETERS_GENERIC) {
        switch (ncols_dst) {
            case 1:
@@ -461,32 +454,11 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
                return 1;
        }
    }
-    if (table_id == MMVQ_PARAMETERS_GB10) {
-        const int generic = calc_nwarps(type, ncols_dst, MMVQ_PARAMETERS_GENERIC);
-        // Only worth the wider block when it actually retires the K loop in half the trips (Observation)
-        if (ncols_dst == 1 && !small_k && halve_iters) {
-            switch (type) {
-                case GGML_TYPE_Q4_0:
-                case GGML_TYPE_Q4_1:
-                case GGML_TYPE_Q5_0:
-                case GGML_TYPE_Q5_1:
-                case GGML_TYPE_Q8_0:
-                case GGML_TYPE_Q4_K:
-                case GGML_TYPE_Q5_K:
-                case GGML_TYPE_Q6_K:
-                case GGML_TYPE_IQ4_NL:
-                    return 2 * generic;
-                default:
-                    break;
-            }
-        }
-        return generic;
-    }
    return 1;
 }

 static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int table_id, bool small_k = false, int nwarps = 1) {
-    if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING || table_id == MMVQ_PARAMETERS_GB10) {
+    if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING) {
        switch (ncols_dst) {
            case 1:
                return small_k ? nwarps : 1;
@@ -505,8 +477,8 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
    return 1;
 }

-template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false, bool halve_iters = false>
-__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id(), small_k, halve_iters)*ggml_cuda_get_physical_warp_size(), 1)
+template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
+__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
 static __global__ void mul_mat_vec_q(
        const void * vx_ptr, const void * vy_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
        const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -523,7 +495,7 @@ static __global__ void mul_mat_vec_q(
    constexpr int qi  = ggml_cuda_type_traits<type>::qi;
    constexpr int vdr = get_vdr_mmvq(type);
    constexpr mmvq_parameter_table_id table_id = get_device_table_id();
-    constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
+    constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id);
    constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
    constexpr int warp_size = ggml_cuda_get_physical_warp_size();

@@ -801,8 +773,8 @@ static __global__ void mul_mat_vec_q_moe(
 template<ggml_type type>
 static std::pair<dim3, dim3> calc_launch_params(
        const int ncols_dst, const int nrows_x, const int nchannels_dst, const int nsamples_or_ntokens,
-        const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false, const bool halve_iters = false) {
-    const int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
+        const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false) {
+    const int nwarps = calc_nwarps(type, ncols_dst, table_id);
    const int rpb = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
    const int64_t nblocks = (nrows_x + rpb - 1) / rpb;
    const dim3 block_nums(nblocks, nchannels_dst, nsamples_or_ntokens);
@@ -810,7 +782,7 @@ static std::pair<dim3, dim3> calc_launch_params(
    return {block_nums, block_dims};
 }

-template<ggml_type type, int c_ncols_dst, bool small_k = false, bool halve_iters = false>
+template<ggml_type type, int c_ncols_dst, bool small_k = false>
 static void mul_mat_vec_q_switch_fusion(
        const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
        const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -825,7 +797,7 @@ static void mul_mat_vec_q_switch_fusion(
    if constexpr (c_ncols_dst == 1) {
        if (has_fusion) {
            const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
-            ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k, halve_iters>, launch_params,
+            ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k>, launch_params,
                 vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
                 channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
                 sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -836,7 +808,7 @@ static void mul_mat_vec_q_switch_fusion(
    GGML_ASSERT(!has_fusion && "fusion only supported for ncols_dst=1");

    const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
-    ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k, halve_iters>, launch_params,
+    ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k>, launch_params,
        vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
        channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
        sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -888,18 +860,16 @@ static void mul_mat_vec_q_switch_ncols_dst(

    const bool has_ids = ids != nullptr;

-    // How the K loop divides up at the baseline block width, both decisions below use these.
-    constexpr int qk                    = ggml_cuda_type_traits<type>::qk;
-    constexpr int qi                    = ggml_cuda_type_traits<type>::qi;
-    constexpr int vdr                   = get_vdr_mmvq(type);
-    const int     blocks_per_row_x      = ncols_x / qk;
-    const int     blocks_per_iter_1warp = vdr * warp_size / qi;
-
    const auto should_use_small_k = [&](int c_ncols_dst) {
        // When K is small, increase rows_per_block to match nwarps so each warp has more work to do
        // Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle.
-        const int  nwarps = calc_nwarps(type, c_ncols_dst, table_id);
-        bool       use    = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;
+        constexpr int qk                    = ggml_cuda_type_traits<type>::qk;
+        constexpr int qi                    = ggml_cuda_type_traits<type>::qi;
+        constexpr int vdr                   = get_vdr_mmvq(type);
+        const int     blocks_per_row_x      = ncols_x / qk;
+        const int     blocks_per_iter_1warp = vdr * warp_size / qi;
+        const int     nwarps                = calc_nwarps(type, c_ncols_dst, table_id);
+        bool          use                   = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;

        constexpr std::array<ggml_type, 2> iq_slow_turing = {
            GGML_TYPE_IQ3_XXS,
@@ -932,28 +902,6 @@ static void mul_mat_vec_q_switch_ncols_dst(
        return use;
    };

-    // Whether doubling nwarps pays off on the ncols_dst == 1 path, where K sets the K loop trip count.
-    const auto should_halve_iters = [&] {
-        if (table_id != MMVQ_PARAMETERS_GB10) {
-            return false;
-        }
-
-        // Expert rows are gathered per token, so a wider block adds reduction work without reuse.
-        if (has_ids) {
-            return false;
-        }
-
-        const int blocks_per_iter = calc_nwarps(type, 1, table_id) * blocks_per_iter_1warp;
-        const int iters           = (blocks_per_row_x + blocks_per_iter - 1) /  blocks_per_iter;
-        const int iters_wide      = (blocks_per_row_x + blocks_per_iter * 2 - 1) / (blocks_per_iter * 2);
-
-        // An odd trip count leaves half the wider block idle for its last iteration, that tail is
-        // only affordable once the loop is long enough to dilute it to an eighth of the work (observation).
-        const int idle = iters_wide * 2 - iters;
-
-        return idle * 8 <= iters_wide * 2;
-    };
-
    if (has_ids && ncols_dst > 1) {
        // Multi-token MUL_MAT_ID path - dedicated MoE kernel
        mul_mat_vec_q_moe_launch<type>(
@@ -966,34 +914,26 @@ static void mul_mat_vec_q_switch_ncols_dst(

    switch (ncols_dst) {
        case 1: {
-            // static, else MSVC lambda capture breaks the constexpr uses below
-            static constexpr int c_ncols_dst = 1;
+            constexpr int c_ncols_dst = 1;

-            // Tag types keep the flags compile-time, so __launch_bounds__ matches what is launched.
-            const auto launch = [&](auto small_k_tag, auto halve_iters_tag) {
-                constexpr bool c_small_k = decltype(small_k_tag)::value;
-                // Types the table does not promote would compile a second, identical kernel.
-                constexpr bool c_promoted =
-                    calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, true) !=
-                    calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, false);
+            bool use_small_k = should_use_small_k(c_ncols_dst);

-                constexpr bool c_halve_iters = decltype(halve_iters_tag)::value && c_promoted;
-
-                const std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
-                                                                              nsamples_dst, warp_size, table_id, c_small_k, c_halve_iters);
-                mul_mat_vec_q_switch_fusion<type, c_ncols_dst, c_small_k, c_halve_iters>(
+            if (use_small_k) {
+                std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
+                                                                        nsamples_dst, warp_size, table_id, true);
+                mul_mat_vec_q_switch_fusion<type, c_ncols_dst, true>(
                    vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
                    channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
                    stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
                    stream);
-            };
-
-            if (should_use_small_k(c_ncols_dst)) {
-                launch(std::true_type{},  std::false_type{});
-            } else if (should_halve_iters()) {
-                launch(std::false_type{}, std::true_type{});
            } else {
-                launch(std::false_type{}, std::false_type{});
+                std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
+                                                                        nsamples_dst, warp_size, table_id);
+                mul_mat_vec_q_switch_fusion<type, c_ncols_dst>(
+                    vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
+                    channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
+                    stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
+                    stream);
            }
        } break;
        case 2: {
@@ -118,17 +118,6 @@ __kernel void flash_attn_f16(
    __local DATA_TYPE4 l_v[BLOCK_N][DV_VEC];

    for (int k_start = 0; k_start < n_kv; k_start += BLOCK_N) {
-#if WG_SIZE > FA_SG
-        // WAR on l_k/l_v: a thread that finishes the compute below early — either
-        // it skipped it (my_query_row >= n_q, the continue) or its subgroup simply
-        // ran ahead — wraps around and reloads the tiles while another subgroup is
-        // still reading them. Any WG that is exactly one lockstep subgroup
-        // (WG_SIZE == FA_SG) cannot diverge and hides this; a WG spanning multiple
-        // subgroups (Intel sg=32, or BLOCK_M > 64 on Adreno) corrupts the result.
-        // All threads reach this each iteration (no-op on the first), so it does
-        // not diverge with the continue. Compiled out when WG == one subgroup.
-        barrier(CLK_LOCAL_MEM_FENCE);
-#endif
        for (int i = tid; i < BLOCK_N * DK_VEC; i += WG_SIZE) {
            const int row = i / DK_VEC;
            const int col = i % DK_VEC;
@@ -119,15 +119,13 @@ __kernel void flash_attn_f32(
    __local DATA_TYPE4 l_v[BLOCK_N][DV_VEC];

    for (int k_start = 0; k_start < n_kv; k_start += BLOCK_N) {
-#if WG_SIZE > FA_SG
-        // WAR on l_k/l_v: a thread that finishes the compute below early — either
-        // it skipped it (my_query_row >= n_q, the continue) or its subgroup simply
-        // ran ahead — wraps around and reloads the tiles while another subgroup is
-        // still reading them. Any WG that is exactly one lockstep subgroup
-        // (WG_SIZE == FA_SG) cannot diverge and hides this; a WG spanning multiple
-        // subgroups (Intel sg=32, or BLOCK_M > 64 on Adreno) corrupts the result.
-        // All threads reach this each iteration (no-op on the first), so it does
-        // not diverge with the continue. Compiled out when WG == one subgroup.
+#if FA_SG < 64
+        // WAR on l_k/l_v: threads with my_query_row >= n_q skip the compute below
+        // (continue) and would race ahead to reload the tiles while active threads
+        // still read them. A single 64-wide Adreno subgroup (WG == sg) runs lockstep
+        // and hides this; a WG that spans multiple narrower subgroups (Intel sg=32)
+        // corrupts the result. All threads reach this each iteration (no-op on the
+        // first), so it does not diverge with the continue. Compiled out at sg=64.
        barrier(CLK_LOCAL_MEM_FENCE);
 #endif
        for (int i = tid; i < BLOCK_N * DK_VEC; i += WG_SIZE) {
@@ -47,7 +47,7 @@ struct rpc_tensor {
    uint64_t data;
    char name[GGML_MAX_NAME];

-    int32_t use_count;
+    char padding[4];
 };

 static_assert(sizeof(rpc_tensor) % 8 == 0, "rpc_tensor size must be multiple of 8");
@@ -447,7 +447,7 @@ static rpc_tensor serialize_tensor(const ggml_tensor * tensor) {

    // Avoid sending uninitialized data over the wire
    memset(result.name, 0, sizeof(result.name));
-    result.use_count = 0;
+    memset(result.padding, 0, sizeof(result.padding));

    snprintf(result.name, GGML_MAX_NAME, "%s", tensor->name);
    return result;
@@ -675,7 +675,7 @@ static void ggml_backend_rpc_synchronize(ggml_backend_t backend) {
    // this is no-op because we don't have any async operations
 }

-static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
+static void add_tensor(ggml_tensor * tensor, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
    if (tensor == nullptr) {
        return;
    }
@@ -684,15 +684,10 @@ static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::ve
    }
    visited.insert(tensor);
    for (int i = 0; i < GGML_MAX_SRC; i++) {
-        add_tensor(tensor->src[i], cgraph, tensors, visited);
+        add_tensor(tensor->src[i], tensors, visited);
    }
-    add_tensor(tensor->view_src, cgraph, tensors, visited);
-    rpc_tensor result = serialize_tensor(tensor);
-    const size_t hash_pos = ggml_hash_find(&cgraph->visited_hash_set, tensor);
-    if (hash_pos != GGML_HASHSET_FULL && ggml_bitset_get(cgraph->visited_hash_set.used, hash_pos)) {
-        result.use_count = cgraph->use_counts[hash_pos];
-    }
-    tensors.push_back(result);
+    add_tensor(tensor->view_src, tensors, visited);
+    tensors.push_back(serialize_tensor(tensor));
 }

 static void serialize_graph(uint32_t device, const ggml_cgraph * cgraph, std::vector<uint8_t> & output) {
@@ -700,7 +695,7 @@ static void serialize_graph(uint32_t device, const ggml_cgraph * cgraph, std::ve
    std::vector<rpc_tensor> tensors;
    std::unordered_set<ggml_tensor*> visited;
    for (uint32_t i = 0; i < n_nodes; i++) {
-        add_tensor(cgraph->nodes[i], cgraph, tensors, visited);
+        add_tensor(cgraph->nodes[i], tensors, visited);
    }
    // serialization format:
    // | device (4 bytes) | n_nodes (4 bytes) | nodes (n_nodes * sizeof(uint64_t) | n_tensors (4 bytes) | tensors (n_tensors * sizeof(rpc_tensor)) |
@@ -1456,10 +1451,6 @@ bool rpc_server::graph_compute(const std::vector<uint8_t> & input) {
            GGML_LOG_ERROR("[%s] failed to create graph node %d (id=%" PRId64 ")\n", __func__, i, id);
            return false;
        }
-        if (graph->nodes[i] != nullptr) {
-            const size_t hash_pos = ggml_hash_insert(&graph->visited_hash_set, graph->nodes[i]);
-            graph->use_counts[hash_pos] = tensor_ptrs.at(id)->use_count;
-        }
    }
    ggml_status status = ggml_backend_graph_compute(backends[device], graph);
    GGML_ASSERT(status == GGML_STATUS_SUCCESS && "Unsuccessful graph computations are not supported with RPC");
@@ -1,119 +0,0 @@
-#include "fwht.hpp"
-
-#include <cmath>
-
-template <int N>
-static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows,
-                        const float scale, const sycl::nd_item<2> & item) {
-    const sycl::sub_group sg = item.get_sub_group();
-
-    const int64_t r = item.get_global_id(0);
-    if (r >= n_rows) {
-        return;
-    }
-
-    src += r * N;
-    dst += r * N;
-
-    constexpr int el_w = N / WARP_SIZE;
-    static_assert(el_w >= 1 && N % WARP_SIZE == 0, "row must be a whole number of sub-group widths");
-
-    float     reg[el_w];
-    const int lane = sg.get_local_linear_id();
-
-#pragma unroll
-    for (int i = 0; i < el_w; ++i) {
-        reg[i] = src[i * WARP_SIZE + lane] * scale;
-    }
-
-    // Butterflies inside the sub-group. The partner of a lane with bit h clear is the
-    // lower index of the pair, so it takes the sum and the upper takes lower - upper.
-#pragma unroll
-    for (int h = 1; h < WARP_SIZE; h *= 2) {
-#pragma unroll
-        for (int j = 0; j < el_w; ++j) {
-            const float val  = reg[j];
-            const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE);
-
-            reg[j] = (lane & h) == 0 ? val + val2 : val2 - val;
-        }
-    }
-
-    // Butterflies across registers: h is a multiple of WARP_SIZE, so the partner of
-    // element i*WARP_SIZE + lane lives in reg[i + h/WARP_SIZE] on the same lane.
-#pragma unroll
-    for (int h = WARP_SIZE; h < N; h *= 2) {
-        const int step = h / WARP_SIZE;
-#pragma unroll
-        for (int j = 0; j < el_w; j += 2 * step) {
-#pragma unroll
-            for (int k = 0; k < step; ++k) {
-                const float x = reg[j + k];
-                const float y = reg[j + k + step];
-
-                reg[j + k]        = x + y;
-                reg[j + k + step] = x - y;
-            }
-        }
-    }
-
-#pragma unroll
-    for (int i = 0; i < el_w; ++i) {
-        dst[i * WARP_SIZE + lane] = reg[i];
-    }
-}
-
-template <int N>
-static void launch_fwht(const float * src, float * dst, const int64_t n_rows, const float scale,
-                        dpct::queue_ptr stream) {
-    constexpr int rows_per_block = 4;
-
-    const int64_t num_blocks = (n_rows + rows_per_block - 1) / rows_per_block;
-
-    // dim 1 is the fastest-varying, so a sub-group is exactly one row's WARP_SIZE lanes.
-    const sycl::range<2> global(num_blocks * rows_per_block, WARP_SIZE);
-    const sycl::range<2> local(rows_per_block, WARP_SIZE);
-
-    stream->parallel_for(sycl::nd_range<2>(global, local),
-                         [=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                             fwht_kernel<N>(src, dst, n_rows, scale, item);
-                         });
-}
-
-bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
-    if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
-        return false;
-    }
-    if (!ggml_are_same_shape(src, dst)) {
-        return false;
-    }
-    if (!ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) {
-        return false;
-    }
-
-    const int     n    = (int) src->ne[0];
-    const int64_t rows = ggml_nrows(src);
-
-    const float *   src_d  = (const float *) src->data;
-    float *         dst_d  = (float *) dst->data;
-    dpct::queue_ptr stream = ctx.stream();
-
-    const float scale = 1.0f / std::sqrt((float) n);
-
-    switch (n) {
-        case 64:
-            launch_fwht<64>(src_d, dst_d, rows, scale, stream);
-            return true;
-        case 128:
-            launch_fwht<128>(src_d, dst_d, rows, scale, stream);
-            return true;
-        case 256:
-            launch_fwht<256>(src_d, dst_d, rows, scale, stream);
-            return true;
-        case 512:
-            launch_fwht<512>(src_d, dst_d, rows, scale, stream);
-            return true;
-        default:
-            return false;
-    }
-}
@@ -1,12 +0,0 @@
-#ifndef GGML_SYCL_FWHT_HPP
-#define GGML_SYCL_FWHT_HPP
-
-#include "common.hpp"
-
-// Fast Walsh-Hadamard transform, the fast path for a MUL_MAT whose src0 ggml has
-// tagged GGML_HINT_SRC0_IS_HADAMARD. src0 is not read at all. Returns false if the
-// shape is not one this can serve, in which case the caller must fall through to the
-// ordinary mat-mul dispatch.
-bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst);
-
-#endif  // GGML_SYCL_FWHT_HPP
@@ -58,7 +58,6 @@
 #include "ggml-sycl/backend.hpp"
 #include "ggml-sycl/common.hpp"
 #include "ggml-sycl/element_wise.hpp"
-#include "ggml-sycl/fwht.hpp"
 #include "ggml-sycl/gemm.hpp"
 #include "ggml-sycl/getrows.hpp"
 #include "ggml-sycl/norm.hpp"
@@ -4474,18 +4473,6 @@ static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor *

 static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
-
-    // Handle HADAMARAD hint given from further up the pipeline and pass it to the correct
-    // kernel.
-    //
-    // The op check is not redundant: this backend also routes MUL_MAT_ID through here with a
-    // stack copy of dst, which carries MUL_MAT_ID's own op_params. ggml_mul_mat_set_hint()
-    // asserts GGML_OP_MUL_MAT for the same reason.
-    if (dst->op == GGML_OP_MUL_MAT && ggml_get_op_params_i32(dst, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
-        ggml_sycl_op_fwht(ctx, src1, dst)) {
-        return;
-    }
-
    const bool split = ggml_backend_buffer_is_sycl_split(src0->buffer);
    int64_t min_compute_capability = INT_MAX;

@@ -1,87 +0,0 @@
-#!/bin/bash
-# Generate the description of a release: the previous release version, the
-# change log and the link to the nightly release corresponding to the commit being released.
-#
-# Usage: make-release-desc.sh <version>
-#   <version>: current release version (v<maj>.<min>.<pat>, the leading v is optional)
-#
-# The previous version is the highest plain semver tag (v<maj>.<min>.<pat>)
-# strictly below <version>. The change log lists all commits between the
-# previous version tag and the release commit, one line per commit.
-#
-# The release commit is the commit <version> points at when the tag exists,
-# HEAD otherwise. The nightly release is the b* tag pointing at that commit
-# (release.yml tags the same commit); the link is only generated when that
-# tag exists.
-#
-# Env (when running in GitHub Actions):
-#   GITHUB_OUTPUT: previous_tag, changelog_title, changelog and nightly are written here
-#   GITHUB_REPOSITORY: owner/repo, used to build the nightly release URL (skipped when unset)
-set -euo pipefail
-
-if [[ $# -ne 1 ]]; then
-    echo "Usage: $(basename "$0") <version>"
-    exit 1
-fi
-VERSION="$1"
-
-# Accept the version with or without the leading v, reject anything else
-if [[ "${VERSION}" =~ ^[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
-    VERSION="v${VERSION}"
-elif [[ ! "${VERSION}" =~ ^v[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
-    echo "Error: invalid version '${VERSION}' (expected v<maj>.<min>.<pat>)"
-    exit 1
-fi
-
-# Make sure all remote tags are available locally (skipped on local runs without origin)
-if ! git fetch --tags origin 2>/dev/null; then
-    echo "Warning: could not fetch tags from origin (local run?)"
-fi
-
-# Release commit: the commit <version> points at when the tag exists, HEAD otherwise.
-if ! RELEASE_COMMIT="$(git rev-parse -q --verify "refs/tags/${VERSION}^{commit}" 2>/dev/null)"; then
-    RELEASE_COMMIT="$(git rev-parse HEAD)"
-fi
-
-echo "Release commit: $(git rev-parse --short "${RELEASE_COMMIT}")"
-
-PREV="$( { git tag --list; echo "${VERSION}"; } \
-    | grep -E '^v[0-9]+\.[0-9]+\.[0-9]+$' \
-    | sort -V \
-    | awk -v cur="${VERSION}" '$0 == cur { exit } { prev = $0 } END { print prev }')"
-
-if [[ -n "${PREV}" ]]; then
-    CHANGELOG="$(git log --oneline "${PREV}..${RELEASE_COMMIT}")"
-    CHANGELOG_TITLE="Change log since ${PREV}"
-else
-    CHANGELOG="(no previous release tag found)"
-    CHANGELOG_TITLE="Change log"
-fi
-
-# Nightly release: the b* tag pointing at the release commit (|| true: no match is not an error)
-NIGHTLY_TAG="$(git tag --points-at "${RELEASE_COMMIT}" | grep -E '(^|-)b[0-9]+(-[0-9a-f]{7})?$' | head -n 1 || true)"
-
-NIGHTLY=""
-if [[ -n "${NIGHTLY_TAG}" ]]; then
-    if [[ -n "${GITHUB_REPOSITORY:-}" ]]; then
-        NIGHTLY_URL="https://github.com/${GITHUB_REPOSITORY}/releases/tag/${NIGHTLY_TAG}"
-        NIGHTLY="**Nightly build:** [${NIGHTLY_TAG}](${NIGHTLY_URL})"
-        echo "Nightly release: ${NIGHTLY_URL}"
-    fi
-else
-    echo "No nightly release found for commit $(git rev-parse --short "${RELEASE_COMMIT}")"
-fi
-
-echo "Previous version: ${PREV:-none}"
-echo "${CHANGELOG}"
-
-if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
-    {
-        echo "previous_tag=${PREV}"
-        echo "changelog_title=${CHANGELOG_TITLE}"
-        echo "nightly=${NIGHTLY}"
-        echo "changelog<<CHANGELOG_EOF"
-        echo "${CHANGELOG}"
-        echo "CHANGELOG_EOF"
-    } >> "${GITHUB_OUTPUT}"
-fi
@@ -1 +1 @@
-8c63e70982c95ceb862e3a1073a2c1beef75d60a
+3834fd814e74e8af277939dabd69ecc780affd21
@@ -1395,11 +1395,6 @@ void llama_model_loader::get_mapping_range(size_t * first, size_t * last, void *
    }
 }

-void llama_model_loader::unmap_weight(const llama_tensor_weight & w) const {
-    if (!use_mmap) { return; }
-    mappings.at(w.idx)->unmap_fragment(w.offs, w.offs + ggml_nbytes(w.tensor));
-}
-
 void llama_model_loader::load_data_for(struct ggml_tensor * cur) const {
    const auto & w = require_weight(ggml_get_name(cur));

@@ -194,9 +194,6 @@ struct llama_model_loader {

    void get_mapping_range(size_t * first, size_t * last, void ** addr, int idx, ggml_context * ctx) const;

-    // release a weight's mmap pages
-    void unmap_weight(const llama_tensor_weight & w) const;
-
    // for backwards compatibility, does not support ggml-backend
    void load_data_for(struct ggml_tensor * cur) const;

@@ -1270,7 +1270,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
            total_size_org += tensor_size;
            total_size_new += new_size;

-            // update the gguf metadata as we go
+            // update the gguf meta data as we go
            gguf_set_tensor_type(ctx_outs[cur_split].get(), metadata[i].name.c_str(), new_type);
            GGML_ASSERT(gguf_get_tensor_size(ctx_outs[cur_split].get(), gguf_find_tensor(ctx_outs[cur_split].get(), metadata[i].name.c_str())) == new_size);
            gguf_set_tensor_data(ctx_outs[cur_split].get(), metadata[i].name.c_str(), new_data);
@@ -1278,10 +1278,6 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
            // write tensor data + padding
            fout.write((const char *) new_data, new_size);
            zeros(fout, GGML_PAD(new_size, align) - new_size);
-
-            // unmap the tensor to free memory
-            if (ml.use_mmap) { ml.unmap_weight(weight); }
-
        } // no --dry-run
    } // main loop

@@ -1241,7 +1241,7 @@ std::vector<std::string> unicode_regex_split(const std::string & text, const std
        { unicode_cpt_flags::LETTER,      "\x41-\x5A\x61-\x7A" }, // A-Za-z
        { unicode_cpt_flags::PUNCTUATION, "\x21-\x23\x25-\x2A\x2C-\x2F\x3A-\x3B\x3F-\x40\\\x5B-\\\x5D\x5F\\\x7B\\\x7D" }, // !-#%-*,-/:-;?-@\[-\]_\{\}
        { unicode_cpt_flags::ACCENT_MARK, "" }, // no sub-128 codepoints
-        { unicode_cpt_flags::SYMBOL,      "\\\x24\\\x2B\x3C-\x3E\x5E\x60\\\x7C\\\x7E" }, // $+<=>^`|~
+        { unicode_cpt_flags::SYMBOL,      "\\\x24\\\x2B\x3C-\x3E\x5E\x60\\\x7C" }, // $+<=>^`|
    };

    // compute collapsed codepoints only if needed by at least one regex
@@ -116,8 +116,6 @@ function(llama_build_and_test source)
    set_property(TEST ${TEST_TARGET} PROPERTY LABELS ${LLAMA_TEST_LABEL})
 endfunction()

-llama_build_and_test(test-unicode.cpp)
-
 # build test-tokenizer-0 target once and add many tests
 llama_build(test-tokenizer-0.cpp)

@@ -312,9 +310,6 @@ llama_build_and_test(test-mtmd-c-api.c)
 target_link_libraries(${LLAMA_TEST_NAME} PRIVATE mtmd)
 unset(LLAMA_TEST_NAME)

-llama_build_and_test(test-mtmd-impl.cpp)
-target_link_libraries(test-mtmd-impl PRIVATE mtmd)
-
 # GGUF model data fetcher library for tests that need real model metadata
 # Only compile when cpp-httplib has SSL support (CPPHTTPLIB_OPENSSL_SUPPORT)
 if (TARGET cpp-httplib)
@@ -1,88 +0,0 @@
-#include "testing.h"
-
-#include "mtmd-image.h"
-
-#include <iostream>
-#include <string>
-#include <utility>
-#include <vector>
-
-// this test file contains:
-// 1. test cases for mtmd helpers
-// 2. test cases for internal mtmd components
-// internal headers can be included here
-
-struct test_registry {
-    using fn_t = void (*)(testing &);
-
-    struct entry {
-        std::string name;
-        fn_t fn;
-    };
-
-    static std::vector<entry> & all() {
-        static std::vector<entry> entries;
-        return entries;
-    }
-
-    test_registry(const char * name, fn_t fn) {
-        all().push_back({ name, fn });
-    }
-};
-
-#define MAKE_TEST(name)                                               \
-    static void name(testing & t);                                    \
-    static const test_registry test_registry_ ## name(#name, &name);  \
-    static void name(testing & t)
-
-
-//
-// mtmd_image
-//
-
-MAKE_TEST(test_image_preprocessor_lfm2) {
-    clip_hparams hparams;
-    hparams.patch_size = 16;
-    hparams.n_merge = 2;
-    hparams.set_limit_image_tokens(64, 256);
-
-    // { image size, expected tiling }
-    const std::vector<std::pair<clip_image_size, bool>> cases = {
-        { {  704, 704 }, false },
-        // 720 / (patch_size * n_merge) is exactly 22.5, so this only matches HF
-        // if round_by_factor rounds half to even (22) instead of away from zero (23)
-        { {  720, 720 }, false },
-        { {  736, 736 }, true  },
-        { { 1024, 977 }, true  },
-        { { 1056, 384 }, false },
-    };
-
-    for (const auto & [size, expected] : cases) {
-        const bool actual = mtmd_image_preprocessor_lfm2::should_tile(hparams, size);
-
-        t.assert_equal(
-            "tiling for " + std::to_string(size.width) + "x" + std::to_string(size.height),
-            std::string(expected ? "tiled" : "single"),
-            std::string(actual   ? "tiled" : "single"));
-    }
-}
-
-//
-// main
-//
-
-int main(int argc, char ** argv) {
-    testing t(std::cout);
-    t.verbose = true;
-
-    // usage: test-mtmd-impl [filter_regex]
-    for (int i = 1; i < argc; i++) {
-        t.set_filter(argv[i]);
-    }
-
-    for (const auto & e : test_registry::all()) {
-        t.test(e.name, e.fn);
-    }
-
-    return t.summary();
-}
@@ -1,24 +0,0 @@
-#include "../src/unicode.h"
-
-#include <cstdio>
-#include <string>
-#include <vector>
-
-int main() {
-    const std::vector<std::string> regex_exprs = {
-        "[~][A-Za-z]+| ?[\\p{S}]+|\\s+",
-    };
-    const std::vector<std::string> expected = { " ~", "foo" };
-    const auto actual = unicode_regex_split(" ~foo", regex_exprs, false);
-
-    if (actual != expected) {
-        fprintf(stderr, "unexpected split:");
-        for (const auto & piece : actual) {
-            fprintf(stderr, " [%s]", piece.c_str());
-        }
-        fprintf(stderr, "\n");
-        return 1;
-    }
-
-    return 0;
-}
@@ -78,8 +78,10 @@ set_target_properties(mtmd PROPERTIES
 )

 target_link_libraries     (mtmd PUBLIC ggml llama)
-target_link_libraries     (mtmd PRIVATE Threads::Threads vendor::hash vendor::miniaudio vendor::stb vendor::sheredom)
+target_link_libraries     (mtmd PRIVATE Threads::Threads vendor-hash)
 target_include_directories(mtmd PUBLIC  .)
+target_include_directories(mtmd PRIVATE ../..)
+target_include_directories(mtmd PRIVATE ../../vendor)
 target_compile_features   (mtmd PRIVATE cxx_std_17)

 if (MTMD_VIDEO)
@@ -90,9 +92,6 @@ if (BUILD_SHARED_LIBS)
    set_target_properties     (mtmd PROPERTIES POSITION_INDEPENDENT_CODE ON)
    target_compile_definitions(mtmd PRIVATE LLAMA_BUILD)
    target_compile_definitions(mtmd PUBLIC  LLAMA_SHARED)
-
-    # export all symbols so that internal components can be tested by test-mtmd-impl
-    set_target_properties     (mtmd PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS ON)
 endif()

 set(MTMD_PUBLIC_HEADERS
@@ -858,9 +858,6 @@ static std::ifstream open_ifstream_binary(const std::string & fname) {
 }
 #endif

-// in test-mtmd-impl, we include woth common.h and this file, and these functions are duplicated
-// this is a quick fix to avoid compilation errors
-#ifndef DIRECTORY_SEPARATOR
 static std::string string_format(const char * fmt, ...) {
    va_list ap;
    va_list ap2;
@@ -918,7 +915,6 @@ inline bool string_ends_with(std::string_view str, std::string_view suffix) {
    return str.size() >= suffix.size() &&
           str.compare(str.size() - suffix.size(), suffix.size(), suffix) == 0;
 }
-#endif

 //
 // gguf utils
@@ -88,22 +88,6 @@ static ggml_tensor * get_rel_pos(ggml_context * ctx0,
    return cur;  // [C, k_size, q_size]
 }

-// ggml_conv_2d with the im2col kept in F32: the F16 im2col it emits since #23660 degrades OCR
-static ggml_tensor * conv_2d_f32(ggml_context * ctx0, ggml_tensor * a, ggml_tensor * b,
-                                 int s0, int s1, int p0, int p1, int d0, int d1) {
-    const ggml_type im2col_type = a->type == GGML_TYPE_F16 ? GGML_TYPE_F16 : GGML_TYPE_F32;
-    ggml_tensor * im2col = ggml_im2col(ctx0, a, b, s0, s1, p0, p1, d0, d1, true, im2col_type); // [N, OH, OW, IC * KH * KW]
-
-    ggml_tensor * result = ggml_mul_mat(ctx0,
-        ggml_reshape_2d(ctx0, im2col, im2col->ne[0], im2col->ne[3] * im2col->ne[2] * im2col->ne[1]),
-        ggml_reshape_2d(ctx0, a, (a->ne[0] * a->ne[1] * a->ne[2]), a->ne[3]));
-
-    result = ggml_reshape_4d(ctx0, result, im2col->ne[1], im2col->ne[2], im2col->ne[3], a->ne[3]); // [OC, N, OH, OW]
-    result = ggml_cont(ctx0, ggml_permute(ctx0, result, 0, 1, 3, 2)); // [N, OC, OH, OW]
-
-    return result;
-}
-

 ggml_tensor * clip_graph_deepseekocr::build_sam(ggml_tensor * inp_raw) {
    // Building SAM
@@ -117,8 +101,7 @@ ggml_tensor * clip_graph_deepseekocr::build_sam(ggml_tensor * inp_raw) {

    ggml_tensor * inpL;

-    inpL = conv_2d_f32(ctx0, model.patch_embed_proj_w, inp_raw,
-                       (int) model.patch_embed_proj_w->ne[0], (int) model.patch_embed_proj_w->ne[1], 0, 0, 1, 1);
+    inpL = ggml_conv_2d_sk_p0(ctx0, model.patch_embed_proj_w, inp_raw);
    inpL = ggml_add(ctx0, inpL, ggml_reshape_3d(ctx0, model.patch_embed_proj_b, 1, 1, n_embd));
    inpL = ggml_cont(ctx0, ggml_permute(ctx0, inpL, 1, 2, 0, 3));

@@ -246,18 +229,18 @@ ggml_tensor * clip_graph_deepseekocr::build_sam(ggml_tensor * inp_raw) {

    cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 2, 0, 1, 3));

-    cur = conv_2d_f32(ctx0, model.neck_0_w, cur, 1, 1, 0, 0, 1, 1);
+    cur = ggml_conv_2d(ctx0, model.neck_0_w, cur, 1, 1, 0, 0, 1, 1);
    cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 1, 2, 0, 3));
    cur = build_norm(cur, model.neck_1_w, model.neck_1_b, NORM_TYPE_NORMAL, sam_eps, -1);
    cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 2, 0, 1, 3));

-    cur = conv_2d_f32(ctx0, model.neck_2_w, cur, 1, 1, 1, 1, 1, 1);
+    cur = ggml_conv_2d(ctx0, model.neck_2_w, cur, 1, 1, 1, 1, 1, 1);
    cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 1, 2, 0, 3));
    cur = build_norm(cur, model.neck_3_w, model.neck_3_b, NORM_TYPE_NORMAL, sam_eps, -1);
    cur = ggml_cont(ctx0, ggml_permute(ctx0, cur, 2, 0, 1, 3));

-    cur = conv_2d_f32(ctx0, model.net_2, cur, 2, 2, 1, 1, 1, 1);
-    cur = conv_2d_f32(ctx0, model.net_3, cur, 2, 2, 1, 1, 1, 1);
+    cur = ggml_conv_2d(ctx0, model.net_2, cur, 2, 2, 1, 1, 1, 1);
+    cur = ggml_conv_2d(ctx0, model.net_3, cur, 2, 2, 1, 1, 1, 1);
    cb(cur, "sam_output", -1);

    ggml_build_forward_expand(gf, cur);
@@ -12,7 +12,7 @@
 #include "mtmd-helper-common.h"
 #include "llama.h"

-#include "hash/hash.h"
+#include "hash.h"

 #include <algorithm>
 #include <cinttypes>
@@ -1013,31 +1013,14 @@ mtmd_image_preproc_out mtmd_image_preprocessor_lfm2::preprocess(const clip_image
    return output;
 }

-bool mtmd_image_preprocessor_lfm2::should_tile(
-        const clip_hparams & hparams,
-        const clip_image_size & original_size) {
-    const int align_size = hparams.patch_size * hparams.n_merge;
-
-    const auto round_by_factor = [align_size](float x) {
-        // see https://github.com/ggml-org/llama.cpp/pull/27057#discussion_r3796264887
-        return static_cast<int>(std::nearbyint(static_cast<double>(x) / align_size)) * align_size;
-    };
-
-    const int h_bar = std::max(hparams.patch_size, round_by_factor(original_size.height));
-    const int w_bar = std::max(hparams.patch_size, round_by_factor(original_size.width));
-
-    return static_cast<double>(h_bar) * static_cast<double>(w_bar) >
-           static_cast<double>(hparams.image_max_pixels) * max_pixels_tolerance;
-}
-
 mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_lfm2::get_slice_instructions(const clip_image_size & original_size) {
    mtmd_image_preprocessor_llava_uhd::slice_instructions inst;
    const int align_size = hparams.patch_size * hparams.n_merge;
    inst.overview_size = img_tool::calc_size_preserved_ratio(
                            original_size,
                            { align_size, hparams.image_min_pixels, hparams.image_max_pixels, 0 });
-
-    const bool needs_tiling = should_tile(hparams, original_size);
+    // tile if either dimension exceeds tile_size with tolerance
+    const bool needs_tiling = original_size.width > tile_size * max_pixels_tolerance || original_size.height > tile_size * max_pixels_tolerance;

    if (!needs_tiling) {
        inst.refined_size = clip_image_size{0, 0};
@@ -148,8 +148,6 @@ struct mtmd_image_preprocessor_lfm2 : mtmd_image_preprocessor_llava_uhd {
    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
    slice_instructions get_slice_instructions(const clip_image_size & original_size) override;

-    static bool should_tile(const clip_hparams & hparams, const clip_image_size & original_size);
-
 private:
    clip_image_size find_closest_aspect_ratio(
            float aspect_ratio,
@@ -2023,6 +2023,7 @@ int llama_perplexity(int argc, char ** argv) {
    }

    const int32_t n_ctx = params.n_ctx;
+
    if (n_ctx <= 0) {
        LOG_ERR("%s: perplexity tool requires '--ctx-size' > 0\n", __func__);
        return 1;
@@ -5,6 +5,8 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_BINARY_DIR})
 set(TARGET server-context)

 add_library(${TARGET} STATIC
+    server-cache-disk.cpp
+    server-cache-disk.h
    server-chat.cpp
    server-chat.h
    server-task.cpp
@@ -31,7 +33,7 @@ endif()

 target_include_directories(${TARGET} PRIVATE ../mtmd)
 target_include_directories(${TARGET} PRIVATE ${CMAKE_SOURCE_DIR})
-target_link_libraries(${TARGET} PUBLIC llama-common mtmd ${CMAKE_THREAD_LIBS_INIT})
+target_link_libraries(${TARGET} PUBLIC llama-common mtmd vendor-hash ${CMAKE_THREAD_LIBS_INIT})

 # llama-server-impl: server logic, reusable by app

@@ -189,7 +189,7 @@ This endpoint is intended to be used internally by the Web UI and subject to cha
 Get a list of tools, each tool has these fields:
 - `tool` (string): the ID name of the tool, to be used in POST call. Example: `read_file`
 - `display_name` (string): the name to be displayed on UI. Example: `Read file`
- `type` (string): `"server"` for a server tool, or `"mcp"` for a tool exposed by an MCP server
+- `type` (string): `"builtin"` for a built-in tool, or `"mcp"` for a tool exposed by an MCP server
 - `permissions` (object): a mapping string --> boolean that indicates the permission required by this tool. This is useful for the UI to ask the user before calling the tool. For now, the only permission supported is `"write"`
 - `definition` (object): the OAI-compat definition of this tool

@@ -164,6 +164,9 @@ For the full list of features, please refer to [server's changelog](https://gith
 | `-ctxcp, --ctx-checkpoints, --swa-checkpoints N` | max number of context checkpoints to create per slot (default: 32)[(more info)](https://github.com/ggml-org/llama.cpp/pull/15293)<br/>(env: LLAMA_ARG_CTX_CHECKPOINTS) |
 | `-cms, --checkpoint-min-step N` | minimum spacing between context checkpoints in tokens (default: 8192, 0 = no minimum)<br/>(env: LLAMA_ARG_CHECKPOINT_MIN_SPACING_NT) |
 | `-cram, --cache-ram N` | set the maximum cache size in MiB (default: 8192, -1 - no limit, 0 - disable)[(more info)](https://github.com/ggml-org/llama.cpp/pull/16391)<br/>(env: LLAMA_ARG_CACHE_RAM) |
+| `-cdisk, --cache-disk PATH` | directory for the disk prompt cache; prompts evicted from the RAM cache are saved here and restored on later requests, including across restarts (default: disabled, requires cache-ram)<br/>(env: LLAMA_ARG_CACHE_DISK) |
+| `--cache-disk-limit N` | total size budget of the disk prompt cache directory in MiB; oldest entries are deleted when exceeded (default: -1, -1 - no limit)<br/>(env: LLAMA_ARG_CACHE_DISK_LIMIT) |
+| `--cache-disk-write-through, --no-cache-disk-write-through` | write prompts to the disk cache every time they are saved to the RAM cache, instead of only when evicted from it (default: disabled)<br/>(env: LLAMA_ARG_CACHE_DISK_WRITE_THROUGH) |
 | `-kvu, --kv-unified, -no-kvu, --no-kv-unified` | use single unified KV buffer shared across all sequences (default: enabled if number of slots is auto)<br/>(env: LLAMA_ARG_KV_UNIFIED) |
 | `--cache-idle-slots, --no-cache-idle-slots` | save idle slots to the prompt cache on new task, and clear them when using unified KV (default: enabled, requires cache-ram)<br/>(env: LLAMA_ARG_CACHE_IDLE_SLOTS) |
 | `--context-shift, --no-context-shift` | whether to use context shift on infinite text generation (default: disabled)<br/>(env: LLAMA_ARG_CONTEXT_SHIFT) |
@@ -196,11 +199,11 @@ For the full list of features, please refer to [server's changelog](https://gith
 | `--ui-config, --webui-config JSON` | JSON that provides default UI settings (overrides UI defaults)<br/>(env: LLAMA_ARG_UI_CONFIG) |
 | `--ui-config-file, --webui-config-file PATH` | JSON file that provides default UI settings (overrides UI defaults)<br/>(env: LLAMA_ARG_UI_CONFIG_FILE) |
 | `--ui-mcp-proxy, --webui-mcp-proxy, --no-ui-mcp-proxy, --no-webui-mcp-proxy` | experimental: whether to enable MCP CORS proxy - do not enable in untrusted environments (default: disabled)<br/>(env: LLAMA_ARG_UI_MCP_PROXY) |
-| `--tools TOOL1,TOOL2,...` | experimental: whether to enable server tools for AI agents - do not enable in untrusted environments (default: no tools)<br/>specify "all" to enable all tools<br/>available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_info<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_TOOLS) |
+| `--tools TOOL1,TOOL2,...` | experimental: whether to enable built-in tools for AI agents - do not enable in untrusted environments (default: no tools)<br/>specify "all" to enable all tools<br/>available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_info<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_TOOLS) |
 | `--tools-runtime OPTION` | experimental: run tools in a separate runtime environment (default: none, use host environment)<br/>available options:<br/>  'docker:<image>', 'podman:<image>': spin up a new container and reuse it for all invocations, clean up on server exit<br/>  'docker-container:<id>', 'podman-container:<id>': use an existing container by ID, won't stop on server exit<br/>  'ssh:<target>': run tools on a remote POSIX host over SSH, key-based auth and a trusted host key are required<br/><br/>(env: LLAMA_ARG_TOOLS_RUNTIME) |
 | `--mcp-servers-config PATH` | experimental: path to JSON file with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_MCP_SERVERS_CONFIG) |
 | `--mcp-servers-json JSON` | experimental: inline JSON with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_MCP_SERVERS_JSON) |
-| `-ag, --agent, -no-ag, --no-agent` | whether to enable CORS proxy and all server tools - do not enable in untrusted environments (default: disabled)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_AGENT) |
+| `-ag, --agent, -no-ag, --no-agent` | whether to enable CORS proxy and all built-in tools - do not enable in untrusted environments (default: disabled)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_AGENT) |
 | `--ui, --webui, --no-ui, --no-webui` | whether to enable the Web UI (default: enabled)<br/>(env: LLAMA_ARG_UI) |
 | `--embedding, --embeddings` | restrict to only support embedding use case; use only with dedicated embedding models (default: disabled)<br/>(env: LLAMA_ARG_EMBEDDINGS) |
 | `--rerank, --reranking` | enable reranking endpoint on server (default: disabled)<br/>(env: LLAMA_ARG_RERANKING) |
@@ -327,6 +330,22 @@ services:
      LLAMA_ARG_PORT: 8080
 ```

+### Prompt disk cache
+
+The server keeps recently used prompts (their processed KV cache state) in RAM, controlled by `--cache-ram`. With `--cache-disk PATH`, a disk tier is added below the RAM cache: entries evicted from RAM are written to the given directory, and all RAM entries are flushed there on graceful shutdown. On later requests - including after a server restart - the longest cached prefix of the incoming prompt is restored from disk instead of being re-processed.
+
+```sh
+llama-server -m model.gguf --cache-disk /path/to/cache --cache-disk-limit 32768
+```
+
+Details:
+
+- Files are named `{compat_hash}-{n_tokens}-{chain_hash}.kvc`, where the hashes identify the server configuration and the exact token prefix the file contains. Lookup is a single directory scan at startup plus one hash pass per prompt - no database is used.
+- The cache is invalidated automatically when the model file, mmproj, LoRA adapters, KV cache types, or rope parameters change (stale files are ignored, and deleted once the size budget is exceeded).
+- `--cache-disk-limit` bounds the total size of the directory in MiB; the oldest files (by modification time) are deleted first, including files left over from other models or configurations. The same directory can be shared by multiple servers.
+- By default, files are only written when an entry is evicted from the RAM cache (or on shutdown). With `--cache-disk-write-through`, every prompt saved to the RAM cache is also written to disk immediately, which is more crash-resilient at the cost of extra I/O.
+- Note that KV cache states can be large (potentially multiple GiB per prompt, depending on the model and prompt length), so make sure the disk budget is sized accordingly.
+
 ### Multimodal support

 Multimodal support was added in [#12898](https://github.com/ggml-org/llama.cpp/pull/12898) and is currently an experimental feature.
@@ -337,9 +356,9 @@ It is currently available in the following endpoints:

 For more details, please refer to [multimodal documentation](../../docs/multimodal.md)

-### Server tools support
+### Built-in tools support

-The server includes a set of server tools that enable the LLM to access the local file system directly from the Web UI.
+The server includes a set of built-in tools that enable the LLM to access the local file system directly from the Web UI.

 To use this feature, start the server with `--tools all`. You can also enable only specific tools by passing a comma-separated list: `--tools name1,name2,...`. Run `--help` for the full list of available tool names.

@@ -1631,9 +1650,9 @@ curl http://localhost:8080/v1/messages/count_tokens \
 {"input_tokens": 10}
 ```

-## Server tools
+## Server built-in tools

-The server exposes a REST API under `/tools` that allows the Web UI to call server tools. This endpoint is intended to be used internally by the Web UI and subject to change or to be removed in the future.
+The server exposes a REST API under `/tools` that allows the Web UI to call built-in tools. This endpoint is intended to be used internally by the Web UI and subject to change or to be removed in the future.

 **Please do NOT use this endpoint in a downstream application**

@@ -0,0 +1,580 @@
+#include "server-cache-disk.h"
+
+#include "common.h"
+#include "llama.h"
+
+#include "xxhash/xxhash.h"
+
+#include <algorithm>
+#include <cinttypes>
+#include <cstdio>
+#include <cstring>
+#include <filesystem>
+#include <fstream>
+#include <functional>
+
+namespace {
+
+constexpr uint32_t SERVER_CACHE_DISK_MAGIC   = 0x3143564B; // "KVC1"
+constexpr uint32_t SERVER_CACHE_DISK_VERSION = 1;
+
+// seed for the chained prefix hash - changing it invalidates all filenames
+constexpr uint64_t SERVER_CACHE_DISK_CHAIN_SEED = 0x6b7663636861696eULL;
+
+struct server_cache_disk_file_header {
+    uint32_t magic       = SERVER_CACHE_DISK_MAGIC;
+    uint32_t version     = SERVER_CACHE_DISK_VERSION;
+    uint64_t compat_hash = 0; // full 64-bit value (the filename only carries the low 32 bits)
+    uint64_t chain_hash  = 0;
+    uint32_t n_tokens    = 0;
+    uint32_t pad         = 0;
+    uint64_t tokens_size = 0; // bytes of the server_tokens::serialize() section
+    uint64_t state_size  = 0; // bytes of the llama_state_seq_get_data section
+};
+
+static_assert(sizeof(server_cache_disk_file_header) == 48, "unexpected header size");
+
+std::string make_filename(uint64_t compat_hash, uint32_t n_tokens, uint64_t chain_hash) {
+    char buf[64];
+    snprintf(buf, sizeof(buf), "%08x-%u-%016" PRIx64 ".kvc", (uint32_t) compat_hash, n_tokens, chain_hash);
+    return buf;
+}
+
+bool parse_filename(const std::string & name, uint32_t & compat32, uint32_t & n_tokens, uint64_t & chain_hash) {
+    if (sscanf(name.c_str(), "%8x-%u-%16" SCNx64 ".kvc", &compat32, &n_tokens, &chain_hash) != 3) {
+        return false;
+    }
+
+    // reject padding/case/suffix variations by requiring the canonical spelling
+    return name == make_filename(compat32, n_tokens, chain_hash);
+}
+
+int64_t file_mtime(const std::filesystem::path & path) {
+    std::error_code ec;
+    const auto t = std::filesystem::last_write_time(path, ec);
+    return ec ? 0 : (int64_t) t.time_since_epoch().count();
+}
+
+uint64_t covered_key(uint32_t n_tokens, uint64_t chain_hash) {
+    const uint64_t buf[2] = { n_tokens, chain_hash };
+    return XXH64(buf, sizeof(buf), 0);
+}
+
+// walk the chained hash over the token list, invoking cb(n, h) at every valid prefix boundary:
+// after each text token and after each complete media chunk (never mid-chunk)
+// returns true if the walk reached n_max
+bool tokens_chain_hash_walk(const server_tokens & tokens, size_t n_max, const std::function<bool(size_t, uint64_t)> & cb) {
+    uint64_t h = SERVER_CACHE_DISK_CHAIN_SEED;
+
+    size_t i = 0;
+
+    try {
+        while (i < n_max) {
+            const llama_token tok = tokens[i];
+
+            if (tok == LLAMA_TOKEN_NULL) {
+                // media chunk - fold in its content id instead of the placeholder token ids,
+                // otherwise different images would hash identically
+                const auto & chunk = tokens.find_chunk(i);
+
+                const char * id    = mtmd_input_chunk_get_id(chunk.get());
+                const size_t n_tok = mtmd_input_chunk_get_n_tokens(chunk.get());
+
+                if (id == nullptr || id[0] == '\0' || n_tok == 0 || i + n_tok > n_max) {
+                    return false;
+                }
+
+                std::vector<uint8_t> buf;
+                buf.reserve(5 + strlen(id));
+                buf.push_back(0x01);
+                for (int b = 0; b < 4; ++b) {
+                    buf.push_back((uint8_t) (n_tok >> (8*b)));
+                }
+                buf.insert(buf.end(), id, id + strlen(id));
+
+                h = XXH64(buf.data(), buf.size(), h);
+
+                i += n_tok;
+            } else {
+                uint8_t buf[5] = { 0x00 };
+                memcpy(buf + 1, &tok, sizeof(tok));
+
+                h = XXH64(buf, sizeof(buf), h);
+
+                i += 1;
+            }
+
+            if (!cb(i, h)) {
+                return false;
+            }
+        }
+    } catch (const std::exception & e) {
+        SRV_WRN("failed to hash token list: %s\n", e.what());
+        return false;
+    }
+
+    return true;
+}
+
+} // namespace
+
+server_prompt_cache_disk::server_prompt_cache_disk(const std::string & dir_, uint64_t compat_hash, bool has_mtmd, int32_t limit_mib, bool write_through) :
+        write_through(write_through),
+        dir(dir_.empty() || dir_.back() == DIRECTORY_SEPARATOR ? dir_ : dir_ + DIRECTORY_SEPARATOR),
+        compat_hash(compat_hash),
+        has_mtmd(has_mtmd),
+        limit_bytes(limit_mib < 0 ? 0 : 1024ull*1024ull*limit_mib) {
+    scan_dir();
+}
+
+void server_prompt_cache_disk::scan_dir() {
+    namespace fs = std::filesystem;
+
+    std::error_code ec;
+
+    for (const auto & ent : fs::directory_iterator(dir, ec)) {
+        if (!ent.is_regular_file(ec)) {
+            continue;
+        }
+
+        const std::string name = ent.path().filename().string();
+
+        // leftover temporary files from a previous crash
+        if (name.size() > 4 && name.compare(name.size() - 4, 4, ".tmp") == 0 && name[0] == '.') {
+            fs::remove(ent.path(), ec);
+            continue;
+        }
+
+        uint32_t compat32  = 0;
+        uint32_t n_tokens  = 0;
+        uint64_t chain     = 0;
+
+        if (!parse_filename(name, compat32, n_tokens, chain)) {
+            continue;
+        }
+
+        server_cache_disk_file file;
+        file.name       = name;
+        file.chain_hash = chain;
+        file.n_tokens   = n_tokens;
+        file.n_bytes    = ent.file_size(ec);
+        file.mtime      = file_mtime(ent.path());
+
+        total_bytes += file.n_bytes;
+
+        if (compat32 == (uint32_t) compat_hash) {
+            index[n_tokens][chain] = std::move(file);
+        } else {
+            foreign.push_back(std::move(file));
+        }
+    }
+
+    SRV_INF("disk prompt cache '%s': %zu usable entries, %zu from other configurations, %.3f MiB total (budget: %.3f MiB)\n",
+            dir.c_str(), n_files(), foreign.size(), total_bytes / (1024.0 * 1024.0), limit_bytes / (1024.0 * 1024.0));
+}
+
+size_t server_prompt_cache_disk::n_files() const {
+    size_t res = 0;
+
+    for (const auto & [n, files] : index) {
+        res += files.size();
+    }
+
+    return res;
+}
+
+server_cache_disk_file * server_prompt_cache_disk::find_file(uint32_t n_tokens, uint64_t chain_hash) {
+    const auto it = index.find(n_tokens);
+    if (it == index.end()) {
+        return nullptr;
+    }
+
+    const auto it_file = it->second.find(chain_hash);
+
+    return it_file == it->second.end() ? nullptr : &it_file->second;
+}
+
+const server_cache_disk_file * server_prompt_cache_disk::lookup(const server_tokens & tokens, size_t n_max) const {
+    if (index.empty()) {
+        return nullptr;
+    }
+
+    // no file can be longer than the largest indexed length - cap the walk
+    n_max = std::min<size_t>(n_max, index.rbegin()->first);
+
+    const server_cache_disk_file * best = nullptr;
+
+    tokens_chain_hash_walk(tokens, n_max, [&](size_t n, uint64_t h) {
+        const auto it = index.find((uint32_t) n);
+        if (it != index.end()) {
+            const auto it_file = it->second.find(h);
+            if (it_file != it->second.end()) {
+                best = &it_file->second;
+            }
+        }
+
+        return true;
+    });
+
+    return best;
+}
+
+void server_prompt_cache_disk::touch(const server_cache_disk_file & file) {
+    std::error_code ec;
+    std::filesystem::last_write_time(dir + file.name, std::filesystem::file_time_type::clock::now(), ec);
+
+    if (auto * f = find_file(file.n_tokens, file.chain_hash)) {
+        f->mtime = file_mtime(dir + file.name);
+    }
+}
+
+void server_prompt_cache_disk::forget(const server_cache_disk_file & file) {
+    // copy the fields first - the reference may point into the index entry being erased
+    const uint32_t n_tokens = file.n_tokens;
+    const uint64_t chain    = file.chain_hash;
+    const uint64_t n_bytes  = file.n_bytes;
+
+    const auto it = index.find(n_tokens);
+    if (it == index.end()) {
+        return;
+    }
+
+    if (it->second.erase(chain) > 0) {
+        total_bytes -= std::min<size_t>(total_bytes, n_bytes);
+    }
+
+    if (it->second.empty()) {
+        index.erase(it);
+    }
+}
+
+void server_prompt_cache_disk::remove_file(const server_cache_disk_file & file) {
+    SRV_WRN("disk prompt cache: removing '%s'\n", file.name.c_str());
+
+    std::error_code ec;
+    std::filesystem::remove(dir + file.name, ec);
+
+    forget(file);
+}
+
+void server_prompt_cache_disk::enforce_budget(const std::string & name_protected) {
+    if (limit_bytes == 0) {
+        return;
+    }
+
+    while (total_bytes > limit_bytes) {
+        // find the oldest file, ours and foreign alike
+        const server_cache_disk_file * oldest = nullptr;
+        bool oldest_foreign = false;
+
+        for (const auto & [n, files] : index) {
+            for (const auto & [h, file] : files) {
+                if (file.name != name_protected && (!oldest || file.mtime < oldest->mtime)) {
+                    oldest = &file;
+                    oldest_foreign = false;
+                }
+            }
+        }
+
+        for (const auto & file : foreign) {
+            if (file.name != name_protected && (!oldest || file.mtime < oldest->mtime)) {
+                oldest = &file;
+                oldest_foreign = true;
+            }
+        }
+
+        if (!oldest) {
+            break;
+        }
+
+        SRV_INF("disk prompt cache: size %.3f MiB over budget %.3f MiB, evicting oldest entry '%s'\n",
+                total_bytes / (1024.0 * 1024.0), limit_bytes / (1024.0 * 1024.0), oldest->name.c_str());
+
+        if (oldest_foreign) {
+            std::error_code ec;
+            std::filesystem::remove(dir + oldest->name, ec);
+
+            total_bytes -= std::min<size_t>(total_bytes, oldest->n_bytes);
+
+            foreign.erase(foreign.begin() + (oldest - foreign.data()));
+        } else {
+            remove_file(*oldest);
+        }
+    }
+}
+
+bool server_prompt_cache_disk::store(const server_tokens & tokens, const std::vector<uint8_t> & state_main) {
+    if (tokens.empty() || state_main.empty()) {
+        return false;
+    }
+
+    std::vector<std::pair<size_t, uint64_t>> bounds;
+
+    if (!tokens_chain_hash_walk(tokens, tokens.size(), [&](size_t n, uint64_t h) { bounds.emplace_back(n, h); return true; }) ||
+        bounds.empty() || bounds.back().first != tokens.size()) {
+        SRV_WRN("%s", "disk prompt cache: token list cannot be hashed, skipping\n");
+        return false;
+    }
+
+    const uint32_t n_tokens = (uint32_t) tokens.size();
+    const uint64_t chain    = bounds.back().second;
+
+    if (auto * existing = find_file(n_tokens, chain)) {
+        SRV_TRC("disk prompt cache: '%s' already exists, refreshing\n", existing->name.c_str());
+        touch(*existing);
+        return true;
+    }
+
+    if (covered.count(covered_key(n_tokens, chain)) > 0) {
+        SRV_TRC(" - prompt with %u tokens is a prefix of an already persisted entry, skipping\n", n_tokens);
+        return true;
+    }
+
+    std::vector<char> tok_data;
+    try {
+        tok_data = tokens.serialize();
+    } catch (const std::exception & e) {
+        SRV_WRN("disk prompt cache: failed to serialize tokens: %s\n", e.what());
+        return false;
+    }
+
+    server_cache_disk_file_header header;
+    header.compat_hash = compat_hash;
+    header.chain_hash  = chain;
+    header.n_tokens    = n_tokens;
+    header.tokens_size = tok_data.size();
+    header.state_size  = state_main.size();
+
+    const std::string name = make_filename(compat_hash, n_tokens, chain);
+
+    char tmp_buf[64];
+    snprintf(tmp_buf, sizeof(tmp_buf), ".%08x-%u.tmp", (uint32_t) (uintptr_t) this, tmp_counter++);
+
+    const std::string path_tmp = dir + tmp_buf;
+    const std::string path     = dir + name;
+
+    {
+        std::ofstream out(path_tmp, std::ios::binary | std::ios::trunc);
+
+        out.write((const char *) &header, sizeof(header));
+        out.write(tok_data.data(), tok_data.size());
+        out.write((const char *) state_main.data(), state_main.size());
+
+        if (!out.good()) {
+            SRV_ERR("disk prompt cache: failed to write '%s'\n", path_tmp.c_str());
+
+            out.close();
+
+            std::error_code ec;
+            std::filesystem::remove(path_tmp, ec);
+
+            return false;
+        }
+    }
+
+    std::error_code ec;
+    std::filesystem::rename(path_tmp, path, ec);
+    if (ec) {
+        SRV_ERR("disk prompt cache: failed to rename '%s' to '%s': %s\n", path_tmp.c_str(), path.c_str(), ec.message().c_str());
+
+        std::filesystem::remove(path_tmp, ec);
+
+        return false;
+    }
+
+    server_cache_disk_file file;
+    file.name       = name;
+    file.chain_hash = chain;
+    file.n_tokens   = n_tokens;
+    file.n_bytes    = sizeof(header) + tok_data.size() + state_main.size();
+    file.mtime      = file_mtime(path);
+
+    total_bytes += file.n_bytes;
+
+    index[n_tokens][chain] = std::move(file);
+
+    for (const auto & [n, h] : bounds) {
+        covered.insert(covered_key((uint32_t) n, h));
+    }
+
+    SRV_INF("disk prompt cache: saved prompt with %u tokens, %.3f MiB to '%s'\n",
+            n_tokens, (sizeof(header) + tok_data.size() + state_main.size()) / (1024.0 * 1024.0), name.c_str());
+    SRV_DBG("%s", "__TEST_TAG_CACHE_DISK_STORE__\n");
+
+    enforce_budget(name);
+
+    return true;
+}
+
+server_prompt_cache_disk::load_status server_prompt_cache_disk::load(
+        server_cache_disk_file file, const server_tokens & tokens_new, llama_context * ctx, int32_t id_slot, server_tokens & tokens_out) {
+    const std::string path = dir + file.name;
+
+    std::error_code ec;
+    const uint64_t n_bytes = std::filesystem::file_size(path, ec);
+
+    if (ec) {
+        // deleted by another process - not an error, just a miss
+        forget(file);
+        return LOAD_MISS;
+    }
+
+    std::ifstream in(path, std::ios::binary);
+    if (!in.good()) {
+        forget(file);
+        return LOAD_MISS;
+    }
+
+    server_cache_disk_file_header header;
+    in.read((char *) &header, sizeof(header));
+
+    if (!in.good() ||
+        header.magic    != SERVER_CACHE_DISK_MAGIC   ||
+        header.version  != SERVER_CACHE_DISK_VERSION ||
+        header.chain_hash  != file.chain_hash ||
+        header.n_tokens    != file.n_tokens   ||
+        header.tokens_size % sizeof(llama_token) != 0 ||
+        sizeof(header) + header.tokens_size + header.state_size != n_bytes) {
+        SRV_WRN("disk prompt cache: '%s' is corrupt\n", file.name.c_str());
+        remove_file(file);
+        return LOAD_MISS;
+    }
+
+    if (header.compat_hash != compat_hash) {
+        // same low 32 bits, different configuration - leave the file for its owner
+        SRV_WRN("disk prompt cache: '%s' belongs to a different configuration, ignoring\n", file.name.c_str());
+        forget(file);
+        return LOAD_MISS;
+    }
+
+    llama_tokens packed(header.tokens_size / sizeof(llama_token));
+    in.read((char *) packed.data(), header.tokens_size);
+
+    if (!in.good()) {
+        SRV_WRN("disk prompt cache: '%s' is truncated\n", file.name.c_str());
+        remove_file(file);
+        return LOAD_MISS;
+    }
+
+    server_tokens loaded;
+    try {
+        loaded = server_tokens::deserialize(packed, has_mtmd);
+    } catch (const std::exception & e) {
+        SRV_WRN("disk prompt cache: failed to deserialize tokens from '%s': %s\n", file.name.c_str(), e.what());
+        remove_file(file);
+        return LOAD_MISS;
+    }
+
+    // the filename hash only proves an exact prefix probabilistically - verify against the actual tokens
+    if (loaded.size() != file.n_tokens ||
+        loaded.get_common_prefix(tokens_new) != file.n_tokens ||
+        !loaded.validate(ctx)) {
+        SRV_WRN("disk prompt cache: token mismatch in '%s' (hash collision?)\n", file.name.c_str());
+        remove_file(file);
+        return LOAD_MISS;
+    }
+
+    std::vector<uint8_t> state;
+    try {
+        state.resize(header.state_size);
+    } catch (const std::bad_alloc &) {
+        SRV_ERR("disk prompt cache: failed to allocate %" PRIu64 " bytes for '%s'\n", header.state_size, file.name.c_str());
+        return LOAD_MISS;
+    }
+
+    in.read((char *) state.data(), state.size());
+
+    if (!in.good()) {
+        SRV_WRN("disk prompt cache: '%s' is truncated\n", file.name.c_str());
+        remove_file(file);
+        return LOAD_MISS;
+    }
+
+    const size_t n = llama_state_seq_set_data_ext(ctx, state.data(), state.size(), id_slot, 0);
+    if (n != state.size()) {
+        SRV_WRN("disk prompt cache: failed to restore state from '%s' (%zu / %zu bytes)\n", file.name.c_str(), n, state.size());
+
+        // the sequence may hold a partial state now - clear it and let the caller recover
+        llama_memory_seq_rm(llama_get_memory(ctx), id_slot, -1, -1);
+
+        return LOAD_FAIL_SEQ_DIRTY;
+    }
+
+    tokens_out = std::move(loaded);
+
+    covered.insert(covered_key(file.n_tokens, file.chain_hash));
+
+    touch(file);
+
+    SRV_INF("disk prompt cache: restored prompt with %u tokens, %.3f MiB from '%s'\n",
+            file.n_tokens, state.size() / (1024.0 * 1024.0), file.name.c_str());
+    SRV_DBG("%s", "__TEST_TAG_CACHE_DISK_HIT__\n");
+
+    return LOAD_OK;
+}
+
+//
+// compat hash
+//
+
+namespace {
+
+template <typename T>
+void hash_pod(std::string & blob, const T & value) {
+    static_assert(std::is_trivially_copyable<T>::value, "hash_pod requires a POD type");
+    blob.append((const char *) &value, sizeof(value));
+}
+
+void hash_str(std::string & blob, const std::string & value) {
+    blob += value;
+    blob += '\0';
+}
+
+// path + size + mtime: conservative, but never misses a changed file
+void hash_file_meta(std::string & blob, const std::string & path) {
+    hash_str(blob, path);
+
+    std::error_code ec;
+
+    const uint64_t size = path.empty() ? 0 : (uint64_t) std::filesystem::file_size(path, ec);
+    hash_pod(blob, ec ? (uint64_t) 0 : size);
+
+    hash_pod(blob, path.empty() ? (int64_t) 0 : file_mtime(path));
+}
+
+} // namespace
+
+uint64_t server_cache_disk_compat_hash(const common_params & params) {
+    std::string blob;
+
+    // format versions
+    hash_pod(blob, (uint32_t) SERVER_CACHE_DISK_VERSION);
+    hash_pod(blob, (uint32_t) LLAMA_STATE_SEQ_VERSION);
+    hash_pod(blob, (uint32_t) server_tokens::SERVER_TOKENS_STATE_VERSION);
+
+    // model identity
+    hash_file_meta(blob, params.model.path);
+    hash_file_meta(blob, params.mmproj.path);
+
+    for (const auto & la : params.lora_adapters) {
+        hash_file_meta(blob, la.path);
+        hash_pod(blob, la.scale);
+    }
+
+    // KV cache layout
+    hash_pod(blob, (int32_t) params.cache_type_k);
+    hash_pod(blob, (int32_t) params.cache_type_v);
+    hash_pod(blob, (uint8_t) params.swa_full);
+
+    // rope params change the KV content for the same tokens
+    hash_pod(blob, params.rope_freq_base);
+    hash_pod(blob, params.rope_freq_scale);
+    hash_pod(blob, (int32_t) params.rope_scaling_type);
+    hash_pod(blob, params.yarn_ext_factor);
+    hash_pod(blob, params.yarn_attn_factor);
+    hash_pod(blob, params.yarn_beta_fast);
+    hash_pod(blob, params.yarn_beta_slow);
+    hash_pod(blob, params.yarn_orig_ctx);
+
+    return XXH64(blob.data(), blob.size(), 0);
+}
@@ -0,0 +1,96 @@
+#pragma once
+
+#include "server-common.h"
+
+#include <cstdint>
+#include <map>
+#include <string>
+#include <unordered_map>
+#include <unordered_set>
+#include <vector>
+
+struct common_params;
+struct llama_context;
+
+// disk-backed prompt cache: a cold tier below the in-RAM server_prompt_cache
+//
+// each entry is one file in a flat directory, named after the exact token prefix it contains:
+//
+//     {compat_hash8}-{n_tokens}-{chain_hash16}.kvc
+//
+//   - compat_hash: hash of everything that invalidates a KV state (model file, mmproj, loras,
+//     cache types, rope params, ...) - see server_cache_disk_compat_hash()
+//   - chain_hash: chained hash over the first n_tokens tokens, so a filename identifies an exact
+//     prefix and lookup is a single rolling-hash pass over the incoming prompt plus an index probe
+//
+// file contents mirror what the RAM cache holds for the target context:
+//
+//     header | server_tokens::serialize() bytes | llama_state_seq_get_data (FLAGS_NONE) bytes
+
+struct server_cache_disk_file {
+    std::string name; // filename inside the cache directory
+
+    uint64_t chain_hash = 0;
+    uint32_t n_tokens   = 0;
+    uint64_t n_bytes    = 0;
+    int64_t  mtime      = 0; // only used for relative ordering during eviction
+};
+
+struct server_prompt_cache_disk {
+    server_prompt_cache_disk(const std::string & dir, uint64_t compat_hash, bool has_mtmd, int32_t limit_mib, bool write_through);
+
+    enum load_status {
+        LOAD_OK,             // state restored into the sequence
+        LOAD_MISS,           // file unusable (corrupt, collision, ...) - sequence untouched
+        LOAD_FAIL_SEQ_DIRTY, // restore failed mid-way - the sequence was cleared and must be re-filled
+    };
+
+    // largest exact-prefix hit for the first n_max tokens, or nullptr on miss
+    const server_cache_disk_file * lookup(const server_tokens & tokens, size_t n_max) const;
+
+    // restore the state from a file into sequence id_slot of ctx
+    // on LOAD_OK, tokens_out receives the cached token list (an exact prefix of tokens_new)
+    load_status load(server_cache_disk_file file, const server_tokens & tokens_new, llama_context * ctx, int32_t id_slot, server_tokens & tokens_out);
+
+    // write one entry; deduplicates against existing files and enforces the size budget
+    bool store(const server_tokens & tokens, const std::vector<uint8_t> & state_main);
+
+    size_t n_files() const;
+    size_t n_bytes_total() const { return total_bytes; }
+
+    const bool write_through;
+
+private:
+    void scan_dir();
+
+    server_cache_disk_file * find_file(uint32_t n_tokens, uint64_t chain_hash);
+
+    void touch (const server_cache_disk_file & file);        // bump mtime so eviction treats it as fresh
+    void forget(const server_cache_disk_file & file);        // drop from the index without touching the filesystem
+    void remove_file(const server_cache_disk_file & file);   // delete from disk and drop from the index
+
+    // delete oldest-mtime files (ours and foreign alike) while over the size budget
+    void enforce_budget(const std::string & name_protected);
+
+    const std::string dir;
+    const uint64_t    compat_hash;
+    const bool        has_mtmd;
+    const size_t      limit_bytes; // 0 = no limit
+
+    // n_tokens -> chain_hash -> file, for our compat hash only
+    std::map<uint32_t, std::unordered_map<uint64_t, server_cache_disk_file>> index;
+
+    // .kvc files with a different compat hash prefix - never opened, but counted toward the budget
+    std::vector<server_cache_disk_file> foreign;
+
+    size_t total_bytes = 0; // ours + foreign
+
+    // (n_tokens, chain_hash) prefixes known to be covered by a file written or loaded this
+    // session - lets store() skip prefixes of already-persisted prompts
+    std::unordered_set<uint64_t> covered;
+
+    uint32_t tmp_counter = 0;
+};
+
+// hash of everything that invalidates a saved KV state for the current server configuration
+uint64_t server_cache_disk_compat_hash(const common_params & params);
@@ -266,8 +266,6 @@ static inline raw_buffer base64_decode(const std::string & encoded_string) {

 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");
@@ -156,6 +156,9 @@ private: // disallow accessing these members directly, risking out-of-sync
    // map_idx_to_media will contain: {5, img0}, {8, img1}

 public:
+    // version of the serialize()/deserialize() format below
+    static constexpr uint32_t SERVER_TOKENS_STATE_VERSION = 1;
+
    server_tokens() = default;
    ~server_tokens() = default;

@@ -275,11 +275,13 @@ struct server_slot {
            llama_state_seq_get_data_ext(ctx_dft, cur->data.drft.data(), cur_size_dft, id, LLAMA_STATE_SEQ_FLAGS_NONE);
        }

+        prompt_cache.disk_store_write_through(*cur);
+
        return true;
    }

    bool prompt_load(server_prompt_cache & prompt_cache, const server_tokens & tokens) {
-        bool res = prompt_cache.load(prompt, tokens, ctx_tgt, ctx_dft, id);
+        bool res = prompt_cache.load(prompt, tokens, ctx_tgt, ctx_dft, id, n_ctx);
        if (!res) {
            SLT_WRN(*this, "%s", "failed to load prompt from cache\n");
        }
@@ -1308,7 +1310,26 @@ private:
            SRV_TRC("%s", "use `--cache-ram 0` to disable the prompt cache\n");

            prompt_cache = std::make_unique<server_prompt_cache>(params_base.cache_ram_mib, n_ctx);
+
+            if (!params_base.cache_disk_path.empty()) {
+                const uint64_t compat_hash = server_cache_disk_compat_hash(params_base);
+
+                SRV_INF("disk prompt cache is enabled, dir: '%s', compat hash: %08x\n",
+                        params_base.cache_disk_path.c_str(), (uint32_t) compat_hash);
+
+                prompt_cache->disk = std::make_unique<server_prompt_cache_disk>(
+                        params_base.cache_disk_path,
+                        compat_hash,
+                        mctx != nullptr,
+                        params_base.cache_disk_limit_mib,
+                        params_base.cache_disk_write_through);
+            }
        } else {
+            if (!params_base.cache_disk_path.empty()) {
+                SRV_ERR("%s", "--cache-disk requires the RAM prompt cache - remove `--cache-ram 0`\n");
+                return false;
+            }
+
            SRV_TRC("%s", "prompt cache is disabled - use `--cache-ram N` to enable it\n");
        }
        SRV_TRC("%s", "for more info see https://github.com/ggml-org/llama.cpp/pull/16391\n");
@@ -4058,6 +4079,11 @@ bool server_context::load_model(common_params & params) {
 void server_context::start_loop() {
    auto & params = impl->params_base;
    impl->queue_tasks.start_loop(params.sleep_idle_seconds * 1000);
+
+    // on graceful shutdown, give the RAM prompt cache entries a chance to survive the restart
+    if (impl->prompt_cache) {
+        impl->prompt_cache->disk_flush();
+    }
 }

 void server_context::terminate() {
@@ -1025,13 +1025,10 @@ void server_models::load(const std::string & name, const load_options & opts) {
            char * buffer = vec_buf.data();
            if (stdout_file) {
                while (fgets(buffer, vec_buf.size(), stdout_file) != nullptr) {
+                    LOG("[%5d] %s", port, buffer);
                    std::string str(buffer);
                    if (string_starts_with(buffer, CMD_CHILD_TO_ROUTER_STATE)) {
-                        LOG_DBG("[%5d] %s", port, buffer); // prevent spamming the log
                        this->handle_child_state(name, str);
-                    } else {
-                        // forward log
-                        LOG("[%5d] %s", port, buffer);
                    }
                }
            } else {
@@ -1750,6 +1750,8 @@ server_prompt_cache_state * server_prompt_cache::alloc(const server_prompt & pro
            SRV_WRN(" - making room for prompt cache entry, removing oldest entry (size = %.3f MiB)\n",
                    states.front().size() / (1024.0 * 1024.0));

+            spill_front();
+
            states.pop_front();
        }
    }
@@ -1787,7 +1789,7 @@ server_prompt_cache_state * server_prompt_cache::alloc(const server_prompt & pro
    return &states.back();
 }

-bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tokens_new, llama_context * ctx_tgt, llama_context * ctx_dft, int32_t id_slot) {
+bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tokens_new, llama_context * ctx_tgt, llama_context * ctx_dft, int32_t id_slot, int32_t n_ctx_slot) {
    const int lcp_best = prompt.tokens.get_common_prefix(tokens_new);

    float f_keep_best = prompt.tokens.size() > 0 ? float(lcp_best) / prompt.tokens.size() : -1.0f; // empty slot: any cache entry wins
@@ -1797,6 +1799,8 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok

    auto it_best = states.end();

+    int lcp_it_best = 0;
+
    // find the most similar cached prompt, that would also preserve the most context
    for (auto it = states.begin(); it != states.end(); ++it) {
        const int lcp_cur = it->prompt.tokens.get_common_prefix(tokens_new);
@@ -1815,7 +1819,41 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
            f_keep_best = f_keep_cur;
            f_sim_best  = f_sim_cur;

-            it_best = it;
+            it_best     = it;
+            lcp_it_best = lcp_cur;
+        }
+    }
+
+    // check the disk tier for an exact-prefix match longer than what RAM (or the slot itself) offers
+    if (disk) {
+        const int lcp_sel = std::max(lcp_best, lcp_it_best);
+
+        const size_t n_max = std::min<size_t>(tokens_new.size(), std::max(0, n_ctx_slot));
+
+        const auto * file = disk->lookup(tokens_new, n_max);
+
+        if (file && (int64_t) file->n_tokens > (int64_t) lcp_sel) {
+            server_tokens tokens_disk;
+
+            const auto status = disk->load(*file, tokens_new, ctx_tgt, id_slot, tokens_disk);
+
+            if (status == server_prompt_cache_disk::LOAD_OK) {
+                // disk entries carry no draft state - clear the draft sequence so it re-prefills
+                if (ctx_dft) {
+                    llama_memory_seq_rm(llama_get_memory(ctx_dft), id_slot, -1, -1);
+                }
+
+                prompt.tokens = std::move(tokens_disk);
+                prompt.checkpoints.clear();
+
+                return true;
+            }
+
+            if (status == server_prompt_cache_disk::LOAD_FAIL_SEQ_DIRTY && it_best == states.end()) {
+                // the slot's sequence was cleared during the failed restore and there is no RAM
+                // candidate to restore over it - the caller has to clear the slot
+                return false;
+            }
        }
    }

@@ -1869,6 +1907,8 @@ void server_prompt_cache::update() {
        while (!states.empty() && size() > limit_size) {
            SRV_WRN(" - cache size limit reached, removing oldest entry (size = %.3f MiB)\n", states.front().size() / (1024.0 * 1024.0));

+            spill_front();
+
            states.pop_front();
        }
    }
@@ -1884,6 +1924,8 @@ void server_prompt_cache::update() {
            SRV_WRN(" - cache token limit (%zu, est: %zu) reached, removing oldest entry (size = %.3f MiB)\n",
                    limit_tokens, limit_tokens_cur, states.front().size() / (1024.0 * 1024.0));

+            spill_front();
+
            states.pop_front();
        }
    }
@@ -1896,3 +1938,39 @@ void server_prompt_cache::update() {
                (const void *)&state, state.prompt.n_tokens(), state.prompt.checkpoints.size(), state.size() / (1024.0 * 1024.0));
    }
 }
+
+void server_prompt_cache::disk_store(const server_prompt_cache_state & state) const {
+    if (!disk || state.data.main.empty()) {
+        return;
+    }
+
+    disk->store(state.prompt.tokens, state.data.main);
+}
+
+void server_prompt_cache::disk_store_write_through(const server_prompt_cache_state & state) const {
+    if (!disk || !disk->write_through) {
+        return;
+    }
+
+    disk_store(state);
+}
+
+void server_prompt_cache::disk_flush() const {
+    if (!disk) {
+        return;
+    }
+
+    SRV_INF("flushing %zu prompt cache entries to disk\n", states.size());
+
+    for (const auto & state : states) {
+        disk_store(state);
+    }
+}
+
+void server_prompt_cache::spill_front() const {
+    if (!disk || states.empty()) {
+        return;
+    }
+
+    disk_store(states.front());
+}
@@ -7,8 +7,10 @@
 #include <unordered_set>
 #include <list>
 #include <map>
+#include <memory>

 // TODO: prevent including the whole server-common.h as we only use server_tokens
+#include "server-cache-disk.h"
 #include "server-common.h"

 using json = nlohmann::ordered_json;
@@ -612,6 +614,10 @@ struct server_prompt_cache {

    std::list<server_prompt_cache_state> states;

+    // optional cold tier - entries evicted from RAM are spilled here and can be restored later,
+    // including across server restarts
+    std::unique_ptr<server_prompt_cache_disk> disk;
+
    // in bytes, 0 = no limit
    size_t limit_size = 0;

@@ -624,9 +630,22 @@ struct server_prompt_cache {

    server_prompt_cache_state * alloc(const server_prompt & prompt, size_t state_size_main, size_t state_size_drft);

-    bool load(server_prompt & prompt, const server_tokens & tokens_new, llama_context * ctx_tgt, llama_context * ctx_dft, int32_t id_slot);
+    bool load(server_prompt & prompt, const server_tokens & tokens_new, llama_context * ctx_tgt, llama_context * ctx_dft, int32_t id_slot, int32_t n_ctx_slot);

    void update();
+
+    // write one RAM cache entry to the disk tier (no-op when the disk tier is disabled)
+    void disk_store(const server_prompt_cache_state & state) const;
+
+    // disk_store, but only when write-through mode is enabled
+    void disk_store_write_through(const server_prompt_cache_state & state) const;
+
+    // spill all RAM entries to the disk tier (e.g. on graceful shutdown)
+    void disk_flush() const;
+
+private:
+    // spill the entry that is about to be evicted
+    void spill_front() const;
 };

 // used exclusively by router mode
@@ -2035,7 +2035,7 @@ void server_tools::setup(const std::vector<std::string> & enabled_tools,
        }
    }

-    // append MCP tools, skipping any that collide with a server tool or another MCP tool of the same "<server>_<tool>" name
+    // append MCP tools, skipping any that collide with a built-in or another MCP tool of the same "<server>_<tool>" name
    if (!mcp_mgr.empty()) {
        std::unordered_set<std::string> seen_names;
        for (auto & t : tools) {
@@ -18,7 +18,7 @@ struct server_tool {

    virtual ~server_tool() = default;
    virtual json get_definition() const = 0;
-    virtual std::string type() const { return "server"; }
+    virtual std::string type() const { return "builtin"; }

    struct stream {
        server_response & qr;
@@ -346,7 +346,7 @@ int llama_server(common_params & params, int argc, char ** argv) {
        ctx_http.get ("/tools",           ex_wrapper(tools.handle_get));
        ctx_http.post("/tools",           ex_wrapper(tools.handle_post));
        if (!params.server_tools.empty()) {
-            warn_names.push_back("server tools (experimental)");
+            warn_names.push_back("built-in tools (experimental)");
        }
        if (!params.server_tools_runtime.empty()) {
            warn_names.push_back("tools runtime (experimental)");
@@ -0,0 +1,268 @@
+import base64
+import glob
+import os
+import shutil
+import tempfile
+import time
+
+import pytest
+import requests
+from utils import *
+
+server = ServerPreset.tinyllama2()
+
+cache_dir: str = ""
+
+
+class LogReader:
+    def __init__(self, path):
+        self.path = path
+        self.pos = 0
+    def drain(self):
+        with open(self.path) as f:
+            f.seek(self.pos)
+            content = f.read()
+            self.pos = f.tell()
+        return content
+    def wait_for(self, tag, timeout=10) -> bool:
+        # the server log is pumped to the file asynchronously - poll for the tag
+        deadline = time.time() + timeout
+        while time.time() < deadline:
+            if tag in self.drain():
+                return True
+            time.sleep(0.25)
+        return False
+
+
+def kvc_files() -> list[str]:
+    return sorted(glob.glob(os.path.join(cache_dir, "*.kvc")))
+
+
+@pytest.fixture(autouse=True)
+def create_server():
+    global server, cache_dir
+    cache_dir = tempfile.mkdtemp(prefix="llama_cache_disk_")
+    server = ServerPreset.tinyllama2()
+    server.n_slots = 1
+    server.temperature = 0.0
+    server.debug = True
+    server.cache_disk = cache_dir
+    fd, server.log_path = tempfile.mkstemp(suffix='.log')
+    os.close(fd)
+    yield
+    shutil.rmtree(cache_dir, ignore_errors=True)
+
+
+PROMPT_A = (
+    "Once upon a time in a land far away, there lived a brave knight "
+    "who traveled across mountains and rivers to find the legendary "
+    "golden sword hidden deep within the enchanted forest of whispers."
+)
+
+PROMPT_B = "The quick brown fox jumps over the lazy dog."
+
+
+def make_prompt_request(prompt, n_predict=0):
+    global server
+    res = server.make_request("POST", "/completion", data={
+        "prompt": prompt,
+        "n_predict": n_predict,  # 0 = evaluate the prompt into the KV cache only
+        "cache_prompt": True,
+    })
+    assert res.status_code == 200
+    return res
+
+
+def test_write_through_and_restart_hit():
+    global server
+    server.cache_disk_write_through = True
+    server.start()
+    log = LogReader(server.log_path)
+
+    res = make_prompt_request(PROMPT_A)
+    prompt_n_full = res.body["timings"]["prompt_n"]
+    assert prompt_n_full > 0
+
+    # nothing is written while the prompt is still live in the slot
+    assert len(kvc_files()) == 0
+
+    # a different prompt takes over the only slot - the previous one is saved
+    # to the RAM cache and, in write-through mode, to disk immediately
+    make_prompt_request(PROMPT_B)
+    assert log.wait_for("__TEST_TAG_CACHE_DISK_STORE__")
+    assert len(kvc_files()) == 1
+
+    # the state must survive a full server restart
+    server.stop()
+    server.start()
+    log = LogReader(server.log_path)
+
+    res = make_prompt_request(PROMPT_A)
+    assert log.wait_for("__TEST_TAG_CACHE_DISK_HIT__")
+    assert res.body["timings"]["prompt_n"] == 1  # only the last token is re-evaluated
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+
+
+def test_spill_on_shutdown_flush():
+    global server
+    server.start()
+    log = LogReader(server.log_path)
+
+    make_prompt_request(PROMPT_A)
+    make_prompt_request(PROMPT_B)  # forces PROMPT_A into the RAM cache
+
+    # without write-through, nothing reaches the disk while running
+    time.sleep(0.5)
+    assert "__TEST_TAG_CACHE_DISK_STORE__" not in log.drain()
+    assert len(kvc_files()) == 0
+
+    # a graceful shutdown flushes the RAM cache entries to disk
+    server.stop()
+    assert len(kvc_files()) == 1
+
+    server.start()
+    log = LogReader(server.log_path)
+
+    res = make_prompt_request(PROMPT_A)
+    assert log.wait_for("__TEST_TAG_CACHE_DISK_HIT__")
+    assert res.body["timings"]["prompt_n"] == 1
+
+
+def test_ram_cache_hit_takes_priority():
+    global server
+    server.cache_disk_write_through = True
+    server.start()
+    log = LogReader(server.log_path)
+
+    make_prompt_request(PROMPT_A)
+    make_prompt_request(PROMPT_B)
+    assert len(kvc_files()) == 1
+
+    # PROMPT_A is in both the RAM cache and on disk - the RAM copy must win
+    # (the disk entry is never longer than the RAM one here)
+    res = make_prompt_request(PROMPT_A)
+    time.sleep(0.5)
+    assert "__TEST_TAG_CACHE_DISK_HIT__" not in log.drain()
+    assert res.body["timings"]["cache_n"] > 0
+
+
+def test_budget_eviction():
+    global server
+    server.n_ctx = 2048
+    server.n_batch = 512
+    server.cache_disk_write_through = True
+    server.cache_disk_limit = 1  # MiB
+    server.start()
+
+    # three long, distinct token-array prompts; each state is close to 1 MiB
+    n_len = 1500
+    for i in range(3):
+        make_prompt_request([100 + i] * n_len)
+
+    # one final small prompt to force the last long prompt out of the slot
+    make_prompt_request(PROMPT_B)
+
+    files = kvc_files()
+    assert len(files) >= 1
+    assert len(files) < 3  # the oldest entries were evicted
+
+    # the budget is respected (a single over-budget file is allowed to remain)
+    if len(files) > 1:
+        assert sum(os.path.getsize(f) for f in files) <= 1024 * 1024
+
+
+def test_corrupt_file_is_removed():
+    global server
+    server.cache_disk_write_through = True
+    server.start()
+
+    make_prompt_request(PROMPT_A)
+    make_prompt_request(PROMPT_B)
+    files = kvc_files()
+    assert len(files) == 1
+
+    server.stop()
+
+    # corrupt the serialized token section (starts right after the 48-byte header)
+    with open(files[0], "r+b") as f:
+        f.seek(48 + 4)
+        f.write(b"\xff\xff\xff\xff")
+
+    server.start()
+    log = LogReader(server.log_path)
+
+    # the request must still succeed, with the prompt fully re-processed
+    res = make_prompt_request(PROMPT_A)
+    time.sleep(0.5)
+    assert "__TEST_TAG_CACHE_DISK_HIT__" not in log.drain()
+    assert res.body["timings"]["prompt_n"] > 1
+
+    # the corrupt file was deleted
+    assert len(kvc_files()) == 0
+
+
+IMG_URL_CAT = "https://huggingface.co/ggml-org/tinygemma3-GGUF/resolve/main/test/91_cat.png"
+
+
+def _get_img_base64(url: str) -> str:
+    response = requests.get(url)
+    response.raise_for_status()
+    return base64.b64encode(response.content).decode("utf-8")
+
+
+@pytest.fixture
+def mmproj_server():
+    global cache_dir
+    os.environ['LLAMA_MEDIA_MARKER'] = '<__media__>'
+    mm_server = ServerPreset.tinygemma3()
+    mm_server.n_slots = 1
+    mm_server.temperature = 0.0
+    mm_server.debug = True
+    # use the full SWA cache so the restored image prefix can be reused
+    mm_server.swa_full = True
+    mm_server.cache_disk = cache_dir
+    mm_server.cache_disk_write_through = True
+    fd, mm_server.log_path = tempfile.mkstemp(suffix='.log')
+    os.close(fd)
+    return mm_server
+
+
+def test_image_prompt_across_restart(mmproj_server):
+    server = mmproj_server
+    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={
+        "n_predict": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    prompt_n_full = res.body["timings"]["prompt_n"]
+
+    res = server.make_request("POST", "/completions", data={
+        "n_predict": 0,
+        "cache_prompt": True,
+        "prompt": "The quick brown fox",
+    })
+    assert res.status_code == 200
+    assert len(kvc_files()) == 1
+
+    server.stop()
+    server.start()
+    log = LogReader(server.log_path)
+
+    # the image KV must be restored from disk in the new process
+    res = server.make_request("POST", "/completions", data={
+        "n_predict": 0,
+        "cache_prompt": True,
+        "prompt": prompt_cat,
+    })
+    assert res.status_code == 200
+    assert log.wait_for("__TEST_TAG_CACHE_DISK_HIT__")
+    assert res.body["timings"]["prompt_n"] == 1
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
@@ -111,6 +111,9 @@ class ServerProcess:
    media_path: str | None = None
    sleep_idle_seconds: int | None = None
    cache_ram: int | None = None
+    cache_disk: str | None = None
+    cache_disk_limit: int | None = None
+    cache_disk_write_through: bool = False
    no_cache_idle_slots: bool = False
    log_path: str | None = None
    ui_mcp_proxy: bool = False
@@ -271,6 +274,12 @@ class ServerProcess:
            server_args.extend(["--sleep-idle-seconds", self.sleep_idle_seconds])
        if self.cache_ram is not None:
            server_args.extend(["--cache-ram", self.cache_ram])
+        if self.cache_disk is not None:
+            server_args.extend(["--cache-disk", self.cache_disk])
+        if self.cache_disk_limit is not None:
+            server_args.extend(["--cache-disk-limit", self.cache_disk_limit])
+        if self.cache_disk_write_through:
+            server_args.append("--cache-disk-write-through")
        if self.no_cache_idle_slots:
            server_args.append("--no-cache-idle-slots")
        if self.ui_mcp_proxy:
@@ -61,9 +61,6 @@ export default ts.config(
 				{ blankLine: 'always', next: ['return', 'throw', 'break', 'continue'], prev: '*' }
 			],

-			// Alphabetical order for enum members
-			'perfectionist/sort-enums': ['error', { type: 'natural' }],
-
 			'perfectionist/sort-objects': ['error', { type: 'natural' }],

 			// Alphabetical order for variable declarations and object keys
@@ -48,7 +48,6 @@
 		containsFileMentionLink,
 		findCommandToken,
 		findMentionToken,
-		getConversationModel,
 		isIMEComposing,
 		isOffsetInCodeBlock,
 		parseClipboardContent,
@@ -191,9 +190,31 @@

 	let isRouter = $derived(serverStore.isRouterMode);
 	let conversationModel = $derived(
-		getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
+		chatStore.getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
 	);
-	let activeModelId = $derived(modelsStore.activeModelId);
+	let activeModelId = $derived.by(() => {
+		const options = modelsStore.models;
+
+		if (!isRouter) {
+			return options.length > 0 ? options[0].model : null;
+		}
+
+		const selectedId = modelsStore.selectedModelId;
+
+		if (selectedId) {
+			const model = options.find((m) => m.id === selectedId);
+
+			if (model) return model.model;
+		}
+
+		if (conversationModel) {
+			const model = options.find((m) => m.model === conversationModel);
+
+			if (model) return model.model;
+		}
+
+		return null;
+	});

 	let hasModelSelected = $derived(
 		!isRouter || !!conversationModel || !!modelsStore.selectedModelId
@@ -35,7 +35,7 @@
 						<span>
 							Run llama-server with <code>{CLI_FLAGS.TOOLS}</code> flag to enable

-							<strong>Server Tools</strong>.
+							<strong>Built-in Tools</strong>.
 						</span>
 					</span>

@@ -2,10 +2,10 @@
 	import ChatFormActionAddButton from './ChatFormActionAddButton.svelte';
 	import ChatFormActionAddDropdown from './ChatFormActionAddDropdown.svelte';
 	import ChatFormActionAddSheet from './ChatFormActionAddSheet.svelte';
-	import { deviceStore } from '$lib/stores';
+	import { isMobile } from '$lib/stores';
 </script>

-{#if deviceStore.isMobile}
+{#if isMobile.current}
 	<ChatFormActionAddSheet>
 		{#snippet trigger({ disabled, onclick })}
 			<ChatFormActionAddButton {disabled} {onclick} />
@@ -1,7 +1,6 @@
 <script lang="ts">
 	import { ModelsSelectorDropdown, ModelsSelectorSheet } from '$lib/components/app';
-	import { conversationsStore, deviceStore, modelsStore, serverStore } from '$lib/stores';
-	import { getConversationModel } from '$lib/utils';
+	import { chatStore, conversationsStore, isMobile, modelsStore, serverStore } from '$lib/stores';

 	interface Props {
 		disabled?: boolean;
@@ -31,7 +30,7 @@
 	let isOffline = $derived(!!serverStore.error);

 	let conversationModel = $derived(
-		getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
+		chatStore.getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
 	);

 	let lastSyncedConversationModel: string | null = null;
@@ -75,7 +74,29 @@
 		}
 	});

-	let activeModelId = $derived(modelsStore.activeModelId);
+	let activeModelId = $derived.by(() => {
+		const options = modelsStore.models;
+
+		if (!isRouter) {
+			return options.length > 0 ? options[0].model : null;
+		}
+
+		const selectedId = modelsStore.selectedModelId;
+
+		if (selectedId) {
+			const model = options.find((m) => m.id === selectedId);
+
+			if (model) return model.model;
+		}
+
+		if (conversationModel) {
+			const model = options.find((m) => m.model === conversationModel);
+
+			if (model) return model.model;
+		}
+
+		return null;
+	});

 	let modelPropsVersion = $state(0); // Used to trigger reactivity after fetch

@@ -149,7 +170,7 @@
 	}
 </script>

-{#if deviceStore.isMobile}
+{#if isMobile.current}
 	<ModelsSelectorSheet
 		disabled={disabled || isOffline}
 		bind:this={selectorModelRef}
@@ -1,14 +1,15 @@
 <script lang="ts">
 	import ContextGaugeDial from './ContextGaugeDial.svelte';
+	import { useContextGauge } from '$lib/hooks/use-context-gauge.svelte';
 	import {
+		chatStore,
+		conversationsStore,
 		gaugeTriggerClick,
 		gaugeTriggerEnter,
 		gaugeTriggerKeydown,
 		gaugeTriggerLeave,
 		gaugeTriggerPointerDown
-	} from './gauge-popup.svelte';
-	import { useContextGauge } from '$lib/hooks/use-context-gauge.svelte';
-	import { chatStore, conversationsStore } from '$lib/stores';
+	} from '$lib/stores';
 	import { untrack } from 'svelte';

 	const gauge = useContextGauge();
@@ -1,9 +1,9 @@
 <script lang="ts">
 	import ContextGaugeDetailRow from './ContextGaugeDetailRow.svelte';
-	import { gaugePopup } from './gauge-popup.svelte';
 	import { ChevronDown } from '@lucide/svelte';
 	import * as Collapsible from '$lib/components/ui/collapsible';
 	import { STATS_UNITS } from '$lib/constants';
+	import { gaugePopup } from '$lib/stores/context-gauge-popup.svelte';

 	interface Props {
 		currentRead: number;
@@ -2,13 +2,8 @@
 	import { colorLevelBgClass, colorLevelTextClass } from './context-gauge';
 	import ContextGaugeDetails from './ContextGaugeDetails.svelte';
 	import ContextGaugeLoadModel from './ContextGaugeLoadModel.svelte';
-	import {
-		gaugeCardEnter,
-		gaugeCardLeave,
-		gaugePopup,
-		gaugePopupClose
-	} from './gauge-popup.svelte';
 	import { useContextGauge } from '$lib/hooks/use-context-gauge.svelte';
+	import { gaugeCardEnter, gaugeCardLeave, gaugePopup, gaugePopupClose } from '$lib/stores';
 	import { formatParameters } from '$lib/utils/formatters';

 	const gauge = useContextGauge();
@@ -62,7 +62,7 @@
 	// it, the picker still opens for manual entry but explains why search is
 	// unavailable instead of firing searches that would only fail. Browse is
 	// hidden too: it resolves the picked folder name through the same tool.
-	const fileSearchKey = $derived(toolsStore.getPermissionKey(BuiltInTool.SERVER_FILE_GLOB_SEARCH));
+	const fileSearchKey = $derived(toolsStore.getPermissionKey(BuiltInTool.FILE_GLOB_SEARCH));
 	const fileSearchEnabled = $derived(
 		fileSearchKey !== null && toolsStore.isToolEnabled(fileSearchKey)
 	);
@@ -212,7 +212,7 @@
 	// so the caller fails visibly instead of committing a bare leaf name.
 	async function resolveNativeName(name: string): Promise<string | null> {
 		try {
-			const res = await ToolsService.executeToolRaw(BuiltInTool.SERVER_FILE_GLOB_SEARCH, {
+			const res = await ToolsService.executeToolRaw(BuiltInTool.FILE_GLOB_SEARCH, {
 				include: buildCaseInsensitiveGlob(name),
 				limit: SEARCH.NATIVE_LIMIT,
 				max_depth: SEARCH.NATIVE_MAX_DEPTH,
@@ -1,5 +1,5 @@
 <script lang="ts">
-	import { deviceStore } from '$lib/stores';
+	import { isMobile } from '$lib/stores';
 	import { autoResizeTextarea } from '$lib/utils';
 	import { onMount } from 'svelte';

@@ -37,7 +37,7 @@
 	}

 	export function focus() {
-		if (deviceStore.isMobile) return;
+		if (isMobile.current) return;

 		textareaElement?.focus({ preventScroll: true });
 	}
@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { CODE_BLOCK, CODE_TOKEN_ATTR, UI_DATA_ATTRS } from '$lib/constants';
 	import { BooleanString, ChatFormInputRichTokenKind, ColorMode } from '$lib/enums';
-	import { deviceStore } from '$lib/stores';
+	import { isMobile } from '$lib/stores';
 	import type { ChatFormInputRichToken } from '$lib/types';
 	import type { SourceHistoryEntry } from '$lib/utils';
 	import {
@@ -750,7 +750,7 @@
 		syncEmptyState();
 		document.addEventListener('selectionchange', handleSelectionChange);

-		if (!deviceStore.isMobile) {
+		if (!isMobile.current) {
 			rootElement?.focus({ preventScroll: true });
 		}
 	});
@@ -792,7 +792,7 @@
 	}

 	export function focus() {
-		if (deviceStore.isMobile) return;
+		if (isMobile.current) return;

 		rootElement?.focus({ preventScroll: true });
 	}
@@ -8,7 +8,7 @@
 	import { BuiltInTool, FileMentionEntryType, GlobSearchType, KeyboardKey } from '$lib/enums';
 	import { useDebouncedSearch } from '$lib/hooks/use-debounced-search.svelte';
 	import { usePickerNavigation } from '$lib/hooks/use-picker-navigation.svelte';
-	import { deviceStore, settingsStore, toolsStore } from '$lib/stores';
+	import { isMobile, settingsStore, toolsStore } from '$lib/stores';
 	import type { FileMentionEntry, GlobEntryResult } from '$lib/types';
 	import { abbreviateHome, runGlobSearchWithChildren } from '$lib/utils';

@@ -51,7 +51,7 @@
 	// When the server does not expose file_glob_search (started without
 	// --tools) or the user disabled it, the picker still opens but explains
 	// why instead of firing searches that would only fail.
-	const fileSearchKey = $derived(toolsStore.getPermissionKey(BuiltInTool.SERVER_FILE_GLOB_SEARCH));
+	const fileSearchKey = $derived(toolsStore.getPermissionKey(BuiltInTool.FILE_GLOB_SEARCH));
 	const fileSearchEnabled = $derived(
 		fileSearchKey !== null && toolsStore.isToolEnabled(fileSearchKey)
 	);
@@ -130,7 +130,7 @@
 		return searchError ? `Search failed - ${searchError}` : 'No matching files or folders';
 	});

-	const showTooltip = $derived(!deviceStore.isMobile);
+	const showTooltip = $derived(!isMobile.current);

 	$effect(() => {
 		if (typeof window === 'undefined') return;
@@ -11,7 +11,7 @@
 	import { setChatMessageActionsContext, setChatMessageEditContext } from '$lib/contexts';
 	import { AgenticSectionType, AttachmentType, MessageRole } from '$lib/enums';
 	import { DatabaseService } from '$lib/services/database.service';
-	import { chatStore, conversationsStore, deviceStore } from '$lib/stores';
+	import { chatStore, conversationsStore, isMobile } from '$lib/stores';
 	import type {
 		ChatMessageActions,
 		ChatMessageDeletionInfo,
@@ -304,7 +304,7 @@

 	// After the system message flow ends, hand focus to the main chat form
 	function focusMainChatForm() {
-		if (deviceStore.isMobile) return;
+		if (isMobile.current) return;

 		document.querySelector<HTMLTextAreaElement>('.chat-screen-form-wrapper textarea')?.focus();
 	}
@@ -35,19 +35,19 @@

 {#if isSearchCall}
 	<ChatMessageToolCallBlockSearchResults {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.BROWSER_GET_DATETIME}
+{:else if section.toolName === BuiltInTool.GET_DATETIME}
 	<ChatMessageToolCallBlockGetDatetime {section} {isStreaming} />
-{:else if section.toolName === BuiltInTool.SERVER_GET_INFO}
+{:else if section.toolName === BuiltInTool.GET_INFO}
 	<ChatMessageToolCallBlockGetInfo {section} {isStreaming} />
-{:else if section.toolName === BuiltInTool.SERVER_READ_FILE}
+{:else if section.toolName === BuiltInTool.READ_FILE}
 	<ChatMessageToolCallBlockReadFile {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.BROWSER_READ_MEDIA}
+{:else if section.toolName === BuiltInTool.READ_MEDIA}
 	<ChatMessageToolCallBlockReadMedia {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.SERVER_EDIT_FILE}
+{:else if section.toolName === BuiltInTool.EDIT_FILE}
 	<ChatMessageToolCallBlockEditFile {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.SERVER_WRITE_FILE}
+{:else if section.toolName === BuiltInTool.WRITE_FILE}
 	<ChatMessageToolCallBlockWriteFile {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.SERVER_EXEC_SHELL_COMMAND}
+{:else if section.toolName === BuiltInTool.EXEC_SHELL_COMMAND}
 	<ChatMessageToolCallBlockExecShellCommand
 		{section}
 		{open}
@@ -56,11 +56,11 @@
 		{attachments}
 		{onToggle}
 	/>
-{:else if section.toolName === BuiltInTool.SERVER_FILE_GLOB_SEARCH}
+{:else if section.toolName === BuiltInTool.FILE_GLOB_SEARCH}
 	<ChatMessageToolCallBlockFileGlobSearch {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.SERVER_GREP_SEARCH}
+{:else if section.toolName === BuiltInTool.GREP_SEARCH}
 	<ChatMessageToolCallBlockGrepSearch {section} {open} {isStreaming} {onToggle} />
-{:else if section.toolName === BuiltInTool.BROWSER_RUN_JAVASCRIPT}
+{:else if section.toolName === BuiltInTool.RUN_JAVASCRIPT}
 	<ChatMessageToolCallBlockRunJavascript {section} {open} {isStreaming} {onToggle} />
 {:else}
 	<ChatMessageToolCallBlockDefault {section} {open} {isStreaming} {attachments} {onToggle} />
@@ -12,7 +12,7 @@
 	import {
 		classifyToolResult,
 		formatJsonPretty,
-		getToolUi,
+		getBuiltinToolUi,
 		parseToolResultWithMedia
 	} from '$lib/utils';
 	import { createBase64DataUrl } from '$lib/utils/data-url';
@@ -27,7 +27,7 @@

 	let { attachments, isStreaming, onToggle, open, section }: Props = $props();

-	const title = $derived(getToolUi(section.toolName)?.label ?? section.toolName ?? '');
+	const title = $derived(getBuiltinToolUi(section.toolName)?.label ?? section.toolName ?? '');
 	const outputKind = $derived(classifyToolResult(section.toolResult));
 	const parsedLines: ToolResultLine[] = $derived(
 		section.toolResult ? parseToolResultWithMedia(section.toolResult, attachments) : []
@@ -6,7 +6,7 @@
 	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
 	import { FileTypeText } from '$lib/enums';
 	import type { AgenticSection } from '$lib/types';
-	import { getToolUi } from '$lib/utils';
+	import { getBuiltinToolUi } from '$lib/utils';

 	interface Props {
 		section: AgenticSection;
@@ -18,7 +18,7 @@
 	let { isStreaming, onToggle, open, section }: Props = $props();

 	const runJsMeta = $derived(parseRunJavascriptMeta(section));
-	const title = $derived(getToolUi(section.toolName)?.label ?? section.toolName ?? '');
+	const title = $derived(getBuiltinToolUi(section.toolName)?.label ?? section.toolName ?? '');
 </script>

 <ToolCallBlock {section} {open} {isStreaming} meta={runJsMeta} {title} {onToggle}>
@@ -14,8 +14,8 @@
 	import { ICON_CLASS_DEFAULT, ICON_CLASS_SPIN } from '$lib/constants';
 	import { AgenticSectionType } from '$lib/enums';
 	import { mcpStore } from '$lib/stores';
-	import type { AgenticSection, ToolUiEntry } from '$lib/types';
-	import { getToolUi } from '$lib/utils';
+	import type { AgenticSection, BuiltinToolUiEntry } from '$lib/types';
+	import { getBuiltinToolUi } from '$lib/utils';
 	import type { Component, Snippet } from 'svelte';

 	type ToolCallBlockMetaWithError = TMeta & { errorMessage?: string };
@@ -82,7 +82,7 @@
 	const showSpinner = $derived(isPending || (isStreamingCall && isStreaming) || extraLiveStreaming);
 	const isCodeStreaming = $derived(isStreaming && (isPending || isStreamingCall));

-	const toolUi: ToolUiEntry | null = $derived(getToolUi(section.toolName));
+	const toolUi: BuiltinToolUiEntry | null = $derived(getBuiltinToolUi(section.toolName));
 	const toolIcon: Component = $derived(
 		spinIconWhenActive && showSpinner ? Loader2 : (toolUi?.icon ?? Wrench)
 	);
@@ -24,7 +24,7 @@ export type EditFileMeta = {
 };

 export function parseEditFileMeta(section: AgenticSection): EditFileMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_EDIT_FILE, section, { partial: true });
+	const args = parseToolArgs(BuiltInTool.EDIT_FILE, section, { partial: true });

 	if (!args) return null;

@@ -14,7 +14,7 @@ export type ExecShellCommandMeta = {
 };

 export function parseExecShellCommandMeta(section: AgenticSection): ExecShellCommandMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_EXEC_SHELL_COMMAND, section);
+	const args = parseToolArgs(BuiltInTool.EXEC_SHELL_COMMAND, section);

 	if (!args) return null;

@@ -19,7 +19,7 @@ export type FileGlobSearchMeta = {
 };

 export function parseFileGlobSearchMeta(section: AgenticSection): FileGlobSearchMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_FILE_GLOB_SEARCH, section);
+	const args = parseToolArgs(BuiltInTool.FILE_GLOB_SEARCH, section);

 	if (!args) return null;

@@ -28,7 +28,7 @@ export type GrepSearchMeta = {
 };

 export function parseGrepSearchMeta(section: AgenticSection): GrepSearchMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_GREP_SEARCH, section);
+	const args = parseToolArgs(BuiltInTool.GREP_SEARCH, section);

 	if (!args) return null;

@@ -16,7 +16,7 @@ export type ReadFileMeta = {
 };

 export function parseReadFileMeta(section: AgenticSection): ReadFileMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_READ_FILE, section, { partial: true });
+	const args = parseToolArgs(BuiltInTool.READ_FILE, section, { partial: true });

 	if (!args) return null;

@@ -16,7 +16,7 @@ export type RunJavascriptMeta = {
 };

 export function parseRunJavascriptMeta(section: AgenticSection): RunJavascriptMeta | null {
-	const args = parseToolArgs(BuiltInTool.BROWSER_RUN_JAVASCRIPT, section);
+	const args = parseToolArgs(BuiltInTool.RUN_JAVASCRIPT, section);

 	if (!args) return null;

@@ -20,7 +20,7 @@ export type WriteFileMeta = {
 };

 export function parseWriteFileMeta(section: AgenticSection): WriteFileMeta | null {
-	const args = parseToolArgs(BuiltInTool.SERVER_WRITE_FILE, section, { partial: true });
+	const args = parseToolArgs(BuiltInTool.WRITE_FILE, section, { partial: true });

 	if (!args) return null;

@@ -61,8 +61,8 @@
 				{:else}
 					{@const source = toolsStore.getToolSource(toolName)}
 					{@const providerName =
-						source === ToolSource.SERVER
-							? TOOL_SERVER_LABELS[ToolSource.SERVER]
+						source === ToolSource.BUILTIN
+							? TOOL_SERVER_LABELS[ToolSource.BUILTIN]
 							: source === ToolSource.CUSTOM
 								? TOOL_SERVER_LABELS[ToolSource.CUSTOM]
 								: 'MCP Tools'}
@@ -21,7 +21,8 @@
 	import {
 		chatStore,
 		conversationsStore,
-		deviceStore,
+		device,
+		isMobile,
 		serverStore,
 		settingsStore
 	} from '$lib/stores';
@@ -31,7 +32,7 @@
 	let { showCenteredEmpty = false } = $props();

 	let disableAutoScroll = $derived(
-		Boolean(settingsStore.config.disableAutoScroll) || deviceStore.isMobile
+		Boolean(settingsStore.config.disableAutoScroll) || isMobile.current
 	);
 	let isMobileUserScrolledUp = $state(false);
 	let mobileScrollDownHint = $state(false);
@@ -51,11 +52,11 @@
 	let hasPropsError = $derived(!!serverStore.error);
 	let isCurrentConversationLoading = $derived(chatStore.isLoading || chatStore.isStreaming());
 	let chatFormBottomPosition = $derived.by(() => {
-		if (!deviceStore.isMobile) return '1rem';
+		if (!isMobile.current) return '1rem';

-		if (deviceStore.isStandalone) return '1.5rem';
+		if (device.isStandalone) return '1.5rem';

-		if (deviceStore.isIOSSafari) return '0.25rem';
+		if (device.isIOSSafari) return '0.25rem';

 		return '0.5rem';
 	});
@@ -83,7 +84,7 @@
 	});

 	function handleMobileScroll() {
-		if (!deviceStore.isMobile) return;
+		if (!isMobile.current) return;

 		const container = scroll.chatScrollContainer;

@@ -183,7 +184,7 @@
 	}

 	function handleSendLikeScroll() {
-		if (!deviceStore.isMobile) {
+		if (!isMobile.current) {
 			autoScroll.enable();
 		}

@@ -196,7 +197,7 @@
 				'.chat-message:nth-last-child(2) .chat-message-user .chat-message-user-bubble'
 			) as HTMLElement | null;

-			if (deviceStore.isMobile) {
+			if (isMobile.current) {
 				// Keep the last user message bubble just above the input on mobile
 				const bubbleHeight = lastUserBubble?.scrollHeight ?? 0;
 				const baseHeight = container.scrollHeight - innerHeight;
@@ -219,7 +220,7 @@
 			}
 		}, 100);

-		if (deviceStore.isMobile) {
+		if (isMobile.current) {
 			autoScroll.setDisabled(disableAutoScroll);
 			mobileScrollDownHint = true;
 			mobileScrollDownHintLockedUntil = Date.now() + 500;
@@ -242,8 +243,7 @@

 	$effect(() => {
 		const shouldDisableAutoScroll =
-			settingsStore.config.disableAutoScroll ||
-			(deviceStore.isMobile && isCurrentConversationLoading);
+			settingsStore.config.disableAutoScroll || (isMobile.current && isCurrentConversationLoading);

 		autoScroll.setDisabled(shouldDisableAutoScroll);

@@ -266,7 +266,7 @@
 			autoScroll.enable();
 		}

-		if (deviceStore.isMobile && isCurrentConversationLoading) {
+		if (isMobile.current && isCurrentConversationLoading) {
 			mobileScrollDownHint = true;
 			mobileScrollDownHintLockedUntil = Date.now() + 500;
 		}
@@ -318,9 +318,9 @@
 		<div
 			class={[
 				'pointer-events-none md:sticky fixed  mt-auto transition-all duration-200',
-				deviceStore.isStandalone
+				device.isStandalone
 					? 'bottom-6 right-4 left-4'
-					: deviceStore.isIOSSafari
+					: device.isIOSSafari
 						? 'bottom-1 left-2 right-2'
 						: 'bottom-2 right-2 left-2',
 				isEmpty ? 'md:bottom-[calc(50dvh-7rem)] 2xl:bottom-[calc(50dvh-4rem)]' : 'md:bottom-4'
@@ -336,7 +336,7 @@
 			{/if}

 			<div class="pointer-events-none flex flex-col gap-6 items-center w-full">
-				{#if (deviceStore.isMobile ? mobileScrollDownHint || isMobileUserScrolledUp : autoScroll.userScrolledUp) && page.url.hash.includes(ROUTES.CHAT) && page.params.id}
+				{#if (isMobile.current ? mobileScrollDownHint || isMobileUserScrolledUp : autoScroll.userScrolledUp) && page.url.hash.includes(ROUTES.CHAT) && page.params.id}
 					<ChatScreenActionScrollDown
 						onclick={() => {
 							mobileScrollDownHint = false;
@@ -3,7 +3,7 @@
 	import { page } from '$app/state';
 	import { ChatForm } from '$lib/components/app';
 	import { useDraftMessages } from '$lib/hooks/use-draft-messages.svelte';
-	import { deviceStore } from '$lib/stores';
+	import { isMobile } from '$lib/stores';
 	import { onMount } from 'svelte';

 	interface Props {
@@ -120,13 +120,13 @@
 	}

 	onMount(() => {
-		if (!deviceStore.isMobile) {
+		if (!isMobile.current) {
 			setTimeout(focusFormUnlessCaptured, 100);
 		}
 	});

 	afterNavigate((navigation) => {
-		if (navigation?.from != null && !deviceStore.isMobile) {
+		if (navigation?.from != null && !isMobile.current) {
 			setTimeout(focusFormUnlessCaptured, 100);
 		}
 	});
@@ -278,7 +278,7 @@ export { default as ChatFormInput } from './ChatForm/ChatFormInput/ChatFormInput
 /**
 * Working directory selector for agent mode. Renders a chip below the chat
 * form; clicking it opens a popover with a directory picker backed by the
- * server's `file_glob_search` server tool (POST /tools). The picked
+ * server's `file_glob_search` built-in tool (POST /tools). The picked
 * directory is exposed via `bind:directory`; changing it records a
 * synthetic "Set working directory to ..." user message into chat history
 * and is enforced on tool calls via the `x-tool-cwd` request header.
@@ -380,7 +380,7 @@ export { default as ChatFormPickerListItemSkeleton } from './ChatForm/ChatFormPi

 /**
 * `@`-triggered file/folder mention picker. Resolves `@<query>` in the chat
- * input to a filesystem match via the server's `file_glob_search` server tool
+ * input to a filesystem match via the server's `file_glob_search` built-in
 * tool, scoped to the conversation cwd (or server home when unset).
 * Selection splices a `[name](file:///<abs path>)` link into the input.
 */
@@ -14,7 +14,7 @@
 	import { useKeyboardShortcuts } from '$lib/hooks/use-keyboard-shortcuts.svelte';
 	import { useMarqueeSelection } from '$lib/hooks/use-marquee-selection.svelte';
 	import { RouterService } from '$lib/services/router.service';
-	import { chatStore, conversationsStore, deviceStore, settingsStore } from '$lib/stores';
+	import { chatStore, conversationsStore, device, isMobile, settingsStore } from '$lib/stores';
 	import { buildConversationTree } from '$lib/utils';
 	import { circIn } from 'svelte/easing';
 	import { SvelteSet } from 'svelte/reactivity';
@@ -36,7 +36,7 @@
 	let logoHovered = $state(false);

 	const isStripExpanded = $derived(isExpandedMode || hoveredTooltip !== null);
-	const isOnMobile = $derived(deviceStore.isMobile);
+	const isOnMobile = $derived(isMobile.current);
 	const alwaysShowOnDesktop = $derived(settingsStore.config.alwaysShowSidebarOnDesktop as boolean);

 	$effect(() => {
@@ -65,7 +65,7 @@
 	});

 	$effect(() => {
-		if (deviceStore.isMobile && page.url.hash.includes(ROUTES.SEARCH)) {
+		if (isMobile.current && page.url.hash.includes(ROUTES.SEARCH)) {
 			isExpandedMode = false;
 		}
 	});
@@ -227,7 +227,7 @@
 	}

 	async function selectConversation(id: string) {
-		if (deviceStore.isMobile) {
+		if (isMobile.current) {
 			scheduleMobileCollapse();
 		}

@@ -315,9 +315,9 @@
 			'fixed md:sticky top-2 left-2 md:left-0 md:ml-2 md:mt-2 pt-2 z-10 w-[calc(100dvw-1rem)]',
 			'md:h-[calc(100dvh-1.125rem)]',
 			isExpandedMode &&
-				(deviceStore.isStandalone
+				(device.isStandalone
 					? 'h-[calc(100dvh-2rem)]'
-					: deviceStore.isIOSDevice
+					: device.isIOSDevice
 						? 'h-[calc(100dvh-0.5rem)]'
 						: 'h-[calc(100dvh-1rem)]'),
 			'rounded-3xl md:rounded-2xl',
@@ -353,7 +353,7 @@

 			{#if isOnMobile || (isExpandedMode && !alwaysShowOnDesktop)}
 				<div
-					class="flex items-center transition-all duration-150 ease-out {deviceStore.isMobile &&
+					class="flex items-center transition-all duration-150 ease-out {isMobile.current &&
 					!isExpandedMode
 						? 'opacity-0 h-0!'
 						: ''}"
@@ -361,7 +361,7 @@
 					out:fade={{ duration: 100 }}
 				>
 					<ActionIcon
-						icon={deviceStore.isMobile ? X : PanelLeftClose}
+						icon={isMobile.current ? X : PanelLeftClose}
 						size="lg"
 						iconSize="h-4.5 w-4.5 md:h-4 md:w-4"
 						class="backdrop-blur-none md:h-9 md:w-9 h-10 w-10 rounded-full mr-1 hover:bg-accent!"
@@ -375,9 +375,9 @@
 		</div>

 		<div
-			class="mt-2 flex min-h-0 flex-1 flex-col gap-4 md:gap-1 {deviceStore.isMobile
+			class="mt-2 flex min-h-0 flex-1 flex-col gap-4 md:gap-1 {isMobile.current
 				? 'transition-[opacity,height] duration-200 ease-out'
-				: ''} {deviceStore.isMobile && !isExpandedMode ? 'opacity-0 !h-0' : ''}"
+				: ''} {isMobile.current && !isExpandedMode ? 'opacity-0 !h-0' : ''}"
 			in:fade={{ duration: 200 }}
 			out:fade={{ duration: 200 }}
 		>
@@ -395,7 +395,7 @@
 					isSearchModeActive = true;
 				}}
 				onNewChat={() => {
-					if (deviceStore.isMobile) {
+					if (isMobile.current) {
 						scheduleMobileCollapse();
 					}
 				}}
@@ -12,7 +12,7 @@
 		SIDEBAR_ACTIONS_ITEMS
 	} from '$lib/constants';
 	import { TooltipSide } from '$lib/enums';
-	import { deviceStore } from '$lib/stores';
+	import { isMobile } from '$lib/stores';
 	import type { Component } from 'svelte';
 	import { onMount } from 'svelte';
 	import { circIn } from 'svelte/easing';
@@ -42,7 +42,7 @@
 	let showIcons = $state(false);
 	let searchInputRef = $state<HTMLInputElement | null>(null);

-	const isOnMobile = $derived(deviceStore.isMobile);
+	const isOnMobile = $derived(isMobile.current);

 	$effect(() => {
 		if (isSearchModeActive && searchInputRef) {
@@ -107,7 +107,7 @@
 	>
 		{#each SIDEBAR_ACTIONS_ITEMS as item, i (item.tooltip)}
 			{@const isActive = isItemActive(item)}
-			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
+			{@const isSearchOnMobile = item.icon === Search && isMobile.current}
 			{@const itemHref = isSearchOnMobile ? ROUTES.SEARCH : item.route}
 			{@const itemOnClick = item.route
 				? () => {
@@ -156,7 +156,7 @@
 	<div class="{className} flex-col gap-1 hidden md:flex">
 		{#each SIDEBAR_ACTIONS_ITEMS as item, i (item.tooltip)}
 			{@const isActive = isItemActive(item)}
-			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
+			{@const isSearchOnMobile = item.icon === Search && isMobile.current}
 			{@const itemOnClick = item.route
 				? () => {
 						onNewChat?.();
@@ -8,7 +8,6 @@
 	} from '$lib/components/app';
 	import SettingsGroup from '$lib/components/app/settings/SettingsGroup.svelte';
 	import { ConversationSelectionMode, FileExtensionText, HtmlInputType } from '$lib/enums';
-	import { ConversationTransferService } from '$lib/services';
 	import { conversationsStore, settingsStore } from '$lib/stores';
 	import { createMessageCountMap } from '$lib/utils';
 	import { fade } from 'svelte/transition';
@@ -148,9 +147,9 @@
 			);

 			if (allData.length === 1) {
-				ConversationTransferService.downloadConversationFile(allData[0]);
+				conversationsStore.downloadConversationFile(allData[0]);
 			} else {
-				ConversationTransferService.downloadConversationsArchive(allData);
+				conversationsStore.downloadConversationsArchive(allData);
 			}

 			exportedConversations = selectedConversations;
@@ -178,7 +177,7 @@
 				if (!file) return;

 				try {
-					const importedData = await ConversationTransferService.parseImportFile(file);
+					const importedData = await conversationsStore.parseImportFile(file);

 					if (importedData.length === 0) {
 						throw new Error('No conversations found in file');
@@ -6,7 +6,7 @@
 	import { ICON_CLASS_DEFAULT } from '$lib/constants';
 	import { ToolSource } from '$lib/enums/tools.enums';
 	import { mcpStore, permissionsStore, toolsStore } from '$lib/stores';
-	import { getToolUi } from '$lib/utils';
+	import { getBuiltinToolUi } from '$lib/utils';
 	import { SvelteSet } from 'svelte/reactivity';

 	let expandedGroups = new SvelteSet<string>();
@@ -69,12 +69,12 @@

 						{#each group.tools as entry (entry.key)}
 							{@const toolName = entry.definition.function.name}
-							{@const toolUi =
-								entry.source === ToolSource.SERVER || entry.source === ToolSource.BROWSER
-									? getToolUi(toolName)
+							{@const builtinUi =
+								entry.source === ToolSource.BUILTIN || entry.source === ToolSource.FRONTEND
+									? getBuiltinToolUi(toolName)
 									: null}
-							{@const displayLabel = toolUi?.label ?? toolName}
-							{@const IconComponent = toolUi?.icon ?? null}
+							{@const displayLabel = builtinUi?.label ?? toolName}
+							{@const IconComponent = builtinUi?.icon ?? null}
 							{@const isEnabled = toolsStore.isToolEnabled(entry.key)}
 							{@const permissionKey = entry.key}
 							{@const isAlwaysAllowed = permissionsStore.hasTool(permissionKey)}
@@ -69,7 +69,7 @@ export { default as SettingsChatFields } from './SettingsChat/SettingsChatFields
 /**
 * **SettingsChatToolsTab** - Tools configuration tab for chat settings
 *
- * Displays available tools grouped by source (server, browser, MCP, custom) with
+ * Displays available tools grouped by source (built-in, MCP, custom) with
 * toggles to enable/disable individual tools and tool groups. Shows MCP
 * server favicons and permission management controls.
 */
@@ -31,11 +31,5 @@ export const API_STREAM = {
 	LOOKUP: './v1/streams/lookup'
 };

-// query params for the resumable stream routes
-export const STREAM_QUERY_PARAMS = {
-	CONV_ID: 'conv_id',
-	FROM: 'from'
-} as const;
-
 /** CORS proxy endpoint path */
 export const CORS_PROXY_ENDPOINT = '/cors-proxy';
@@ -2,9 +2,7 @@ import { CLI_FLAGS } from './cli-flags.constants';
 import { BuiltInTool, JsonSchemaType, ToolCallType } from '$lib/enums';
 import type { OpenAIToolDefinition } from '$lib/types';

-// get_info is served by the server, but the browser falls back to this
-// implementation when the server does not provide it - same wire name.
-export const BROWSER_INFO_TOOL_NAME = BuiltInTool.SERVER_GET_INFO;
+export const BROWSER_INFO_TOOL_NAME = BuiltInTool.GET_INFO;

 /** UA token to OS name, first match wins - Android and iOS UAs also carry the Linux / Mac OS X tokens */
 export const BROWSER_INFO_OS_UA_PATTERNS: readonly [RegExp, string][] = [
--- a/Show More
+++ b/Show More