demo, wip

.devops/openvino.Dockerfile

@@ -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; \

.github/actions/linux-setup-vulkan/action.yml

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

.github/actions/windows-setup-cuda/action.yml

@@ -6,7 +6,8 @@ inputs:
     required: true
   cuda_arch:
     description: "CUDA target architecture"
-    required: true
+    required: false
+    default: "x64"
 
 runs:
   using: "composite"

.github/workflows/build-cache.yml

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

.github/workflows/build-cpu.yml

@@ -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: |

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

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

.github/workflows/build-openvino.yml

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

.github/workflows/build-rpc.yml

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

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

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

.github/workflows/build-vulkan.yml

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

.github/workflows/docker.yml

@@ -394,11 +394,6 @@ jobs:
     name: Create shared tags from digests
     needs: [prepare_matrices, push_to_registry, create_tag]
     runs-on: ubuntu-24.04
-    permissions:
-      contents: read
-      packages: write
-      id-token: write
-      attestations: write
     strategy:
       fail-fast: false
       matrix:
@@ -433,7 +428,6 @@ jobs:
           password: ${{ secrets.GITHUB_TOKEN }}
 
       - name: Create tags from digests
-        id: create_tags
         shell: bash
         run: |
           set -euo pipefail
@@ -445,7 +439,6 @@ jobs:
           SRC_TAG="${{ needs.create_tag.outputs.source_tag }}"
           BUILD_DATE="${{ steps.build_date.outputs.date }}"
           COMMIT_SHA="${{ steps.checkout.outputs.commit }}"
-          echo "image_repo=${IMAGE_REPO}" >> "$GITHUB_OUTPUT"
           TAGS="${{ matrix.config.tag }}"
           ARCHES="${{ matrix.config.arches }}"
           DIGEST_GLOB="/tmp/digests/*.tsv"
@@ -512,16 +505,6 @@ jobs:
 
               echo "Creating ${merged_versioned_tag} from ${refs[*]}"
               docker buildx imagetools create "${annotations[@]}" --tag "${merged_versioned_tag}" "${refs[@]}"
-
-              if [[ "$tag_name" == "${TAGS%% *}" ]]; then
-                  local digest
-                  digest="$(docker buildx imagetools inspect "${merged_versioned_tag}" --format '{{.Manifest.Digest}}')"
-                  if [[ ! "$digest" =~ ^sha256:[0-9a-f]{64}$ ]]; then
-                      echo "Invalid digest for ${merged_versioned_tag}: ${digest}" >&2
-                      exit 1
-                  fi
-                  echo "${image_type}_digest=${digest}" >> "$GITHUB_OUTPUT"
-              fi
           }
 
           for tag in $TAGS; do
@@ -545,24 +528,3 @@ jobs:
           done
         env:
           GITHUB_REPOSITORY_OWNER: '${{ github.repository_owner }}'
-
-      - name: Attest full image
-        if: ${{ matrix.config.full }}
-        uses: actions/attest@v4
-        with:
-          subject-name: ${{ steps.create_tags.outputs.image_repo }}
-          subject-digest: ${{ steps.create_tags.outputs.full_digest }}
-
-      - name: Attest light image
-        if: ${{ matrix.config.light }}
-        uses: actions/attest@v4
-        with:
-          subject-name: ${{ steps.create_tags.outputs.image_repo }}
-          subject-digest: ${{ steps.create_tags.outputs.light_digest }}
-
-      - name: Attest server image
-        if: ${{ matrix.config.server }}
-        uses: actions/attest@v4
-        with:
-          subject-name: ${{ steps.create_tags.outputs.image_repo }}
-          subject-digest: ${{ steps.create_tags.outputs.server_digest }}

.github/workflows/make-release.yml

@@ -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: |

.github/workflows/release.yml

@@ -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
@@ -1569,8 +1569,6 @@ jobs:
     # https://docs.github.com/en/actions/security-for-github-actions/security-guides/automatic-token-authentication#modifying-the-permissions-for-the-github_token
     permissions:
         contents: write # for creating release
-        id-token: write
-        attestations: write
 
     runs-on: ubuntu-slim
 
@@ -1664,11 +1662,6 @@ jobs:
         run: |
           tar -czvf release/llama-${{ steps.tag.outputs.name }}-ui.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./ui-dist .
 
-      - name: Attest release artifacts
-        uses: actions/attest@v4
-        with:
-          subject-path: 'release/*'
-
       - name: Create and push git tag
         run: |
           TAG="${{ steps.tag.outputs.name }}"

CMakeLists.txt

@@ -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)

README.md

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

build-xcframework.sh

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

ci/run.sh

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

common/CMakeLists.txt

@@ -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)

common/arg.cpp

@@ -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"},
@@ -4658,12 +4690,6 @@ void common_params_add_preset_options(std::vector<common_arg> & args) {
         [](common_params &, int) { /* unused */ }
     ).set_env(COMMON_ARG_PRESET_STOP_TIMEOUT).set_preset_only());
 
-    args.push_back(common_arg(
-        {"dedup-cache-models"}, "0|1",
-        "in server router mode, hide a cached model from the model list when this preset resolves to the same model file",
-        [](common_params &, const std::string &) { /* unused */ }
-    ).set_env(COMMON_ARG_PRESET_DEDUP_CACHE_MODELS).set_preset_only());
-
     // args.push_back(common_arg(
     //     {"pin"},
     //     "in server router mode, do not unload this model if models_max is exceeded",

common/arg.h

@@ -11,9 +11,8 @@
 #include <memory>
 
 // pseudo-env variable to identify preset-only arguments
-#define COMMON_ARG_PRESET_LOAD_ON_STARTUP    "__PRESET_LOAD_ON_STARTUP"
-#define COMMON_ARG_PRESET_STOP_TIMEOUT       "__PRESET_STOP_TIMEOUT"
-#define COMMON_ARG_PRESET_DEDUP_CACHE_MODELS "__PRESET_DEDUP_CACHE_MODELS"
+#define COMMON_ARG_PRESET_LOAD_ON_STARTUP "__PRESET_LOAD_ON_STARTUP"
+#define COMMON_ARG_PRESET_STOP_TIMEOUT    "__PRESET_STOP_TIMEOUT"
 
 //
 // CLI argument parsing

common/common.cpp

@@ -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);

common/common.h

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

common/download.cpp

@@ -989,26 +989,6 @@ std::vector<common_cached_model_info> common_list_cached_models() {
     return result;
 }
 
-std::string common_download_resolve_path(const std::string & hf_repo_with_tag, const std::string & hf_file) {
-    auto [repo, tag] = common_download_split_repo_tag(hf_repo_with_tag);
-
-    auto files = hf_cache::get_cached_files(repo);
-    if (files.empty()) {
-        return "";
-    }
-
-    if (!hf_file.empty()) {
-        for (const auto & f : files) {
-            if (f.path == hf_file) {
-                return f.local_path;
-            }
-        }
-        return "";
-    }
-
-    return find_best_model(files, tag).local_path;
-}
-
 bool common_download_remove(const std::string & hf_repo_with_tag) {
     namespace fs = std::filesystem;
 

common/download.h

@@ -85,10 +85,6 @@ std::vector<std::string> common_download_get_all_parts(const std::string & url);
 // returns list of cached models
 std::vector<common_cached_model_info> common_list_cached_models();
 
-// resolve the local cached file path for a HF repo without network access (hf_file, if given, must match exactly)
-// returns an empty string if the model is not present in the cache
-std::string common_download_resolve_path(const std::string & hf_repo_with_tag, const std::string & hf_file = "");
-
 // download single file from url to local path
 // returns status code or -1 on error
 // skip_etag: if true, don't read/write .etag files (for HF cache where filename is the hash)

docs/backend/OPENVINO.md

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

examples/gguf-hash/CMakeLists.txt

@@ -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)

examples/gguf-hash/gguf-hash.cpp

@@ -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;
 
 

ggml/CMakeLists.txt

@@ -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/")

ggml/include/ggml-rpc.h

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

ggml/include/ggml.h

@@ -1981,14 +1981,6 @@ extern "C" {
             float                 beta_fast,
             float                 beta_slow);
 
-    // set the offset dims for RoPE
-    // a must be GGML_OP_ROPE or GGML_OP_ROPE_BACK
-    // vision RoPE is not supported
-    // example: (marking: x = rotated, 0 = unrotated)
-    //     n_embd = 10, n_dims = 4, offset = 2 --> [00xxxx0000]
-    GGML_API struct ggml_tensor * ggml_rope_set_offset(
-            struct ggml_tensor  * a,
-            int                   n_offs);
 
     // clamp
     // in-place, returns view(a)

ggml/src/ggml-cann/ggml-cann.cpp

@@ -2534,9 +2534,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
             }
         case GGML_OP_ROPE:
             {
-                if (((const int32_t *) op->op_params)[15] != 0) {
-                    return false; // FIXME: support ggml_rope_set_offset
-                }
                 if (op->src[0]->ne[0] > 896) {
                     return false;
                 }

ggml/src/ggml-cpu/ops.cpp

@@ -5979,8 +5979,6 @@ static void ggml_compute_forward_rope_flt(
     memcpy(&beta_slow,   (int32_t *) dst->op_params + 10, sizeof(float));
     memcpy(&sections,    (int32_t *) dst->op_params + 11, sizeof(int)*4);
 
-    const int n_offs = ((int32_t *) dst->op_params)[15];
-
     GGML_TENSOR_UNARY_OP_LOCALS
 
     //printf("ne0: %d, ne1: %d, ne2: %d, ne3: %d\n", ne0, ne1, ne2, ne3);
@@ -5997,10 +5995,6 @@ static void ggml_compute_forward_rope_flt(
     GGML_ASSERT(n_dims <= ne0);
     GGML_ASSERT(n_dims % 2 == 0);
 
-    GGML_ASSERT(n_offs >= 0);
-    GGML_ASSERT(n_offs % 2 == 0);
-    GGML_ASSERT(n_offs + n_dims <= ne0);
-
     // rows per thread
     const int dr = (nr + nth - 1)/nth;
 
@@ -6026,7 +6020,6 @@ static void ggml_compute_forward_rope_flt(
 
     if (is_vision) {
         GGML_ASSERT(n_dims == ne0/2);
-        GGML_ASSERT(n_offs == 0);
     }
 
     const float * freq_factors = NULL;
@@ -6075,12 +6068,12 @@ static void ggml_compute_forward_rope_flt(
 
                 switch (mode) {
                     case GGML_ROPE_TYPE_NORMAL:
-                        rotate_pairs<T>(n_dims, 1, cache, src + n_offs, dst_data + n_offs, 1);
+                        rotate_pairs<T>(n_dims, 1, cache, src, dst_data, 1);
                         break;
                     case GGML_ROPE_TYPE_NEOX:
                     case GGML_ROPE_TYPE_MROPE:
                     case GGML_ROPE_TYPE_IMROPE:
-                        rotate_pairs<T>(n_dims, n_dims/2, cache, src + n_offs, dst_data + n_offs);
+                        rotate_pairs<T>(n_dims, n_dims/2, cache, src, dst_data);
                         break;
                     case GGML_ROPE_TYPE_VISION:
                         rotate_pairs<T>(ne0, n_dims, cache, src, dst_data);
@@ -6091,11 +6084,7 @@ static void ggml_compute_forward_rope_flt(
 
                 if (!is_vision) {
                     // fill the remain channels with data from src tensor
-                    for (int64_t i0 = 0; i0 < ne0; i0 += 2) {
-                        if (i0 == n_offs) {
-                            i0 += n_dims - 2; // skip the rotated channels
-                            continue;
-                        }
+                    for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) {
                         const T * const src = (T *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
                         T * dst_data  = (T *)((char *)  dst->data + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 

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

@@ -2723,12 +2723,6 @@ static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm
         return false;
     }
 
-    // ggml_rope_set_offset is not yet supported in the fused kernel
-    const int n_offs = ((const int32_t *) rope->op_params)[15];
-    if (n_offs != 0) {
-        return false;
-    }
-
     return true;
 }
 

ggml/src/ggml-cuda/mmvq.cu

@@ -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: {

ggml/src/ggml-cuda/rope.cu

@@ -53,7 +53,6 @@ static __global__ void rope_norm(const T *            x,
                                  const int            s2,
                                  const int            s3,
                                  const int            n_dims,
-                                 const int            n_offs,
                                  const int32_t *      pos,
                                  const float          freq_scale,
                                  const float          ext_factor,
@@ -62,8 +61,7 @@ static __global__ void rope_norm(const T *            x,
                                  const float          theta_scale,
                                  const float *        freq_factors,
                                  const int64_t *      row_indices,
-                                 const int            set_rows_stride,
-                                 const bool           inplace) {
+                                 const int            set_rows_stride) {
     const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y);
 
     if (i0 >= ne00) {
@@ -94,24 +92,19 @@ static __global__ void rope_norm(const T *            x,
             ggml_cuda_memcpy_1<4>(dst + idst, &v);
         }
     };
-    if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
+    if (i0 >= n_dims) {
         store_coaelsced(x[ix + 0], x[ix + 1]);
         return;
     }
 
-    const int iw = i0 - n_offs; // relative idx
+    const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
 
-    const float theta_base = pos[i2]*powf(theta_scale, iw/2.0f);
-
-    const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
+    const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
 
     float cos_theta;
     float sin_theta;
 
-    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
+    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
 
     const float x0 = x[ix + 0];
     const float x1 = x[ix + 1];
@@ -132,7 +125,6 @@ static __global__ void rope_neox(const T *            x,
                                  const int            s2,
                                  const int            s3,
                                  const int            n_dims,
-                                 const int            n_offs,
                                  const int32_t *      pos,
                                  const float          freq_scale,
                                  const float          ext_factor,
@@ -141,8 +133,7 @@ static __global__ void rope_neox(const T *            x,
                                  const float          theta_scale,
                                  const float *        freq_factors,
                                  const int64_t *      row_indices,
-                                 const int            set_rows_stride,
-                                 const bool           inplace) {
+                                 const int            set_rows_stride) {
     ggml_cuda_pdl_lc();
     const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y);
 
@@ -167,33 +158,27 @@ static __global__ void rope_neox(const T *            x,
         idst += row_indices[i2] * set_rows_stride;
     }
 
-    if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
+    if (i0 >= n_dims) {
         dst[idst + i0 / 2 + 0] = ggml_cuda_cast<D>(x[ix + i0 / 2 + 0]);
         dst[idst + i0 / 2 + 1] = ggml_cuda_cast<D>(x[ix + i0 / 2 + 1]);
 
         return;
     }
 
-    const int iw = i0 - n_offs; // relative idx
+    const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
 
-    const float theta_base = pos[i2]*powf(theta_scale, iw/2.0f);
-
-    const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
+    const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
 
     float cos_theta;
     float sin_theta;
 
-    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
+    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
 
-    // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
-    const float x0 = x[ix + n_offs/2 + 0];
-    const float x1 = x[ix + n_offs/2 + n_dims/2];
+    const float x0 = x[ix + 0];
+    const float x1 = x[ix + n_dims/2];
 
-    dst[idst + n_offs/2 + 0]          = ggml_cuda_cast<D>(x0 * cos_theta - x1 * sin_theta);
-    dst[idst + n_offs/2 + n_dims / 2] = ggml_cuda_cast<D>(x0 * sin_theta + x1 * cos_theta);
+    dst[idst + 0]          = ggml_cuda_cast<D>(x0 * cos_theta - x1 * sin_theta);
+    dst[idst + n_dims / 2] = ggml_cuda_cast<D>(x0 * sin_theta + x1 * cos_theta);
 }
 
 template <bool forward, bool has_ff, typename T>
@@ -209,7 +194,6 @@ static __global__ void rope_multi(const T *            x,
                                   const int            s2,
                                   const int            s3,
                                   const int            n_dims,
-                                  const int            n_offs,
                                   const int32_t *      pos,
                                   const float          freq_scale,
                                   const float          ext_factor,
@@ -218,8 +202,7 @@ static __global__ void rope_multi(const T *            x,
                                   const float          theta_scale,
                                   const float *        freq_factors,
                                   const mrope_sections sections,
-                                  const bool           is_imrope,
-                                  const bool           inplace) {
+                                  const bool           is_imrope) {
     const int i0 = 2 * (blockDim.y * blockIdx.y + threadIdx.y);
 
     if (i0 >= ne00) {
@@ -236,58 +219,52 @@ static __global__ void rope_multi(const T *            x,
     const int ix   = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03;
 
     ggml_cuda_pdl_sync();
-    if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
+    if (i0 >= n_dims) {
         dst[idst + i0/2 + 0] = x[ix + i0/2 + 0];
         dst[idst + i0/2 + 1] = x[ix + i0/2 + 1];
 
         return;
     }
 
-    const int iw = i0 - n_offs; // relative idx
-
     const int sect_dims = sections.v[0] + sections.v[1] + sections.v[2] + sections.v[3];
     const int sec_w = sections.v[1] + sections.v[0];
-    const int sector = (iw / 2) % sect_dims;
+    const int sector = (i0 / 2) % sect_dims;
 
     float theta_base = 0.0;
     if (is_imrope) {
         if (sector % 3 == 1 && sector < 3 * sections.v[1]) {         // h
-            theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, i0 / 2.0f);
         } else if (sector % 3 == 2 && sector < 3 * sections.v[2]) {  // w
-            theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, i0 / 2.0f);
         } else if (sector % 3 == 0 && sector < 3 * sections.v[0]) {  // t
-            theta_base = pos[i2] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2] * powf(theta_scale, i0 / 2.0f);
         } else {
-            theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, i0 / 2.0f);
         }
     } else {
         if (sector < sections.v[0]) {
-            theta_base = pos[i2] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2] * powf(theta_scale, i0 / 2.0f);
         } else if (sector >= sections.v[0] && sector < sec_w) {
-            theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, i0 / 2.0f);
         } else if (sector >= sec_w && sector < sec_w + sections.v[2]) {
-            theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, i0 / 2.0f);
         } else if (sector >= sec_w + sections.v[2]) {
-            theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, iw / 2.0f);
+            theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, i0 / 2.0f);
         }
     }
 
-    const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
+    const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
 
     float cos_theta;
     float sin_theta;
 
-    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
+    rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
 
-    // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
-    const float x0 = x[ix + n_offs/2 + 0];
-    const float x1 = x[ix + n_offs/2 + n_dims/2];
+    const float x0 = x[ix + 0];
+    const float x1 = x[ix + n_dims/2];
 
-    dst[idst + n_offs/2 + 0]        = x0*cos_theta - x1*sin_theta;
-    dst[idst + n_offs/2 + n_dims/2] = x0*sin_theta + x1*cos_theta;
+    dst[idst + 0]        = x0*cos_theta - x1*sin_theta;
+    dst[idst + n_dims/2] = x0*sin_theta + x1*cos_theta;
 }
 
 template <bool forward, bool has_ff, typename T>
@@ -367,7 +344,6 @@ static void rope_norm_cuda(const T *            x,
                            const int            s2,
                            const int            s3,
                            const int            n_dims,
-                           const int            n_offs,
                            const int            nr,
                            const int32_t *      pos,
                            const float          freq_scale,
@@ -378,7 +354,6 @@ static void rope_norm_cuda(const T *            x,
                            const float *        freq_factors,
                            const int64_t *      row_indices,
                            const int            set_rows_stride,
-                           const bool           inplace,
                            cudaStream_t         stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -389,12 +364,12 @@ static void rope_norm_cuda(const T *            x,
 
     if (freq_factors == nullptr) {
         rope_norm<forward, false><<<block_nums, block_dims, 0, stream>>>(
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
     } else {
         rope_norm<forward, true><<<block_nums, block_dims, 0, stream>>>(
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
     }
 }
 
@@ -411,7 +386,6 @@ static void rope_neox_cuda(const T *            x,
                            const int            s2,
                            const int            s3,
                            const int            n_dims,
-                           const int            n_offs,
                            const int            nr,
                            const int32_t *      pos,
                            const float          freq_scale,
@@ -422,7 +396,6 @@ static void rope_neox_cuda(const T *            x,
                            const float *        freq_factors,
                            const int64_t *      row_indices,
                            const int            set_rows_stride,
-                           const bool           inplace,
                            cudaStream_t         stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -434,12 +407,12 @@ static void rope_neox_cuda(const T *            x,
 
     if (freq_factors == nullptr) {
         ggml_cuda_kernel_launch(rope_neox<forward, false, T, D>, launch_params,
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
     } else {
         ggml_cuda_kernel_launch(rope_neox<forward, true, T, D>, launch_params,
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
     }
 }
 
@@ -456,7 +429,6 @@ static void rope_multi_cuda(const T *            x,
                             const int            s2,
                             const int            s3,
                             const int            n_dims,
-                            const int            n_offs,
                             const int            nr,
                             const int32_t *      pos,
                             const float          freq_scale,
@@ -467,7 +439,6 @@ static void rope_multi_cuda(const T *            x,
                             const float *        freq_factors,
                             const mrope_sections sections,
                             const bool           is_imrope,
-                            const bool           inplace,
                             cudaStream_t         stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -479,13 +450,13 @@ static void rope_multi_cuda(const T *            x,
     if (freq_factors == nullptr) {
         const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
         ggml_cuda_kernel_launch(rope_multi<forward, false, T>, launch_params,
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope);
     } else {
         const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
         ggml_cuda_kernel_launch(rope_multi<forward, true, T>, launch_params,
-            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
-            attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope, inplace);
+            x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
+            attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope);
     }
 }
 
@@ -581,12 +552,8 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
     const int mode       = ((int32_t *) dst->op_params)[2];
     //const int n_ctx      = ((int32_t *) dst->op_params)[3];
     const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
-    const int n_offs     = ((int32_t *) dst->op_params)[15];
     mrope_sections sections;
 
-    // when dst aliases src0, the channels outside the rotated window already hold the correct data
-    const bool inplace = dst_d == src0->data;
-
     // RoPE alteration for extended context
     float freq_base;
     float freq_scale;
@@ -614,7 +581,6 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
 
     if (is_vision) {
         GGML_ASSERT(n_dims == ne00/2);
-        GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row
     }
 
     const int32_t * pos = (const int32_t *) src1_d;
@@ -631,31 +597,31 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
     if (is_neox) {
         if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
             rope_neox_cuda<forward, float, float>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                  s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                  s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                   ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                  set_rows_stride, inplace, stream);
+                                                  set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
             rope_neox_cuda<forward, float, half>((const float *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                 s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                 s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                  ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                 set_rows_stride, inplace, stream);
+                                                 set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
             rope_neox_cuda<forward, half, half>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                 ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                set_rows_stride, inplace, stream);
+                                                set_rows_stride, stream);
         } else {
             GGML_ABORT("fatal error");
         }
     } else if (is_mrope && !is_vision) {
         if (src0->type == GGML_TYPE_F32) {
             rope_multi_cuda<forward>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02, s03, s1,
-                                     s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
-                                     corr_dims, freq_factors, sections, is_imrope, inplace, stream);
+                                     s2, s3, n_dims, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
+                                     corr_dims, freq_factors, sections, is_imrope, stream);
         } else if (src0->type == GGML_TYPE_F16) {
             rope_multi_cuda<forward>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02, s03, s1,
-                                     s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
-                                     corr_dims, freq_factors, sections, is_imrope, inplace, stream);
+                                     s2, s3, n_dims, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
+                                     corr_dims, freq_factors, sections, is_imrope, stream);
         } else {
             GGML_ABORT("fatal error");
         }
@@ -674,19 +640,19 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
     } else {
         if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
             rope_norm_cuda<forward, float, float>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                  s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                  s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                   ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                  set_rows_stride, inplace, stream);
+                                                  set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
             rope_norm_cuda<forward, float, half>((const float *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                 s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                 s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                  ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                 set_rows_stride, inplace, stream);
+                                                 set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
             rope_norm_cuda<forward, half, half>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
-                                                s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
+                                                s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
                                                 ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                                                set_rows_stride, inplace, stream);
+                                                set_rows_stride, stream);
         } else {
             GGML_ABORT("fatal error");
         }

ggml/src/ggml-et/ggml-et.cpp

@@ -1061,11 +1061,9 @@ static bool ggml_backend_et_device_supports_op(ggml_backend_dev_t dev, const ggm
                 const bool zero_view_offset = op->src[0]->view_src == nullptr || op->src[0]->view_offs == 0;
                 const bool has_sections = ggml_get_op_params_i32(op, 11) > 0 || ggml_get_op_params_i32(op, 12) > 0 ||
                                           ggml_get_op_params_i32(op, 13) > 0;
-                // FIXME: support ggml_rope_set_offset
-                const bool zero_rot_offset  = ggml_get_op_params_i32(op, 15) == 0;
 
                 supported =
-                    zero_view_offset && zero_rot_offset && ndims <= 512 &&
+                    zero_view_offset && ndims <= 512 &&
                     (is_normal || (is_neox && ndims % 16 == 0) || (is_imrope && ndims % 16 == 0 && has_sections));
             } else {
                 supported = false;

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

@@ -3180,10 +3180,6 @@ static bool ggml_hexagon_supported_argsort(const struct ggml_hexagon_session * s
 static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
     const int32_t * op_params = &op->op_params[0];
 
-    if (op_params[15] != 0) {
-        return false; // FIXME: support ggml_rope_set_offset
-    }
-
     int mode = op_params[2];
 
     // n_dims == ne0/2, so the rotation spans the full row

ggml/src/ggml-metal/ggml-metal-impl.h

@@ -329,7 +329,6 @@ typedef struct {
     uint64_t nb3;
     int32_t  n_past;
     int32_t  n_dims;
-    int32_t  n_offs;
     int32_t  n_ctx_orig;
     float    freq_base;
     float    freq_scale;
@@ -342,7 +341,6 @@ typedef struct {
     int32_t  sect_2;
     int32_t  sect_3;
     bool     src2;
-    bool     inplace;
 } ggml_metal_kargs_rope;
 
 typedef struct {

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -3884,11 +3884,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
     const int sect_2 = ((const int32_t *) op->op_params)[13];
     const int sect_3 = ((const int32_t *) op->op_params)[14];
 
-    const int n_offs = ((const int32_t *) op->op_params)[15];
-
-    // when dst aliases src0, the channels outside the rotated window already hold the correct data
-    const bool inplace = op->data == op->src[0]->data;
-
     ggml_metal_kargs_rope args = {
         /*.ne00        =*/ ne00,
         /*.ne01        =*/ ne01,
@@ -3908,7 +3903,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
         /*.nb3         =*/ nb3,
         /*.n_past      =*/ n_past,
         /*.n_dims      =*/ n_dims,
-        /*.n_offs      =*/ n_offs,
         /*.n_ctx_orig  =*/ n_ctx_orig,
         /*.freq_base   =*/ freq_base,
         /*.freq_scale  =*/ freq_scale,
@@ -3921,7 +3915,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
         /* sect_2      =*/ sect_2,
         /* sect_3      =*/ sect_3,
         /* src2        =*/ op->src[2] != nullptr,
-        /* inplace     =*/ inplace,
     };
 
     auto pipeline = ggml_metal_library_get_pipeline_rope(lib, op);

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

@@ -656,13 +656,13 @@ void dequantize_q5_1_t4(device const block_q5_1 * xb, short il, thread type4 & r
 
 template <typename type4x4>
 void dequantize_q8_0(device const block_q8_0 *xb, short il, thread type4x4 & reg) {
-    device const packed_char4 * qs = (device const packed_char4 *) xb->qs;
+    device const int8_t * qs = ((device const int8_t *)xb->qs);
     const float d = xb->d;
 
     float4x4 reg_f;
 
-    for (int i = 0; i < 4; ++i) {
-        reg_f[i] = float4(qs[4*il + i]) * d;
+    for (int i = 0; i < 16; i++) {
+        reg_f[i/4][i%4] = (qs[i + 16*il] * d);
     }
 
     reg = (type4x4) reg_f;
@@ -670,10 +670,12 @@ void dequantize_q8_0(device const block_q8_0 *xb, short il, thread type4x4 & reg
 
 template <typename type4>
 void dequantize_q8_0_t4(device const block_q8_0 *xb, short il, thread type4 & reg) {
-    device const packed_char4 * qs = (device const packed_char4 *) xb->qs;
+    device const int8_t * qs = ((device const int8_t *)xb->qs);
     const float d = xb->d;
 
-    reg = (type4) (float4(qs[il]) * d);
+    for (int i = 0; i < 4; i++) {
+        reg[i] = (qs[4*(il%4) + i + 16*(il/4)] * d);
+    }
 }
 
 template <typename type4x4>
@@ -4686,15 +4688,14 @@ kernel void kernel_rope_norm(
     float sin_theta;
 
     for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
-        if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
-            const int iw = i0 - args.n_offs; // relative idx
-            const int ic = iw/2;
+        if (i0 < args.n_dims) {
+            const int ic = i0/2;
 
-            const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
+            const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
 
             const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
 
-            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
+            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
 
             device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
             device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + i0*args.nb0);
@@ -4705,10 +4706,6 @@ kernel void kernel_rope_norm(
             dst_data[0] = x0*cos_theta - x1*sin_theta;
             dst_data[1] = x0*sin_theta + x1*cos_theta;
         } else {
-            if (args.inplace) {
-                continue;
-            }
-
             device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
             device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + i0*args.nb0);
 
@@ -4744,18 +4741,17 @@ kernel void kernel_rope_neox(
     float sin_theta;
 
     for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
-        if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
-            const int iw = i0 - args.n_offs; // relative idx
-            const int ic = iw/2;
+        if (i0 < args.n_dims) {
+            const int ic = i0/2;
 
-            const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
+            const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
 
             const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
 
-            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
+            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
 
-            device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + (args.n_offs + ic)*args.nb00);
-            device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + (args.n_offs + ic)*args.nb0);
+            device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + ic*args.nb00);
+            device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + ic*args.nb0);
 
             const float x0 = src[0];
             const float x1 = src[args.n_dims/2];
@@ -4763,10 +4759,6 @@ kernel void kernel_rope_neox(
             dst_data[0]             = x0*cos_theta - x1*sin_theta;
             dst_data[args.n_dims/2] = x0*sin_theta + x1*cos_theta;
         } else {
-            if (args.inplace) {
-                continue;
-            }
-
             device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
             device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + i0*args.nb0);
 
@@ -4801,9 +4793,8 @@ kernel void kernel_rope_multi(
     float sin_theta;
 
     for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
-        if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
-            const int iw = i0 - args.n_offs; // relative idx
-            const int ic = iw/2;
+        if (i0 < args.n_dims) {
+            const int ic = i0/2;
 
             // mrope theta calculations
             // note: the rest is the same as kernel_rope_neox
@@ -4836,14 +4827,14 @@ kernel void kernel_rope_multi(
             }
             // end of mrope
 
-            const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
+            const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
 
             const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
 
-            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
+            rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
 
-            device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + (args.n_offs + ic)*args.nb00);
-            device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + (args.n_offs + ic)*args.nb0);
+            device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + ic*args.nb00);
+            device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + ic*args.nb0);
 
             const float x0 = src[0];
             const float x1 = src[args.n_dims/2];
@@ -4851,10 +4842,6 @@ kernel void kernel_rope_multi(
             dst_data[0]             = x0*cos_theta - x1*sin_theta;
             dst_data[args.n_dims/2] = x0*sin_theta + x1*cos_theta;
         } else {
-            if (args.inplace) {
-                continue;
-            }
-
             device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
             device       T * dst_data  = (device T *)( dst + i3*args.nb3  + i2*args.nb2  + i1*args.nb1  + i0*args.nb0);
 

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

@@ -7376,9 +7376,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
         case GGML_OP_DIAG_MASK_INF:
             return op->ne[3] == 1;
         case GGML_OP_ROPE: {
-            if (((const int32_t *) op->op_params)[15] != 0) {
-                return false; // FIXME: support ggml_rope_set_offset
-            }
             const int mode = ((const int32_t *) op->op_params)[2];
             const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE;
             const bool is_vision = mode == GGML_ROPE_TYPE_VISION;

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

@@ -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;

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

@@ -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) {

ggml/src/ggml-openvino/ggml-openvino.cpp

@@ -1227,10 +1227,6 @@ static bool is_op_unsupported_case(const ggml_tensor * op) {
         const int32_t * op_params = op->op_params;
         const int n_dims = op_params[1];
         const int mode = op_params[2];
-        if (op_params[15] != 0) {
-            // FIXME: support ggml_rope_set_offset
-            return true;
-        }
         if (mode != GGML_ROPE_TYPE_NORMAL && mode != GGML_ROPE_TYPE_NEOX && mode != GGML_ROPE_TYPE_IMROPE) {
             // GGML_LOG_WARN("OpenVINO backend does not support ROPE with mode %d\n", mode);
             return true;

ggml/src/ggml-rpc/ggml-rpc.cpp

@@ -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");

ggml/src/ggml-sycl/fwht.cpp

@@ -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;
-    }
-}

ggml/src/ggml-sycl/fwht.hpp

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

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

@@ -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;
 
@@ -6235,8 +6222,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
         }
         case GGML_OP_ROPE:
         case GGML_OP_ROPE_BACK:
-            // FIXME: support ggml_rope_set_offset
-            return ((const int32_t *) op->op_params)[15] == 0;
         case GGML_OP_IM2COL:
         case GGML_OP_IM2COL_3D:
         case GGML_OP_UPSCALE:

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

@@ -962,7 +962,6 @@ struct vk_device_struct {
     vk_pipeline pipeline_cpy_f32_quant[GGML_TYPE_COUNT];
     vk_pipeline pipeline_cpy_quant_f32[GGML_TYPE_COUNT];
     vk_pipeline pipeline_cpy_transpose_16, pipeline_cpy_transpose_32;
-    vk_pipeline pipeline_cpy_transpose_02_16, pipeline_cpy_transpose_02_32;
     // [src0 0=fp32,1=fp16][dst]
     vk_pipeline pipeline_set_rows_i32[2][GGML_TYPE_COUNT];
     vk_pipeline pipeline_set_rows_i64[2][GGML_TYPE_COUNT];
@@ -1645,7 +1644,6 @@ struct vk_op_rope_push_constants {
     uint32_t rope_mode;
     uint32_t nrows;
     uint32_t n_dims;
-    uint32_t n_offs;
     float freq_scale;
     float freq_base;
     float ext_factor;
@@ -5527,8 +5525,6 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
 
     ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_32, "cpy_transpose_32", cpy_transpose_32_len, cpy_transpose_32_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_16, "cpy_transpose_16", cpy_transpose_16_len, cpy_transpose_16_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
-    ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_02_32, "cpy_transpose_02_32", cpy_transpose_02_32_len, cpy_transpose_02_32_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
-    ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_02_16, "cpy_transpose_02_16", cpy_transpose_02_16_len, cpy_transpose_02_16_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
 
     ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_Q1_0], "cpy_f32_q1_0", cpy_f32_q1_0_len, cpy_f32_q1_0_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_Q2_0], "cpy_f32_q2_0", cpy_f32_q2_0_len, cpy_f32_q2_0_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
@@ -8935,18 +8931,6 @@ static vk_pipeline ggml_vk_get_cpy_pipeline(ggml_backend_vk_context * ctx, const
         }
     }
 
-    // Same, for a 0<->2 swap: src dim2 is the innermost dimension.
-    bool transpose02 = dst && !contig && src->nb[2] == ggml_type_size(to) &&
-                       ggml_is_contiguous(dst) && ggml_are_same_shape(dst, src);
-
-    if (transpose02 && src->type == to) {
-        if (ggml_type_size(to) == 4) {
-            return ctx->device->pipeline_cpy_transpose_02_32;
-        } else if (ggml_type_size(to) == 2) {
-            return ctx->device->pipeline_cpy_transpose_02_16;
-        }
-    }
-
     if (src->type == GGML_TYPE_F32 && to == GGML_TYPE_F32) {
         if (contig) {
             return ctx->device->pipeline_contig_cpy_f32_f32;
@@ -12208,16 +12192,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
                 elements = { ne, 1, 1 };
             }
 
-            if (pipeline == ctx->device->pipeline_cpy_transpose_02_32 ||
-                pipeline == ctx->device->pipeline_cpy_transpose_02_16) {
-                // 32x32 tiles over dims 0 and 2; dim1 and dim3 are the batch
-                elements[0] = (uint32_t)CEIL_DIV(dst->ne[0], 32);
-                elements[1] = (uint32_t)CEIL_DIV(dst->ne[2], 32);
-                elements[2] = (uint32_t)(dst->ne[1]*dst->ne[3]);
-                elements[0] = std::min(elements[0], ctx->device->properties.limits.maxComputeWorkGroupCount[0]);
-                elements[1] = std::min(elements[1], ctx->device->properties.limits.maxComputeWorkGroupCount[1]);
-                elements[2] = std::min(elements[2], ctx->device->properties.limits.maxComputeWorkGroupCount[2]);
-            } else if (pipeline == ctx->device->pipeline_cpy_transpose_32 ||
+            if (pipeline == ctx->device->pipeline_cpy_transpose_32 ||
                 pipeline == ctx->device->pipeline_cpy_transpose_16) {
                 // 32x32 tiles
                 elements[0] = (uint32_t)CEIL_DIV(dst->ne[0], 32);
@@ -13145,7 +13120,6 @@ static uint32_t ggml_vk_rms_partials_size(ggml_backend_vk_context * ctx, const g
 static vk_op_rope_push_constants ggml_vk_make_rope_constants(const ggml_tensor *dst, const ggml_tensor *src0, const bool has_ff, bool backprop, const uint32_t set_rows_stride) {
     const int n_dims        = ((const int32_t *) dst->op_params)[1];
     const int mode          = ((const int32_t *) dst->op_params)[2];
-    const int n_offs        = ((const int32_t *) dst->op_params)[15];
     // const int n_ctx         = ((const int32_t *) dst->op_params)[3];
     const int n_ctx_orig    = ((const int32_t *) dst->op_params)[4];
     const float freq_base   = ((const float *)   dst->op_params)[5];
@@ -13175,7 +13149,7 @@ static vk_op_rope_push_constants ggml_vk_make_rope_constants(const ggml_tensor *
     uint32_t nb13 = dst->nb[3] / ggml_type_size(dst->type);
 
     vk_op_rope_push_constants rope {
-        (uint32_t)mode, (uint32_t)ggml_nrows(src0), (uint32_t)n_dims, (uint32_t)n_offs, freq_scale,
+        (uint32_t)mode, (uint32_t)ggml_nrows(src0), (uint32_t)n_dims, freq_scale,
         freq_base, ext_factor, attn_factor, {corr_dims[0], corr_dims[1]}, theta_scale, has_ff,
         { sections[0], sections[1], sections[2], sections[3] }, is_imrope, backprop, set_rows_stride,
 
@@ -19221,10 +19195,6 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
                     tensor_clone = ggml_rope_ext_back(ggml_ctx, src_clone[0], src_clone[1], src_clone[2], n_dims, mode, n_ctx_orig_ggml, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
                 }
             }
-            const int n_offs = ((int32_t *) tensor->op_params)[15];
-            if (n_offs != 0) {
-                tensor_clone = ggml_rope_set_offset(tensor_clone, n_offs);
-            }
         } else if (tensor->op == GGML_OP_UNARY) {
             switch (ggml_get_unary_op(tensor)) {
             case GGML_UNARY_OP_EXP:

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

@@ -1,61 +0,0 @@
-#version 450
-
-#include "types.glsl"
-#include "generic_unary_head.glsl"
-
-// workgroup does 32x32 tile, but uses 32x8 threads
-#define TILE_DIM 32
-layout(local_size_x = 32, local_size_y = 8, local_size_z = 1) in;
-
-// +1 padding avoids shared-memory bank conflicts on the transposed read
-shared uint sh[TILE_DIM][TILE_DIM + 1];
-
-void iter(uvec3 wg_id) {
-    const uint tile_i0 = wg_id.x;   // tiles dst ne10 (== src ne00)
-    const uint tile_i2 = wg_id.y;   // tiles dst ne12 (== src ne02)
-
-    const uint tid_col = gl_LocalInvocationID.x;
-    const uint tid_row = gl_LocalInvocationID.y;
-
-    const uint i1 = wg_id.z % p.ne11;
-    const uint i3 = wg_id.z / p.ne11;
-    const uint i01 = i1;
-    const uint i03 = i3;
-
-    [[unroll]] for (uint y = 0; y < 4; ++y) {
-        const uint i00 = tile_i0 * TILE_DIM + tid_row + 8 * y;
-        const uint i02 = tile_i2 * TILE_DIM + tid_col;
-        if (i00 < p.ne00 && i01 < p.ne01 && i02 < p.ne02 && i03 < p.ne03) {
-            const uint src_idx = i00 * p.nb00 + i01 * p.nb01 + i02 * p.nb02 + i03 * p.nb03;
-            sh[tid_row + 8 * y][tid_col] = uint(data_a[get_aoffset() + src_idx]);
-        }
-    }
-
-    barrier();
-
-    [[unroll]] for (uint y = 0; y < 4; ++y) {
-        const uint i0 = tile_i0 * TILE_DIM + tid_col;
-        const uint i2 = tile_i2 * TILE_DIM + tid_row + 8 * y;
-        if (i0 < p.ne10 && i1 < p.ne11 && i2 < p.ne12 && i3 < p.ne13) {
-            const uint dst_idx = i0 * p.nb10 + i1 * p.nb11 + i2 * p.nb12 + i3 * p.nb13;
-            data_d[get_doffset() + dst_idx] = D_TYPE(sh[tid_col][tid_row + 8 * y]);
-        }
-    }
-}
-
-#define CEIL_DIV(a, b) (((a) + (b) - 1) / (b))
-
-void main() {
-    bool need_barrier = false;
-    for (uint z = gl_WorkGroupID.z; z < p.ne11 * p.ne13; z += gl_NumWorkGroups.z) {
-        for (uint y = gl_WorkGroupID.y; y < CEIL_DIV(p.ne12, TILE_DIM); y += gl_NumWorkGroups.y) {
-            for (uint x = gl_WorkGroupID.x; x < CEIL_DIV(p.ne10, TILE_DIM); x += gl_NumWorkGroups.x) {
-                if (need_barrier) {
-                    barrier();
-                }
-                need_barrier = true;
-                iter(uvec3(x, y, z));
-            }
-        }
-    }
-}

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

@@ -50,21 +50,19 @@ void rope_norm(const uint i0, const uint i1, const uint i2, const uint i3, rope_
     }
     idst += p.d_offset;
 
-    if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
+    if (i0 >= p.n_dims) {
         rope_data_d[idst + 0] = ROPE_D_TYPE(rope_data_a[ix + 0]);
         rope_data_d[idst + 1] = ROPE_D_TYPE(rope_data_a[ix + 1]);
 
         return;
     }
 
-    const uint iw = i0 - p.n_offs; // relative idx
+    const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, i0/2.0f);
 
-    const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, iw/2.0f);
-
-    const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
+    const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
 
     float cos_theta, sin_theta;
-    rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
+    rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
 
     const float x0 = float(rope_data_a[ix + 0]);
     const float x1 = float(rope_data_a[ix + 1]);
@@ -89,28 +87,25 @@ void rope_neox(const uint i0, const uint i1, const uint i2, const uint i3, rope_
     }
     idst += p.d_offset;
 
-    if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
+    if (i0 >= p.n_dims) {
         rope_data_d[idst + i0/2 + 0] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 0]);
         rope_data_d[idst + i0/2 + 1] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 1]);
 
         return;
     }
 
-    const uint iw = i0 - p.n_offs; // relative idx
+    const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, i0/2.0f);
 
-    const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, iw/2.0f);
-
-    const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
+    const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
 
     float cos_theta, sin_theta;
-    rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
+    rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
 
-    // idst/ix point at channel i0/2; the first channel of the rotated pair is p.n_offs + iw/2 = i0/2 + p.n_offs/2
-    const float x0 = float(rope_data_a[ix + p.n_offs/2 + 0]);
-    const float x1 = float(rope_data_a[ix + p.n_offs/2 + p.n_dims/2]);
+    const float x0 = float(rope_data_a[ix + 0]);
+    const float x1 = float(rope_data_a[ix + p.n_dims/2]);
 
-    rope_data_d[idst + p.n_offs/2 + 0]          = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
-    rope_data_d[idst + p.n_offs/2 + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
+    rope_data_d[idst + 0]          = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
+    rope_data_d[idst + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
 }
 
 
@@ -130,56 +125,53 @@ void rope_multi(const uint i0, const uint i1, const uint i2, const uint i3, rope
     }
     idst += p.d_offset;
 
-    if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
+    if (i0 >= p.n_dims) {
         rope_data_d[idst + i0/2 + 0] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 0]);
         rope_data_d[idst + i0/2 + 1] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 1]);
 
         return;
     }
 
-    const uint iw = i0 - p.n_offs; // relative idx
-
     const int sect_dims = p.sections[0] + p.sections[1] + p.sections[2] + p.sections[3];
     const int sec_w = p.sections[1] + p.sections[0];
-    const uint sector = (iw / 2) % sect_dims;
+    const uint sector = (i0 / 2) % sect_dims;
 
     float theta_base = 0.0;
     if (p.is_imrope != 0) {
         if (sector % 3 == 1 && sector < 3 * p.sections[1]) {
-            theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, i0/2.0f);
         } else if (sector % 3 == 2 && sector < 3 * p.sections[2]) {
-            theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, i0/2.0f);
         } else if (sector % 3 == 0 && sector < 3 * p.sections[0]) {
-            theta_base = rope_data_pos[i2]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2]*pow(p.theta_scale, i0/2.0f);
         } else {
-            theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, i0/2.0f);
         }
     } else {
         if (sector < p.sections[0]) {
-            theta_base = rope_data_pos[i2]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2]*pow(p.theta_scale, i0/2.0f);
         }
         else if (sector >= p.sections[0] && sector < sec_w) {
-            theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, i0/2.0f);
         }
         else if (sector >= sec_w && sector < sec_w + p.sections[2]) {
-            theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, i0/2.0f);
         }
         else if (sector >= sec_w + p.sections[2]) {
-            theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, iw/2.0f);
+            theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, i0/2.0f);
         }
     }
 
-    const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
+    const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
 
     float cos_theta, sin_theta;
-    rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
+    rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
 
-    // idst/ix point at channel i0/2; the first channel of the rotated pair is p.n_offs + iw/2 = i0/2 + p.n_offs/2
-    const float x0 = float(rope_data_a[ix + p.n_offs/2 + 0]);
-    const float x1 = float(rope_data_a[ix + p.n_offs/2 + p.n_dims/2]);
+    const float x0 = float(rope_data_a[ix + 0]);
+    const float x1 = float(rope_data_a[ix + p.n_dims/2]);
 
-    rope_data_d[idst + p.n_offs/2 + 0]          = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
-    rope_data_d[idst + p.n_offs/2 + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
+    rope_data_d[idst + 0]          = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
+    rope_data_d[idst + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
 }
 
 void rope_vision(const uint i0, const uint i1, const uint i2, const uint i3, rope_params p) {

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

@@ -5,7 +5,6 @@ struct rope_params {
     uint rope_mode;
     uint nrows;
     uint n_dims;
-    uint n_offs;
     float freq_scale;
     float freq_base;
     float ext_factor;

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

@@ -826,8 +826,6 @@ void process_shaders() {
 
     string_to_spv("cpy_transpose_16", "copy_transpose.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});
     string_to_spv("cpy_transpose_32", "copy_transpose.comp", {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}});
-    string_to_spv("cpy_transpose_02_16", "copy_transpose_02.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});
-    string_to_spv("cpy_transpose_02_32", "copy_transpose_02.comp", {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}});
 
     for (std::string t : {"q1_0", "q2_0", "q4_0", "q4_1", "q5_0", "q5_1", "q8_0", "iq4_nl"}) {
         string_to_spv("cpy_f32_" + t, "copy_to_quant.comp", {{"DATA_A_" + to_uppercase(t), "1"}, {"S_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});

ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp

@@ -954,11 +954,10 @@ struct ggml_webgpu_mul_mat_vec_pipeline_key {
     int       vectorized;
     uint32_t  num_cols;
     bool      use_mmvq;
-    bool      src_overlap;
 
     bool operator==(const ggml_webgpu_mul_mat_vec_pipeline_key & other) const {
         return src0_type == other.src0_type && src1_type == other.src1_type && vectorized == other.vectorized &&
-               num_cols == other.num_cols && use_mmvq == other.use_mmvq && src_overlap == other.src_overlap;
+               num_cols == other.num_cols && use_mmvq == other.use_mmvq;
     }
 };
 
@@ -970,7 +969,6 @@ struct ggml_webgpu_mul_mat_vec_pipeline_key_hash {
         ggml_webgpu_hash_combine(seed, key.vectorized);
         ggml_webgpu_hash_combine(seed, key.num_cols);
         ggml_webgpu_hash_combine(seed, key.use_mmvq);
-        ggml_webgpu_hash_combine(seed, key.src_overlap);
         return seed;
     }
 };
@@ -979,7 +977,6 @@ struct ggml_webgpu_mul_mat_vec_shader_decisions {
     uint32_t wg_size;
     uint32_t outputs_per_wg;
     uint32_t vec_size;
-    bool     src_overlap = false;
 };
 
 struct ggml_webgpu_quantize_q8_pipeline_key {
@@ -1001,11 +998,10 @@ struct ggml_webgpu_mul_mat_pipeline_key {
     ggml_type src1_type;
     int       vectorized;
     int       use_subgroup_matrix;
-    bool      src_overlap;
 
     bool operator==(const ggml_webgpu_mul_mat_pipeline_key & other) const {
         return src0_type == other.src0_type && src1_type == other.src1_type && vectorized == other.vectorized &&
-               use_subgroup_matrix == other.use_subgroup_matrix && src_overlap == other.src_overlap;
+               use_subgroup_matrix == other.use_subgroup_matrix;
     }
 };
 
@@ -1016,7 +1012,6 @@ struct ggml_webgpu_mul_mat_pipeline_key_hash {
         ggml_webgpu_hash_combine(seed, key.src1_type);
         ggml_webgpu_hash_combine(seed, key.vectorized);
         ggml_webgpu_hash_combine(seed, key.use_subgroup_matrix);
-        ggml_webgpu_hash_combine(seed, key.src_overlap);
         return seed;
     }
 };
@@ -1039,7 +1034,6 @@ struct ggml_webgpu_mul_mat_shader_decisions {
     uint32_t subgroup_matrix_n;
 
     uint32_t mul_mat_wg_size;
-    bool     src_overlap = false;
 };
 
 /** MUL_MAT_ID **/
@@ -1956,7 +1950,7 @@ class ggml_webgpu_shader_lib {
         return quantize_q8_pipelines[key];
     }
 
-    webgpu_pipeline get_mul_mat_vec_pipeline(const ggml_webgpu_shader_lib_context & context, bool src_overlap) {
+    webgpu_pipeline get_mul_mat_vec_pipeline(const ggml_webgpu_shader_lib_context & context) {
         ggml_webgpu_mul_mat_vec_pipeline_key key = {};
         key.src0_type                            = context.src0->type;
         key.src1_type                            = context.src1->type;
@@ -1967,7 +1961,6 @@ class ggml_webgpu_shader_lib {
         key.num_cols   = context.dst->ne[1];
         key.use_mmvq =
             ggml_webgpu_can_use_mmvq(context.src0, context.src1, context.supports_dot_product, context.vendor);
-        key.src_overlap = src_overlap;
 
         auto it = mul_mat_vec_pipelines.find(key);
         if (it != mul_mat_vec_pipelines.end()) {
@@ -2075,11 +2068,6 @@ class ggml_webgpu_shader_lib {
             defines.push_back("Q8_1_T");
         }
 
-        if (key.src_overlap) {
-            defines.push_back("SRC_OVERLAP");
-            variant += "_src_overlap";
-        }
-
         defines.push_back(std::string("WG_SIZE=") + std::to_string(wg_size));
         defines.push_back(std::string("OUTPUTS_PER_WG=") + std::to_string(outputs_per_wg));
         defines.push_back(context.supports_subgroups ? "USE_SUBGROUP_REDUCTION" : "USE_WORKGROUP_REDUCTION");
@@ -2101,7 +2089,7 @@ class ggml_webgpu_shader_lib {
         return mul_mat_vec_pipelines[key];
     }
 
-    webgpu_pipeline get_mul_mat_fast_pipeline(const ggml_webgpu_shader_lib_context & context, bool src_overlap) {
+    webgpu_pipeline get_mul_mat_fast_pipeline(const ggml_webgpu_shader_lib_context & context) {
         ggml_webgpu_mul_mat_pipeline_key key = {};
         key.src0_type                        = context.src0->type;
         key.src1_type                        = context.src1->type;
@@ -2110,7 +2098,6 @@ class ggml_webgpu_shader_lib {
                                       1 :
                                       0;
         key.use_subgroup_matrix = context.supports_subgroup_matrix;
-        key.src_overlap         = src_overlap;
 
         auto it = mul_mat_fast_pipelines.find(key);
         if (it != mul_mat_fast_pipelines.end()) {
@@ -2229,11 +2216,6 @@ class ggml_webgpu_shader_lib {
             variant += "_vectorized";
         }
 
-        if (key.src_overlap) {
-            defines.push_back("SRC_OVERLAP");
-            variant += "_src_overlap";
-        }
-
         if (!key.use_subgroup_matrix) {
             defines.push_back("WORKGROUP_SIZE_M=" + std::to_string(WEBGPU_MUL_MAT_WG_SIZE_M) + "u");
             defines.push_back("WORKGROUP_SIZE_N=" + std::to_string(WEBGPU_MUL_MAT_WG_SIZE_N) + "u");

ggml/src/ggml-webgpu/ggml-webgpu.cpp

@@ -1628,65 +1628,48 @@ static webgpu_encoded_op ggml_webgpu_mul_mat(webgpu_context & ctx,
     // Get or create pipeline
     webgpu_pipeline                   pipeline;
     std::vector<webgpu_dispatch_desc> dispatches;
-    const bool src_overlap = ggml_webgpu_tensor_binding_overlap(ctx->global_ctx, src0, src1) && !use_mmvq;
 
     if (use_mat_vec) {
         if (use_mmvq) {
             ggml_webgpu_quantize_q8_dispatch(ctx, src0, src1, dst, dispatches);
         }
-        pipeline = ctx->shader_lib->get_mul_mat_vec_pipeline(shader_lib_ctx, src_overlap);
+        pipeline = ctx->shader_lib->get_mul_mat_vec_pipeline(shader_lib_ctx);
     } else {
-        pipeline = ctx->shader_lib->get_mul_mat_fast_pipeline(shader_lib_ctx, src_overlap);
-    }
-
-    uint32_t offset_src0   = (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src0) / ggml_type_size(src0->type));
-    uint32_t offset_src1   = (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src1) / ggml_type_size(src1->type));
-    size_t   merged_offset = 0;
-    size_t   merged_size   = 0;
-    if (src_overlap) {
-        const ggml_webgpu_merged_binding_range merged_range =
-            ggml_webgpu_tensor_merged_binding_range(ctx, { src0, src1 });
-        merged_offset = merged_range.offset;
-        merged_size   = merged_range.size;
-        offset_src0   = ggml_webgpu_tensor_merged_element_offset(src0, merged_range);
-        offset_src1   = ggml_webgpu_tensor_merged_element_offset(src1, merged_range);
+        pipeline = ctx->shader_lib->get_mul_mat_fast_pipeline(shader_lib_ctx);
     }
 
     // Build params
-    std::vector<uint32_t> params = { offset_src0,
-                                     offset_src1,
-                                     (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
-                                     (uint32_t) dst->ne[0],
-                                     (uint32_t) dst->ne[1],
-                                     (uint32_t) src0->ne[0],
-                                     (uint32_t) (src0->nb[1] / ggml_type_size(src0->type)),
-                                     (uint32_t) (src1->nb[1] / ggml_type_size(src1->type)),
-                                     (uint32_t) (src0->nb[2] / ggml_type_size(src0->type)),
-                                     (uint32_t) (src1->nb[2] / ggml_type_size(src1->type)),
-                                     (uint32_t) (src0->nb[3] / ggml_type_size(src0->type)),
-                                     (uint32_t) (src1->nb[3] / ggml_type_size(src1->type)),
-                                     (uint32_t) src0->ne[2],
-                                     (uint32_t) src0->ne[3],
-                                     (uint32_t) (src1->ne[2] / src0->ne[2]),
-                                     (uint32_t) (src1->ne[3] / src0->ne[3]) };
+    std::vector<uint32_t> params = {
+        (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src0) / ggml_type_size(src0->type)),
+        (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src1) / ggml_type_size(src1->type)),
+        (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
+        (uint32_t) dst->ne[0],
+        (uint32_t) dst->ne[1],
+        (uint32_t) src0->ne[0],
+        (uint32_t) (src0->nb[1] / ggml_type_size(src0->type)),
+        (uint32_t) (src1->nb[1] / ggml_type_size(src1->type)),
+        (uint32_t) (src0->nb[2] / ggml_type_size(src0->type)),
+        (uint32_t) (src1->nb[2] / ggml_type_size(src1->type)),
+        (uint32_t) (src0->nb[3] / ggml_type_size(src0->type)),
+        (uint32_t) (src1->nb[3] / ggml_type_size(src1->type)),
+        (uint32_t) src0->ne[2],
+        (uint32_t) src0->ne[3],
+        (uint32_t) (src1->ne[2] / src0->ne[2]),
+        (uint32_t) (src1->ne[3] / src0->ne[3])
+    };
 
     // Build bind group entries
     std::vector<wgpu::BindGroupEntry> entries = {};
+
+    entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
     if (use_mmvq) {
-        entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
         auto & mmvq_qq8_entry = dispatches[0].bind_group_entries[1];
         entries.push_back(ggml_webgpu_make_bind_group_entry(1, ggml_webgpu_tensor_buf(dst), mmvq_qq8_entry.offset,
                                                             mmvq_qq8_entry.size));
-        entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
-    } else if (src_overlap) {
-        entries.push_back(
-            ggml_webgpu_make_bind_group_entry(0, ggml_webgpu_tensor_buf(src0), merged_offset, merged_size));
-        entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 1, dst));
     } else {
-        entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
         entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 1, src1));
-        entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
     }
+    entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
 
     // Calculate workgroup dimensions
     uint32_t       wg_x           = 1;
@@ -4472,9 +4455,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
             supports_op = (op->type == GGML_TYPE_F32 && src0->type == GGML_TYPE_F32) && ggml_is_contiguous_rows(src0);
             break;
         case GGML_OP_ROPE:
-            // FIXME: support ggml_rope_set_offset
-            supports_op =
-                (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) && ((const int32_t *) op->op_params)[15] == 0;
+            supports_op = op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16;
             break;
         case GGML_OP_GLU:
             switch (ggml_get_glu_op(op)) {

ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl

@@ -1,7 +1,3 @@
-#ifndef SRC0
-#define SRC0 src0
-#endif
-
 #ifdef BYTE_HELPERS
 fn get_byte(value: u32, index: u32) -> u32 {
     return (value >> (index * 8)) & 0xFF;
@@ -50,7 +46,7 @@ fn load_f16_as_f32_at_src(byte_offset: u32) -> f32 {
 
 #ifdef DECLARE_BYTE_LOADERS_SRC0
 fn load_u16_at_src0(byte_offset: u32) -> u32 {
-    let word = SRC0[byte_offset / 4u];
+    let word = src0[byte_offset / 4u];
     let shift = (byte_offset & 0x2u) * 8u;
     return (word >> shift) & 0xFFFFu;
 }
@@ -59,14 +55,14 @@ fn load_u16_at_src0(byte_offset: u32) -> u32 {
 // Caller extracts the 16-bit half it needs via & 0xFFFFu or >> 16u.
 // this is used in k-quants for better performance
 fn load_u32_at_src0_aligned(byte_offset: u32) -> u32 {
-    return SRC0[(byte_offset & ~3u) / 4u];
+    return src0[(byte_offset & ~3u) / 4u];
 }
 
 fn load_u32_at_src0(byte_offset: u32) -> u32 {
     let word_idx = byte_offset / 4u;
     let shift = (byte_offset & 0x3u) * 8u;
-    let lo = SRC0[word_idx];
-    let hi = SRC0[word_idx + 1u];
+    let lo = src0[word_idx];
+    let hi = src0[word_idx + 1u];
     let shifted = (lo >> shift) | (hi << (32u - shift));
     return select(shifted, lo, shift == 0u);
 }
@@ -77,7 +73,7 @@ fn load_f16_at_src0(byte_offset: u32) -> f16 {
 }
 
 fn load_f16_as_f32_at_src0(byte_offset: u32) -> f32 {
-    let word = SRC0[byte_offset / 4u];
+    let word = src0[byte_offset / 4u];
     let shift = (byte_offset & 0x2u) * 8u;
     let d_bits = (word >> shift) & 0xFFFFu;
     return unpack2x16float(d_bits)[0];

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl

@@ -1,10 +1,3 @@
-#ifndef SRC0
-#define SRC0 src0
-#endif
-#ifndef SRC1
-#define SRC1 src1
-#endif
-
 #ifdef VEC
 #define VEC_SIZE 4
 #define SHMEM_TYPE vec4<f16>
@@ -46,7 +39,7 @@ fn init_shmem_src0(thread_id: u32, batch_offset: u32, offset_m: u32, k_outer: u3
         let src0_idx = batch_offset + global_m * params.stride_01 + global_k;
         let src0_val = select( // taking a slight performance hit to avoid oob
             SRC0_TYPE(0.0),
-            SRC0[src0_idx/VEC_SIZE],
+            src0[src0_idx/VEC_SIZE],
             global_m < params.m && global_k < params.k);
         store_shmem(SHMEM_TYPE(src0_val), elem_idx);
     }
@@ -64,7 +57,7 @@ fn init_shmem_src1(thread_id: u32, batch_offset: u32, offset_n: u32, k_outer: u3
         let src1_idx = batch_offset + global_n * params.stride_11 + global_k;
         let src1_val = select(
             SRC1_TYPE(0.0),
-            SRC1[src1_idx/VEC_SIZE],
+            src1[src1_idx/VEC_SIZE],
             global_n < params.n && global_k < params.k);
         store_shmem(SHMEM_TYPE(src1_val), TILE_SRC0_SHMEM + elem_idx);
     }

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_reg_tile.wgsl

@@ -1,12 +1,8 @@
 enable f16;
 
 #define DECLARE_BYTE_LOADERS_SRC0
-
-#ifdef SRC_OVERLAP
-#define SRC0 merged_src
-#define SRC1 merged_src
-#endif
 #include "common_decls.tmpl"
+
 #include "mul_mat_decls.tmpl"
 
 #ifdef VEC
@@ -40,17 +36,11 @@ struct MulMatParams {
     broadcast3: u32
 };
 
-#ifdef SRC_OVERLAP
-@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
-#define DST_BINDING 1
-#else
 @group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>; // M rows, K columns
 @group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>; // K rows, N columns (transposed)
-#define DST_BINDING 2
-#endif
+@group(0) @binding(2) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
 
-@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
-@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
+@group(0) @binding(3) var<uniform> params: MulMatParams;
 
 fn get_local_n(thread_id: u32) -> u32 {
     return thread_id / WORKGROUP_SIZE_M;

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_subgroup_matrix.wgsl

@@ -4,10 +4,6 @@ enable subgroups;
 enable chromium_experimental_subgroup_matrix;
 
 #define DECLARE_BYTE_LOADERS_SRC0
-#ifdef SRC_OVERLAP
-#define SRC0 merged_src
-#define SRC1 merged_src
-#endif
 #include "common_decls.tmpl"
 
 #include "mul_mat_decls.tmpl"
@@ -52,17 +48,11 @@ struct MulMatParams {
 };
 
 // SRC0_TYPE and SRC1_TYPE are defined in mul_mat_decls, which is included
-#ifdef SRC_OVERLAP
-@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
-#define DST_BINDING 1
-#else
 @group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>; // M rows, K columns
 @group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>; // K rows, N columns (transposed)
-#define DST_BINDING 2
-#endif
+@group(0) @binding(2) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
 
-@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
-@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
+@group(0) @binding(3) var<uniform> params: MulMatParams;
 
 const WG_M_SG_TILE_SIZE = SUBGROUP_M * SUBGROUP_MATRIX_M * SUBGROUP_MATRIX_M_SIZE;
 const WG_N_SG_TILE_SIZE = SUBGROUP_N * SUBGROUP_MATRIX_N * SUBGROUP_MATRIX_N_SIZE;

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec.wgsl

@@ -7,11 +7,6 @@ enable f16;
 requires packed_4x8_integer_dot_product;
 #endif
 
-#ifdef SRC_OVERLAP
-#define SRC0 merged_src
-#define SRC1 merged_src
-#endif
-
 #define DECLARE_BYTE_LOADERS_SRC0
 #include "common_decls.tmpl"
 
@@ -40,22 +35,17 @@ struct MulMatParams {
     broadcast3: u32
 };
 
-#if defined(MMVQ)
 @group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>;
+
+#ifdef MMVQ
 @group(0) @binding(1) var<storage, read_write> src1q: array<q8_1>;
-#define DST_BINDING 2
-#elif defined(SRC_OVERLAP)
-@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
-#define DST_BINDING 1
 #else
-@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>;
 @group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>;
-#define DST_BINDING 2
 #endif
 
-@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<f32>;
+@group(0) @binding(2) var<storage, read_write> dst: array<f32>;
 // "mul_mat_vec_acc.tmpl" requires params.k, params.m, params.stride_01
-@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
+@group(0) @binding(3) var<uniform> params: MulMatParams;
 
 // Flattened as [row][thread] to keep each row's reduction contiguous in memory.
 var<workgroup> partial_sums: array<f32, OUTPUTS_PER_WG * WG_SIZE>;

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl

@@ -1,10 +1,3 @@
-#ifndef SRC0
-#define SRC0 src0
-#endif
-#ifndef SRC1
-#define SRC1 src1
-#endif
-
 #ifdef U32_DEQUANT_HELPERS
 #define SRC0_TYPE u32
 
@@ -50,13 +43,13 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
     for (var k = thread_id; k < k_vec; k += WG_SIZE) {
         var x_vals: array<SRC1_TYPE, NUM_COLS>;
         for (var col = 0u;col < NUM_COLS;col += 1) {
-            x_vals[col] = SRC1[src1_idx_base_vec + col * (params.stride_11 / VEC_SIZE) + k];
+            x_vals[col] = src1[src1_idx_base_vec + col * (params.stride_11 / VEC_SIZE) + k];
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
             let output_row = row_base + row;
             if (output_row < params.m) {
                 let src0_idx = (src0_batch_offset + output_row * params.stride_01) / VEC_SIZE + k;
-                let w = SRC0[src0_idx];
+                let w = src0[src0_idx];
                 for (var col = 0u;col < NUM_COLS;col += 1) {
                     acc[col][row] += inner_dot(w, x_vals[col]);
                 }
@@ -83,7 +76,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -123,8 +116,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -167,8 +160,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -212,8 +205,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -260,8 +253,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -309,7 +302,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -354,7 +347,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -416,10 +409,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 4u; i++) {
-                x_block[col][i]       = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4u]  = f32(SRC1[x_base + col * params.stride_11 + 32u + i]);
-                x_block[col][i + 8u]  = f32(SRC1[x_base + col * params.stride_11 + 64u + i]);
-                x_block[col][i + 12u] = f32(SRC1[x_base + col * params.stride_11 + 96u + i]);
+                x_block[col][i]       = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4u]  = f32(src1[x_base + col * params.stride_11 + 32u + i]);
+                x_block[col][i + 8u]  = f32(src1[x_base + col * params.stride_11 + 64u + i]);
+                x_block[col][i + 12u] = f32(src1[x_base + col * params.stride_11 + 96u + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -525,8 +518,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 8u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 8u] = f32(SRC1[x_base + col * params.stride_11 + 32u + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 8u] = f32(src1[x_base + col * params.stride_11 + 32u + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -617,10 +610,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         for (var col = 0u; col < NUM_COLS;col += 1) {
             let col_base = x_base + col * params.stride_11;
             for (var i = 0u; i < 4u; i++) {
-                x_block[col][i]       = f32(SRC1[col_base + i]);
-                x_block[col][i + 4u]  = f32(SRC1[col_base + 32u + i]);
-                x_block[col][i + 8u]  = f32(SRC1[col_base + 128u + i]);
-                x_block[col][i + 12u] = f32(SRC1[col_base + 160u + i]);
+                x_block[col][i]       = f32(src1[col_base + i]);
+                x_block[col][i + 4u]  = f32(src1[col_base + 32u + i]);
+                x_block[col][i + 8u]  = f32(src1[col_base + 128u + i]);
+                x_block[col][i + 12u] = f32(src1[col_base + 160u + i]);
             }
         }
 
@@ -720,10 +713,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         for (var col = 0u; col < NUM_COLS;col += 1) {
             let col_base = x_base + col * params.stride_11;
             for (var i = 0u; i < 4u; i++) {
-                x_block[col][i]       = f32(SRC1[col_base + i]);
-                x_block[col][i + 4u]  = f32(SRC1[col_base + 32u + i]);
-                x_block[col][i + 8u]  = f32(SRC1[col_base + 128u + i]);
-                x_block[col][i + 12u] = f32(SRC1[col_base + 160u + i]);
+                x_block[col][i]       = f32(src1[col_base + i]);
+                x_block[col][i + 4u]  = f32(src1[col_base + 32u + i]);
+                x_block[col][i + 8u]  = f32(src1[col_base + 128u + i]);
+                x_block[col][i + 12u] = f32(src1[col_base + 160u + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -830,10 +823,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         for (var col = 0u; col < NUM_COLS;col += 1) {
             let col_base = x_base + col * params.stride_11;
             for (var l = 0u; l < 4u; l++) {
-                x_block[col][l]       = f32(SRC1[col_base + l]);
-                x_block[col][l + 4u]  = f32(SRC1[col_base + 32u + l]);
-                x_block[col][l + 8u]  = f32(SRC1[col_base + 64u + l]);
-                x_block[col][l + 12u] = f32(SRC1[col_base + 96u + l]);
+                x_block[col][l]       = f32(src1[col_base + l]);
+                x_block[col][l + 4u]  = f32(src1[col_base + 32u + l]);
+                x_block[col][l + 8u]  = f32(src1[col_base + 64u + l]);
+                x_block[col][l + 12u] = f32(src1[col_base + 96u + l]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -906,7 +899,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -967,7 +960,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1046,7 +1039,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1108,7 +1101,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1175,7 +1168,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1241,7 +1234,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1309,7 +1302,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1374,8 +1367,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2u; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4u] = f32(SRC1[x_base + col * params.stride_11 + i + 16u]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4u] = f32(src1[x_base + col * params.stride_11 + i + 16u]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1425,7 +1418,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, 16>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < 16u; i++) {
-                x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
+                x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1483,8 +1476,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1528,8 +1521,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
         var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
         for (var col = 0u; col < NUM_COLS;col += 1) {
             for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
-                x_block[col][i]     = f32(SRC1[x_base + col * params.stride_11 + i]);
-                x_block[col][i + 8] = f32(SRC1[x_base + col * params.stride_11 + i + 8]);
+                x_block[col][i]     = f32(src1[x_base + col * params.stride_11 + i]);
+                x_block[col][i + 8] = f32(src1[x_base + col * params.stride_11 + i + 8]);
             }
         }
         for (var row = 0u; row < OUTPUTS_PER_WG; row++) {

ggml/src/ggml.c

@@ -4200,7 +4200,7 @@ static struct ggml_tensor * ggml_rope_impl(
 
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
-    int32_t params[16] = { /*n_past*/ 0, n_dims, mode, /*n_ctx*/ 0, n_ctx_orig };
+    int32_t params[15] = { /*n_past*/ 0, n_dims, mode, /*n_ctx*/ 0, n_ctx_orig };
     memcpy(params +  5, &freq_base,    sizeof(float));
     memcpy(params +  6, &freq_scale,   sizeof(float));
     memcpy(params +  7, &ext_factor,   sizeof(float));
@@ -4212,8 +4212,6 @@ static struct ggml_tensor * ggml_rope_impl(
     } else {
         memset(params + 11, 0,         sizeof(int32_t) * GGML_MROPE_SECTIONS);
     }
-    params[15] = 0; // n_offs, set via ggml_rope_set_offset()
-
     ggml_set_op_params(result, params, sizeof(params));
 
     result->op     = GGML_OP_ROPE;
@@ -4424,20 +4422,6 @@ struct ggml_tensor * ggml_rope_multi_back(
     result->op = GGML_OP_ROPE_BACK;
     return result;
 }
-
-struct ggml_tensor * ggml_rope_set_offset(
-        struct ggml_tensor  * a,
-        int                   n_offs) {
-    GGML_ASSERT(a->op == GGML_OP_ROPE || a->op == GGML_OP_ROPE_BACK);
-    GGML_ASSERT(n_offs >= 0);
-
-    const int32_t mode = ggml_get_op_params_i32(a, 2);
-    GGML_ASSERT(mode != GGML_ROPE_TYPE_VISION);
-
-    ggml_set_op_params_i32(a, 15, n_offs);
-    return a;
-}
-
 // ggml_clamp
 
 struct ggml_tensor * ggml_clamp(

gguf-py/gguf/gguf_reader.py

@@ -32,10 +32,6 @@ from gguf.constants import (
     GGUFEndian,
 )
 
-# limits mirroring ggml/src/gguf.cpp (not part of gguf.h)
-GGUF_MAX_STRING_LENGTH  = 1024 * 1024 * 1024
-GGUF_MAX_ARRAY_ELEMENTS = 1024 * 1024 * 1024
-
 logger = logging.getLogger(__name__)
 
 READER_SUPPORTED_VERSIONS = [2, GGUF_VERSION]
@@ -171,10 +167,6 @@ class GGUFReader:
         offs += self._push_field(ReaderField(offs, 'GGUF.tensor_count', [temp_counts[:1]], [0], [GGUFValueType.UINT64]))
         offs += self._push_field(ReaderField(offs, 'GGUF.kv_count', [temp_counts[1:]], [0], [GGUFValueType.UINT64]))
         tensor_count, kv_count = temp_counts
-        if tensor_count > GGUF_MAX_ARRAY_ELEMENTS:
-            raise ValueError(f'Tensor count {tensor_count} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
-        if kv_count > GGUF_MAX_ARRAY_ELEMENTS:
-            raise ValueError(f'KV count {kv_count} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
         offs = self._build_fields(offs, kv_count)
 
         # Build Tensor Info Fields
@@ -225,10 +217,6 @@ class GGUFReader:
 
     def _get_str(self, offset: int) -> tuple[npt.NDArray[np.uint64], npt.NDArray[np.uint8]]:
         slen = self._get(offset, np.uint64)
-        if int(slen[0]) > GGUF_MAX_STRING_LENGTH:
-            raise ValueError(f'String length {int(slen[0])} exceeds maximum {GGUF_MAX_STRING_LENGTH}')
-        if offset + 8 + int(slen[0]) > self.data.nbytes:
-            raise ValueError(f'String length {int(slen[0])} exceeds remaining file size {self.data.nbytes - offset - 8}')
         return slen, self._get(offset + 8, np.uint8, slen[0])
 
     def _get_field_parts(
@@ -253,8 +241,6 @@ class GGUFReader:
             raw_itype = self._get(offs, np.uint32)
             offs += int(raw_itype.nbytes)
             alen = self._get(offs, np.uint64)
-            if int(alen[0]) > GGUF_MAX_ARRAY_ELEMENTS:
-                raise ValueError(f'Array length {int(alen[0])} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
             offs += int(alen.nbytes)
             aparts: list[npt.NDArray[Any]] = [raw_itype, alen]
             data_idxs: list[int] = []

scripts/make-release-desc.sh

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

scripts/sync-ggml.last

@@ -1 +1 @@
-8c63e70982c95ceb862e3a1073a2c1beef75d60a
+3834fd814e74e8af277939dabd69ecc780affd21

src/llama-model-loader.cpp

@@ -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));
 

src/llama-model-loader.h

@@ -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;
 

src/llama-quant.cpp

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

src/unicode.cpp

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

tests/CMakeLists.txt

@@ -152,7 +152,6 @@ llama_build(test-recurrent-state-rollback.cpp)
 
 if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
     # these tests are disabled on Windows because they use internal functions not exported with LLAMA_API (when building with shared libraries)
-    llama_build_and_test(test-unicode.cpp)
     llama_build_and_test(test-sampling.cpp)
     llama_build_and_test(test-reasoning-budget.cpp)
     llama_build_and_test(test-grammar-parser.cpp)
@@ -311,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)

tests/test-backend-ops.cpp

@@ -3061,36 +3061,28 @@ struct test_cpy : public test_case {
 };
 
 // GGML_OP_CONT
-// permute = {0, 0, 0, 0} means no permutation: the source is transposed (or
-// view-sliced). A non-identity permute applies ggml_permute before ggml_cont.
 struct test_cont : public test_case {
     const ggml_type type;
     const std::array<int64_t, 4> ne;
     bool use_view_slice;
-    const std::array<int64_t, 4> permute;
 
     std::string vars() override {
-        return VARS_TO_STR4(type, ne, use_view_slice, permute);
+        return VARS_TO_STR3(type, ne, use_view_slice);
     }
 
     test_cont(ggml_type type = GGML_TYPE_F32,
             std::array<int64_t, 4> ne = {10, 10, 10, 1},
-            bool use_view_slice = false,
-            std::array<int64_t, 4> permute = {0, 0, 0, 0})
-        : type(type), ne(ne), use_view_slice(use_view_slice), permute(permute) {}
+            bool use_view_slice = false)
+        : type(type), ne(ne), use_view_slice(use_view_slice) {}
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         ggml_tensor * src = ggml_new_tensor(ctx, type, 4, ne.data());
         ggml_set_param(src);
         ggml_set_name(src, "src");
 
-        const bool permuted = permute[0] != 0 || permute[1] != 0 || permute[2] != 0 || permute[3] != 0;
 
         ggml_tensor * dst;
-        if (permuted) {
-            dst = ggml_permute(ctx, src, permute[0], permute[1], permute[2], permute[3]);
-            ggml_set_name(dst, "src_permuted");
-        } else if (use_view_slice) {
+        if (use_view_slice) {
             dst = ggml_view_4d(ctx, src, src->ne[0], 1, src->ne[2], src->ne[3],
                 src->nb[1], src->nb[2], src->nb[3], src->nb[0] * (src->ne[1] - 1));
             ggml_set_name(dst, "src_view_slice");
@@ -4478,10 +4470,9 @@ struct test_mul_mat : public test_case {
     const std::array<int64_t, 4> per; // permutation of dimensions
     const int64_t k_v; // size of k in memory, resulting in a non-contiguous view for k_v > k, no view for k_v == 0
     const uint32_t o; // number of outputs
-    const bool src_overlap; // a and b are overlapping views of the same tensor
 
     std::string vars() override {
-        return VARS_TO_STR11(type_a, type_b, m, n, k, bs, nr, per, k_v, o, src_overlap);
+        return VARS_TO_STR10(type_a, type_b, m, n, k, bs, nr, per, k_v, o);
     }
 
     double max_nmse_err() override {
@@ -4510,8 +4501,8 @@ struct test_mul_mat : public test_case {
             std::array<int64_t, 2> bs = {10, 10},
             std::array<int64_t, 2> nr = {2, 2},
             std::array<int64_t, 4> per = {0, 1, 2, 3},
-            int64_t k_v = 0, uint32_t o = 1, bool src_overlap = false)
-        : type_a(type_a), type_b(type_b), m(m), n(n), k(k), bs(bs), nr(nr), per(per), k_v(k_v), o(o), src_overlap(src_overlap) {}
+            int64_t k_v = 0, uint32_t o = 1)
+        : type_a(type_a), type_b(type_b), m(m), n(n), k(k), bs(bs), nr(nr), per(per), k_v(k_v), o(o) {}
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         // C^T = A * B^T: (k, m) * (k, n) => (m, n)
@@ -4544,18 +4535,6 @@ struct test_mul_mat : public test_case {
             b = ggml_permute(ctx, b, per[0], per[1], per[2], per[3]);
             ggml_set_name(a, "a_permuted");
             ggml_set_name(b, "b_permuted");
-        } else if (src_overlap) {
-            GGML_ASSERT(type_a == type_b);
-            GGML_ASSERT(k_v == 0);
-
-            // a and b are interleaved views of the same tensor: (e.g. fused QKV in MiniMax-01)
-            ggml_tensor * base = ggml_new_tensor_4d(ctx, type_a, 2*k, std::max(m, n), bs[0]*nr[0], bs[1]*nr[1]);
-            ggml_set_name(base, "base");
-
-            a = ggml_view_4d(ctx, base, k, m, bs[0],       bs[1],       base->nb[1], base->nb[2], base->nb[3], 0);
-            b = ggml_view_4d(ctx, base, k, n, bs[0]*nr[0], bs[1]*nr[1], base->nb[1], base->nb[2], base->nb[3], k*ggml_type_size(type_a));
-            ggml_set_name(a, "a");
-            ggml_set_name(b, "b");
         } else {
             const int64_t k_physical = k_v == 0 ? k : k_v;
             a = ggml_new_tensor_4d(ctx, type_a, k_physical, m, bs[0],       bs[1]);
@@ -5352,27 +5331,24 @@ struct test_rope : public test_case {
     int v; // view (1 : non-contiguous a)
     bool forward;
     bool inplace;
-    int n_offs; // offset of the rotated dims window, set via ggml_rope_set_offset()
 
     std::string vars() override {
         // forward can be inferred from the op, does not need to be printed
-        return VARS_TO_STR12(type, ne_a, n_dims, mode, n_ctx, fs, ef, af, ff, v, inplace, n_offs);
+        return VARS_TO_STR11(type, ne_a, n_dims, mode, n_ctx, fs, ef, af, ff, v, inplace);
     }
 
     test_rope(ggml_type type = GGML_TYPE_F32,
             std::array<int64_t, 4> ne_a = {10, 5, 3, 1},
             int n_dims = 10, int mode = GGML_ROPE_TYPE_NORMAL, int n_ctx = 512, float fs = 1.0f,
-            float ef = 0.0f, float af = 0.0f, bool ff = false, int v = 0, bool forward = true, bool inplace = false,
-            int n_offs = 0)
-        : type(type), ne_a(ne_a), n_dims(n_dims), mode(mode), n_ctx(n_ctx), fs(fs), ef(ef), af(af), ff(ff), v(v), forward(forward), inplace(inplace), n_offs(n_offs) {}
+            float ef = 0.0f, float af = 0.0f, bool ff = false, int v = 0, bool forward = true, bool inplace = false)
+        : type(type), ne_a(ne_a), n_dims(n_dims), mode(mode), n_ctx(n_ctx), fs(fs), ef(ef), af(af), ff(ff), v(v), forward(forward), inplace(inplace) {}
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         ggml_tensor * a;
         if (v & 1) {
             auto ne = ne_a; ne[0] *= 2; ne[1] *= 4; ne[2] *= 3;
             a = ggml_new_tensor(ctx, type, 4, ne.data());
-            if (forward && n_offs == 0) {
-                // FIXME: support gradients with n_offs > 0
+            if (forward) {
                 ggml_set_param(a);
             }
             ggml_set_name(a, "a");
@@ -5385,8 +5361,7 @@ struct test_rope : public test_case {
             // non-aligned buffer offset, which exercises backends' alignment paths.
             auto ne = ne_a; ne[0] *= 2;
             a = ggml_new_tensor(ctx, type, 4, ne.data());
-            if (forward && n_offs == 0) {
-                // FIXME: support gradients with n_offs > 0
+            if (forward) {
                 ggml_set_param(a);
             }
             ggml_set_name(a, "a");
@@ -5397,8 +5372,7 @@ struct test_rope : public test_case {
             ggml_set_name(a, "view_of_a");
         } else {
             a = ggml_new_tensor(ctx, type, 4, ne_a.data());
-            if (forward && n_offs == 0) {
-                // FIXME: support gradients with n_offs > 0
+            if (forward) {
                 ggml_set_param(a);
             }
             ggml_set_name(a, "a");
@@ -5459,9 +5433,6 @@ struct test_rope : public test_case {
                 out = ggml_rope_ext_back(ctx, a, pos, freq, n_dims, mode, 0, 10000.0f, fs, ef, af, 1.0f, 1.0f);
             }
         }
-        if (n_offs != 0) {
-            out = ggml_rope_set_offset(out, n_offs);
-        }
         ggml_set_name(out, "out");
 
         return out;
@@ -8921,20 +8892,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         }
     }
 
-    for (ggml_type type_dst : { GGML_TYPE_F32, GGML_TYPE_F16 }) {
-        for (std::array<int64_t, 4> ne : std::initializer_list<std::array<int64_t, 4>>{
-                {10, 10, 10, 1}, {33, 5, 7, 1}, {64, 3, 65, 1}, {2, 3, 5, 7},
-                // large, tile-aligned and tile-unaligned, matching the perf cases
-                {1024, 64, 64, 1}, {2304, 64, 64, 1}, {1000, 33, 65, 1} }) {
-            for (std::array<int64_t, 4> perm : std::initializer_list<std::array<int64_t, 4>>{
-                    {2, 1, 0, 3},   // 0<->2 swap
-                    {1, 2, 0, 3},   // 3-cycle
-                    {0, 2, 1, 3} }) {
-                test_cases.emplace_back(new test_cont(type_dst, ne, false, perm));
-            }
-        }
-    }
-
     auto add_test_bin_bcast = [&](ggml_type type, std::array<int64_t, 4> ne, std::array<int, 4> nr, bool perm1 = false, bool src_overlap = false) {
         for (auto op : {ggml_add, ggml_sub, ggml_mul, ggml_div}) {
             test_cases.emplace_back(new test_bin_bcast(op, type, ne, nr, 1, perm1, src_overlap));
@@ -9286,7 +9243,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 1056, 1, 67,  {1,  1}, {4, 1}, {0, 2, 1, 3}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 16, 32, 32, { 1,  1}, {1, 1}, {0, 1, 2, 3}, 64, 3));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 64, 77, 77, {12,1}, {1,1}));
-    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 32, 4, 96, {3, 2}, {1, 1}, {0, 1, 2, 3}, 0, 1, true));
 
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 576, 512, 576, {1,1}, {1,1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 1, 2048, 8192, {1,  1}, {1, 1}));
@@ -9629,20 +9585,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         }
     }
 
-    // rotated dims window at an offset (ggml_rope_set_offset), not supported for vision mode
-    for (ggml_type type : {GGML_TYPE_F32, GGML_TYPE_F16}) {
-        for (bool fw : {true, false}) { // fw == forward
-            for (bool ff : {false, true}) {
-                test_cases.emplace_back(new test_rope(type, {128, 32, 2, 1}, 32, GGML_ROPE_TYPE_NORMAL, 512, 1.4245f, 0.7465f, 1.4245f, ff, 0, fw, false, 32));
-                test_cases.emplace_back(new test_rope(type, {128, 32, 2, 1}, 32, GGML_ROPE_TYPE_NEOX,   512, 1.4245f, 0.7465f, 1.4245f, ff, 0, fw, false, 32));
-                test_cases.emplace_back(new test_rope(type, {128, 12, 2, 1}, 24, GGML_ROPE_TYPE_MROPE,  512, 1.4245f, 0.7465f, 1.4245f, ff, 0, fw, false, 32));
-                test_cases.emplace_back(new test_rope(type, {128, 12, 2, 1}, 24, GGML_ROPE_TYPE_IMROPE, 512, 1.4245f, 0.7465f, 1.4245f, ff, 0, fw, false, 32));
-            }
-        }
-        // inplace with an offset
-        test_cases.emplace_back(new test_rope(type, {128, 32, 2, 1}, 32, GGML_ROPE_TYPE_NEOX, 512, 1.4245f, 0.7465f, 1.4245f, false, 0, true, true, 32));
-    }
-
     for (int v : { 0, 1, 2, 3 }) {
         for (int dim : { 0, 1, 2, 3, }) {
             test_cases.emplace_back(new test_concat(GGML_TYPE_F32, {11, 12, 13, 14}, 7, dim, v));
@@ -10100,17 +10042,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
-    // CONT of a 0<->2 permute at DeepSeek-V4 lightning-indexer shapes:
-    // indexer_kq is [n_kv, n_tokens, n_head=64] and gets ggml_cont(ggml_permute(.., 2,1,0,3)).
-    for (int64_t n_kv : { 1024, 1280, 2048, 2304 }) {
-        test_cases.emplace_back(new test_cont(
-            GGML_TYPE_F32, {n_kv, 64, 64, 1}, false, {2, 1, 0, 3}));
-    }
-    for (int64_t n_kv : { 2048, 2304 }) {
-        test_cases.emplace_back(new test_cont(
-            GGML_TYPE_F32, {n_kv, 512, 64, 1}, false, {2, 1, 0, 3}));
-    }
-
     // Conv2d: K=CRS=NPQ=4096 matmul performance
     uint32_t                        iwh_idx  = 0;
     uint32_t                        kwh_idx  = 1;

tests/test-llama-archs.cpp

@@ -456,7 +456,7 @@ static bool arch_supported(const llm_arch arch) {
 
     // FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
 #ifdef GGML_USE_WEBGPU
-    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA) {
+    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_MINIMAX_01) {
         return false;
     }
 #endif // GGML_USE_WEBGPU

tests/test-mtmd-impl.cpp

@@ -1,158 +0,0 @@
-#include "testing.h"
-
-#include "mtmd-image.h"
-#include "mtmd-internal.h"
-
-#include <iostream>
-#include <stdexcept>
-#include <string>
-#include <tuple>
-#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"));
-    }
-}
-
-//
-// mtmd temporal merge
-//
-
-MAKE_TEST(test_temporal_merge_grouping) {
-    std::vector<mtmd::bitmap_ptr> pool; // keeps the bitmaps alive until the end of the test
-
-    // spec chars:
-    //   v = video frame, w = video frame of another size, a = audio, i = plain image, t = text
-    auto make_parts = [&pool](const std::string & spec) {
-        std::vector<mtmd_input_part> parts;
-        for (char c : spec) {
-            if (c == 't') {
-                parts.push_back({ "hello", nullptr });
-                continue;
-            }
-            mtmd_bitmap * bm = nullptr;
-            switch (c) {
-                case 'v': bm = mtmd_bitmap_init(100, 100, nullptr);   break;
-                case 'w': bm = mtmd_bitmap_init(200, 200, nullptr);   break;
-                case 'a': bm = mtmd_bitmap_init_from_audio(100, nullptr); break;
-                case 'i': bm = mtmd_bitmap_init(100, 100, nullptr);   break;
-                default: throw std::runtime_error(std::string("unknown spec char: ") + c);
-            }
-            mtmd_bitmap_set_mergeable(bm, c != 'i');
-            pool.emplace_back(bm);
-            parts.push_back({ "", bm });
-        }
-        return parts;
-    };
-
-    // { parts, n_merge, expected size of each group }
-    const std::vector<std::tuple<std::string, int, std::string>> cases = {
-        { "vv",   2, "2"    },
-        { "vvv",  2, "21"   },
-        { "vvvv", 2, "22"   },
-        { "vvi",  2, "21"   },
-        { "tvvt", 2, "2"    },
-        { "vtv",  2, "11"   }, // text in between breaks the merge
-        { "vw",   2, "11"   }, // different sizes cannot be merged
-        { "aa",   2, "11"   }, // audio is never merged
-        { "ii",   2, "11"   }, // two unrelated images must stay separated
-        { "iv",   2, "11"   },
-        { "vi",   2, "11"   },
-        { "vv",   1, "11"   }, // model without temporal merge
-    };
-
-    for (const auto & [spec, n_merge, expected] : cases) {
-        auto parts  = make_parts(spec);
-        auto groups = mtmd_group_mergeable_bitmaps(parts, n_merge);
-
-        std::string actual;
-        for (const auto & group : groups) {
-            actual += std::to_string(group.size());
-        }
-
-        const std::string name = "\"" + spec + "\" with n_merge=" + std::to_string(n_merge);
-        t.assert_equal("groups for " + name, expected, actual);
-
-        size_t n_bitmap_parts = 0;
-        for (const auto & p : parts) {
-            n_bitmap_parts += p.bitmap != nullptr ? 1 : 0;
-        }
-        t.assert_equal("remaining bitmap parts for " + name, groups.size(), n_bitmap_parts);
-    }
-}
-
-//
-// 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();
-}

tests/test-unicode.cpp

@@ -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;
-}

tools/mtmd/CMakeLists.txt

@@ -17,7 +17,6 @@ add_library(mtmd
             mtmd-audio.cpp
             mtmd-image.cpp
             mtmd.h
-            mtmd-internal.h
             mtmd-helper.cpp
             mtmd-helper-gen.cpp
             mtmd-helper-common.h
@@ -79,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)
@@ -91,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

tools/mtmd/README-dev.md

@@ -21,7 +21,7 @@ A typical pipeline of the core libmtmd is as follows:
 - A bitmap (RGB image or PCM audio) is created
 - Bitmap and the text prompt is provided to `mtmd_tokenize()` that breaks the input into chunks
     - The tokenizer function first expands a "lazy" bitmap if it finds one. Typically, this is used by video, so that one media token corresponds to one input bitmap
-    - For models that support "fused" temporal frames like Qwen-VL, the tokenizer tries to merge pair of consecutive frames into one batch. Only bitmaps marked by `mtmd_bitmap_set_mergeable()` are merged
+    - For models that support "fused" temporal frames like Qwen-VL, the tokenizer tries to merge pair of consecutive frames into one batch
     - The preprocessor will then be called, which produces a list of chunks
     - Depending on the model itself, special tokens will be injected to separate image chunks (i.e. llava-uhd-style models)
 - Multiple bitmaps may be batched together to form a larger `mtmd_batch()`

tools/mtmd/clip-impl.h

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

tools/mtmd/models/deepseekocr.cpp

@@ -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);

tools/mtmd/mtmd-helper.cpp

@@ -12,7 +12,7 @@
 #include "mtmd-helper-common.h"
 #include "llama.h"
 
-#include "hash/hash.h"
+#include "hash.h"
 
 #include <algorithm>
 #include <cinttypes>
@@ -727,9 +727,7 @@ struct mtmd_helper_video {
 
         LOG_DBG("%s: frame %d read OK\n", __func__, current_frame);
         current_frame++;
-        mtmd_bitmap * frame = mtmd_bitmap_init(info.width, info.height, frame_buf.data());
-        mtmd_bitmap_set_mergeable(frame, true);
-        return frame;
+        return mtmd_bitmap_init(info.width, info.height, frame_buf.data());
     }
 
     int32_t read_next(mtmd_bitmap ** out_bitmap, char ** out_text) {

tools/mtmd/mtmd-image.cpp

@@ -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};

tools/mtmd/mtmd-image.h

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

tools/mtmd/mtmd-internal.h

@@ -1,19 +0,0 @@
-#pragma once
-
-#include "mtmd.h"
-
-#include <string>
-#include <vector>
-
-// !!! Internal header, to be used by mtmd and its unit tests only !!!
-
-#define MTMD_INTERNAL_HEADER
-
-// bitmap is null for text parts
-struct mtmd_input_part {
-    std::string text;
-    const mtmd_bitmap * bitmap;
-};
-
-// [QWEN_VIDEO] merged parts are erased from `parts`, so one group always maps to one part
-std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_input_part> & parts, int n_merge);

tools/mtmd/mtmd.cpp

@@ -1,7 +1,6 @@
 #include "clip.h"
 #include "clip-impl.h"
 #include "mtmd.h"
-#include "mtmd-internal.h"
 #include "mtmd-audio.h"
 #include "mtmd-image.h"
 #include "debug/mtmd-debug.h"
@@ -150,7 +149,6 @@ struct mtmd_bitmap {
     uint32_t ny = 0;
     std::string id; // optional user-defined id, for ex: can be set to image hash, useful for KV cache tracking
     bool is_audio = false; // true if the bitmap is audio
-    bool mergeable = false; // [QWEN_VIDEO] set only on frames of the same video
 
     // lazy-loaded bitmap
     mtmd_bitmap_lazy_callback lazy_callback = nullptr;
@@ -188,9 +186,7 @@ struct mtmd_bitmap {
 
     bool can_merge_with(const mtmd_bitmap & other) const {
         // [QWEN_VIDEO] can (temporal) merge if both are images with same size
-        return mergeable && other.mergeable
-            && !is_audio && !other.is_audio
-            && nx == other.nx && ny == other.ny;
+        return !is_audio && !other.is_audio && nx == other.nx && ny == other.ny;
     }
 
   private:
@@ -1080,25 +1076,6 @@ void mtmd_free(mtmd_context * ctx) {
     delete ctx;
 }
 
-std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_input_part> & parts, int n_merge) {
-    std::vector<std::vector<const mtmd_bitmap *>> output;
-    for (size_t i = 0; i < parts.size(); i++) {
-        if (parts[i].bitmap == nullptr) {
-            continue; // text part
-        }
-        const bool has_next = n_merge > 1 && i + 1 < parts.size() && parts[i + 1].bitmap != nullptr;
-        if (has_next && parts[i].bitmap->can_merge_with(*parts[i + 1].bitmap)) {
-            LOG_DBG("%s: merging 2 frames at part index %zu and %zu\n", __func__, i, i + 1);
-            output.push_back({parts[i].bitmap, parts[i + 1].bitmap});
-            parts.erase(parts.begin() + i + 1);
-            continue;
-        }
-        LOG_DBG("%s: no merging for part index %zu\n", __func__, i);
-        output.push_back({parts[i].bitmap});
-    }
-    return output;
-}
-
 struct mtmd_tokenizer {
     mtmd_context * ctx;
 
@@ -1107,7 +1084,10 @@ struct mtmd_tokenizer {
     bool parse_special;
     const llama_vocab * vocab;
 
-    using part = mtmd_input_part;
+    struct part {
+        std::string text;
+        const mtmd_bitmap * bitmap;
+    };
     std::vector<part> parts;
     // these will be freed when mtmd_tokenizer finishes
     std::vector<mtmd::bitmap> bm_from_lazy; // TODO @ngxson : refactor, free bm_from_lazy progressively
@@ -1212,7 +1192,34 @@ struct mtmd_tokenizer {
             GGML_ASSERT(n_merge_frames <= 2 && "we only support merging maximum 2 images for now; open an issue if this model supports merging more");
         }
 
-        auto merged_bitmaps = mtmd_group_mergeable_bitmaps(parts, n_merge_frames);
+        // Build merged_bitmaps: each entry is a group of 1 or 2 bitmaps.
+        // For consecutive mergeable bitmap parts, merge them and collapse the second part out of this->parts.
+        std::vector<std::vector<const mtmd_bitmap *>> merged_bitmaps;
+        if (n_merge_frames > 1) {
+            for (size_t i = 0; i < parts.size(); ++i) {
+                if (parts[i].bitmap == nullptr) {
+                    continue;
+                }
+                if (i + 1 < parts.size() && parts[i + 1].bitmap != nullptr) {
+                    const mtmd_bitmap * bm_a = parts[i].bitmap;
+                    const mtmd_bitmap * bm_b = parts[i + 1].bitmap;
+                    if (bm_a->can_merge_with(*bm_b)) {
+                        LOG_DBG("%s: merging 2 frames at part index %zu and %zu\n", __func__, i, i + 1);
+                        merged_bitmaps.push_back({bm_a, bm_b});
+                        parts.erase(parts.begin() + i + 1); // collapse the second bitmap part
+                        continue;
+                    }
+                }
+                LOG_DBG("%s: no merging for part index %zu\n", __func__, i);
+                merged_bitmaps.push_back({parts[i].bitmap});
+            }
+        } else {
+            for (const auto & p : parts) {
+                if (p.bitmap != nullptr) {
+                    merged_bitmaps.push_back({p.bitmap});
+                }
+            }
+        }
 
         size_t i_bm = 0;
         for (const auto & p : parts) {
@@ -2193,10 +2200,6 @@ void mtmd_bitmap_set_id(mtmd_bitmap * bitmap, const char * id) {
     }
 }
 
-void mtmd_bitmap_set_mergeable(mtmd_bitmap * bitmap, bool mergeable) {
-    bitmap->mergeable = mergeable;
-}
-
 mtmd_bitmap * mtmd_bitmap_init_lazy(mtmd_context * ctx,
                                     const char * id,
                                     void * user_data,

tools/mtmd/mtmd.h

@@ -154,8 +154,7 @@ MTMD_API const char * mtmd_get_marker(const mtmd_context * ctx);
 //     length of data must be nx * ny * 3
 //     the data is in RGBRGBRGB... format
 //     note: some video-capable models (i.e. qwen-vl) can merge consecutive bitmaps
-//           into one chunk; mtmd_tokenize() handles this, but remember to set
-//           mtmd_bitmap_set_mergeable(true) for every frame
+//           into one chunk, mtmd_tokenize() will automatically handle this
 // if bitmap is audio:
 //     length of data must be n_samples * sizeof(float)
 //     the data is in float format (PCM F32)
@@ -176,8 +175,6 @@ MTMD_API void                  mtmd_bitmap_free       (mtmd_bitmap * bitmap);
 // these getters/setters are dedicated functions, so you can for example calculate the hash of the image based on mtmd_bitmap_get_data()
 MTMD_API const char * mtmd_bitmap_get_id(const mtmd_bitmap * bitmap);
 MTMD_API void         mtmd_bitmap_set_id(mtmd_bitmap * bitmap, const char * id);
-// if true, this bitmap can be merged (temporal merge) with an adjacent mergeable bitmap by certain video input models
-MTMD_API void         mtmd_bitmap_set_mergeable(mtmd_bitmap * bitmap, bool mergeable);
 
 // mtmd_bitmap lazy
 //

tools/perplexity/perplexity.cpp

@@ -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;

tools/server/CMakeLists.txt

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

tools/server/README-dev.md

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

tools/server/README.md

@@ -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**
 
@@ -1759,7 +1778,6 @@ The precedence rule for preset options is as follows:
 We also offer additional options that are exclusive to presets (these aren't treated as command-line arguments):
 - `load-on-startup` (boolean): Controls whether the model loads automatically when the server starts
 - `stop-timeout` (int, seconds): After requested unload, wait for this many seconds before forcing termination (default: 10)
-- `dedup-cache-models` (boolean): When the preset uses `hf-repo` pointing to a model that is already downloaded, hide the corresponding cached model entry from `GET /models` (the preset entry remains visible). Set it in the `[*]` section to apply to all presets.
 
 ### Routing requests
 

tools/server/server-cache-disk.cpp

@@ -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);
+}

tools/server/server-cache-disk.h

@@ -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);

tools/server/server-common.cpp

@@ -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");

tools/server/server-common.h

@@ -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;
 

tools/server/server-context.cpp

@@ -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() {

tools/server/server-models.cpp

@@ -555,40 +555,6 @@ void server_models::load_models() {
         return source_map.count(name) ? source_map.at(name) : SERVER_MODEL_SOURCE_PRESET;
     };
 
-    // hide cache models whose resolved file is already used by a preset with dedup-cache-models enabled
-    std::set<std::string> hidden_models;
-    {
-        std::set<std::string> preset_paths;
-        for (const auto & [name, preset] : custom_presets) {
-            std::string val;
-            if (!preset.get_option(COMMON_ARG_PRESET_DEDUP_CACHE_MODELS, val) || !common_arg_utils::is_truthy(val)) {
-                continue;
-            }
-            std::string hf_repo;
-            if (!preset.get_option("LLAMA_ARG_HF_REPO", hf_repo) || hf_repo.empty()) {
-                continue;
-            }
-            std::string hf_file;
-            preset.get_option("LLAMA_ARG_HF_FILE", hf_file);
-            std::string path = common_download_resolve_path(hf_repo, hf_file);
-            if (!path.empty()) {
-                preset_paths.insert(path);
-            }
-        }
-        if (!preset_paths.empty()) {
-            for (const auto & [name, preset] : cached_models) {
-                if (get_source(name) != SERVER_MODEL_SOURCE_CACHE) {
-                    continue; // merged with another source, not a pure cache entry
-                }
-                std::string path = common_download_resolve_path(name);
-                if (!path.empty() && preset_paths.count(path)) {
-                    SRV_INF("hiding cache model name=%s (deduplicated by a preset)\n", name.c_str());
-                    hidden_models.insert(name);
-                }
-            }
-        }
-    }
-
     // Helpers that read `mapping` - must be called while holding the lock.
     std::unordered_set<std::string> custom_names;
     for (const auto & [name, preset] : custom_presets) custom_names.insert(name);
@@ -624,11 +590,6 @@ void server_models::load_models() {
             }
         }
     };
-    auto apply_hidden = [&]() {
-        for (auto & [name, inst] : mapping) {
-            inst.meta.hidden = hidden_models.count(name) > 0;
-        }
-    };
     // update_args() injects HOST/PORT/ALIAS, so strip them before comparing presets
     auto preset_options_for_compare = [](common_preset p) {
         p.unset_option("LLAMA_ARG_HOST");
@@ -669,7 +630,6 @@ void server_models::load_models() {
             add_model(std::move(meta));
         }
         apply_stop_timeout();
-        apply_hidden();
         log_available_models();
 
         std::vector<std::string> models_to_load;
@@ -846,7 +806,6 @@ void server_models::load_models() {
         }
 
         apply_stop_timeout();
-        apply_hidden();
 
         // clear reload flag before unlocking for autoload - load() blocks on !is_reloading,
         // so clearing it here (while still locked) prevents a deadlock in the autoload calls below
@@ -1066,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 {
@@ -1970,9 +1926,6 @@ void server_models_routes::init_routes() {
         auto all_models = models.get_all_meta();
         std::time_t t = std::time(0);
         for (const auto & meta : all_models) {
-            if (meta.hidden) {
-                continue; // cache model deduplicated by a preset
-            }
             json status {
                 {"value",  server_model_status_to_string(meta.status)},
                 {"args",   meta.args},

tools/server/server-models.h

@@ -84,7 +84,6 @@ struct server_model_meta {
     int exit_code = 0; // exit code of the model instance process (only valid if status == FAILED)
     int stop_timeout = 0; // seconds to wait before force-killing the model instance during shutdown
     mtmd_caps multimodal; // multimodal capabilities
-    bool hidden = false; // hidden from GET /models, but still accept if requested
 
     bool is_ready() const {
         return status == SERVER_MODEL_STATUS_LOADED;

tools/server/server-task.cpp

@@ -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());
+}

tools/server/server-task.h

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

tools/server/server-tools.cpp

@@ -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) {

tools/server/server-tools.h

@@ -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;