ui: Move styles/ to $lib scope (#26950 )

* refactor: Move `styles/` to `src/lib` and remove legacy alias * chore: Add newline
ui: Stores architecture improvements (#26910 )
2026-08-13 08:16:55 +02:00 · 2026-08-13 08:13:52 +02:00 · 2026-08-13 08:11:30 +02:00 · 2026-08-13 08:07:34 +02:00 · 2026-08-13 07:57:24 +03:00 · 2026-08-13 06:53:56 +02:00
482 changed files with 9814 additions and 4020 deletions
@@ -8,8 +8,26 @@ inputs:
 runs:
  using: "composite"
  steps:
-    - name: Setup ROCm
-      uses: ./.github/actions/install-exe
-      with:
-        url: https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ inputs.version }}-Win11-For-HIP.exe
-        args: -install
+    - name: Install ROCm with Wheels
+      shell: pwsh
+      run: |
+        $ErrorActionPreference = "Stop"
+        write-host "Setting up Python virtual environment"
+
+        # Create the venv directly at the cache location to avoid relocation issues
+        New-Item -Path "C:\TheRock\build" -ItemType Directory -Force | Out-Null
+        python -m venv C:\TheRock\build\.venv
+        & C:\TheRock\build\.venv\Scripts\Activate.ps1
+
+        write-host "Upgrading pip"
+        python -m pip install --upgrade pip
+
+        write-host "Installing ROCm wheels for multi-arch support"
+        # Install ROCm wheels for multi-arch support (this may take several minutes)
+        python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ inputs.version }}"
+
+        # Pre-expand the devel tree so it is included in the cache
+        write-host "Initializing ROCm devel tree"
+        rocm-sdk init
+        if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
+        write-host "Completed ROCm wheel installation to C:\TheRock\build"
@@ -119,27 +119,27 @@ jobs:
          version_major: ${{ env.OPENVINO_VERSION_MAJOR }}
          version_full: ${{ env.OPENVINO_VERSION_FULL }}

-  windows-2022-rocm-cache:
-    runs-on: windows-2022
+  # windows-2022-rocm-cache:
+  #   runs-on: windows-2022

-    env:
-      # Make sure this is in sync with build.yml
-      HIPSDK_INSTALLER_VERSION: "26.Q1"
+  #   env:
+  #     # Make sure this is in sync with release.yml and build-cuda-windows.yml
+  #     ROCM_VERSION: "7.14.0"

-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
+  #   steps:
+  #     - name: Clone
+  #       id: checkout
+  #       uses: actions/checkout@v6

-      - name: Setup Cache
-        uses: actions/cache@v5
-        id: cache-rocm
-        with:
-          path: C:\Program Files\AMD\ROCm
-          key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
+  #     - name: Setup Cache
+  #       uses: actions/cache@v5
+  #       id: cache-rocm
+  #       with:
+  #         path: C:\TheRock\build
+  #         key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}

-      - name: Setup ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
-        uses: ./.github/actions/windows-setup-rocm
-        with:
-          version: ${{ env.HIPSDK_INSTALLER_VERSION }}
+  #     - name: Setup ROCm
+  #       if: steps.cache-rocm.outputs.cache-hit != 'true'
+  #       uses: ./.github/actions/windows-setup-rocm
+  #       with:
+  #         version: ${{ env.ROCM_VERSION }}
@@ -5,7 +5,7 @@ on:

 jobs:
  linux:
-    runs-on: [self-hosted, Linux, CPU]
+    runs-on: [self-hosted, Linux]
    steps:
      - uses: actions/checkout@v6
        with:
@@ -21,15 +21,21 @@ jobs:
                -DLLAMA_BUILD_TOOLS=OFF \
                -DLLAMA_BUILD_EXAMPLES=OFF \
                -DLLAMA_BUILD_APP=OFF \
+                -DLLAMA_BUILD_IS_DEV=OFF \
                -DCMAKE_BUILD_TYPE=Release
-          cmake --build build --config Release
+          cmake --build build --config Release -j $(nproc)
          cmake --install build --prefix "$PREFIX" --config Release

          export LLAMA_CONFIG="$PREFIX"/lib/cmake/llama/llama-config.cmake
          tclsh <<'EOF'
          set build(commit)  [string trim [exec git rev-parse --short HEAD]]
          set build(number)  [string trim [exec git rev-list  --count HEAD]]
-          set build(version) "0.0.$build(number)"
+
+          set cmakelists [read [open "CMakeLists.txt" r]]
+          regexp {set\(LLAMA_VERSION_MAJOR\s+(\d+)\)} $cmakelists -> major
+          regexp {set\(LLAMA_VERSION_MINOR\s+(\d+)\)} $cmakelists -> minor
+          regexp {set\(LLAMA_VERSION_PATCH\s+(\d+)\)} $cmakelists -> patch
+          set build(version) "$major.$minor.$patch"

          set llamaconfig [read [open "$env(LLAMA_CONFIG)" r]]
          set checks [list "set\\(LLAMA_VERSION     \\s+$build(version)\\)" \
@@ -48,4 +54,4 @@ jobs:

          cd examples/simple-cmake-pkg
          cmake -S . -B build -DCMAKE_PREFIX_PATH="$PREFIX"/lib/cmake
-          cmake --build build
+          cmake --build build -j $(nproc)
@@ -94,8 +94,10 @@ jobs:
        id: cmake_build
        run: |
          cmake -B build \
+            -DGGML_NATIVE=OFF \
            -DLLAMA_FATAL_WARNINGS=ON \
-            -DGGML_RPC=ON
+            -DGGML_RPC=ON \
+            -DGGML_NATIVE=OFF
          time cmake --build build --config Release -j $(nproc)

      - name: Test
@@ -83,7 +83,7 @@ jobs:

    env:
      # Make sure this is in sync with build-cache.yml
-      HIPSDK_INSTALLER_VERSION: "26.Q1"
+      ROCM_VERSION: "7.14.0"

    strategy:
      matrix:
@@ -97,36 +97,53 @@ jobs:
        id: checkout
        uses: actions/checkout@v6

-      - name: Grab rocWMMA package
-        id: grab_rocwmma
-        run: |
-          curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2.1/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70201-81~24.04_amd64.deb"
-          7z x rocwmma.deb
-          7z x data.tar
-
-      - name: Use ROCm Installation Cache
-        uses: actions/cache@v5
-        id: cache-rocm
-        with:
-          path: C:\Program Files\AMD\ROCm
-          key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
+      # - name: Cache ROCm Installation
+      #   uses: actions/cache@v5
+      #   id: cache-rocm
+      #   with:
+      #     path: C:\TheRock\build
+      #     key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}

      - name: Setup ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
+        # if: steps.cache-rocm.outputs.cache-hit != 'true'
        uses: ./.github/actions/windows-setup-rocm
        with:
-          version: ${{ env.HIPSDK_INSTALLER_VERSION }}
+          version: ${{ env.ROCM_VERSION }}
+
+      - name: Setup ROCm Environment
+        run: |
+          $ErrorActionPreference = "Stop"
+
+          # Activate venv from cache or fresh install
+          & C:\TheRock\build\.venv\Scripts\Activate.ps1
+
+          # Expand the devel tree (idempotent; no-op if already done during install)
+          rocm-sdk init
+          if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
+
+          # Get ROCm installation paths using the rocm-sdk CLI tool
+          $rocmPath = (rocm-sdk path --root)
+          if (-not $rocmPath) { throw "rocm-sdk path --root returned empty - devel package may not be installed" }
+          $rocmPath = $rocmPath.Trim()
+          $cmakePath = (rocm-sdk path --cmake).Trim()
+          $binPath = (rocm-sdk path --bin).Trim()
+          write-host "ROCm root: $rocmPath"
+
+          echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
+          echo "CMAKE_PREFIX_PATH=$cmakePath" >> $env:GITHUB_ENV
+          echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
+          echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
+          echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
+          echo "$binPath" >> $env:GITHUB_PATH
+
+          # Keep venv in PATH for subsequent steps
+          echo "C:\TheRock\build\.venv\Scripts" >> $env:GITHUB_PATH

      - name: Verify ROCm
        id: verify
        run: |
-          # Find and test ROCm installation
-          $clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1
-          if (-not $clangPath) {
-            Write-Error "ROCm installation not found"
-            exit 1
-          }
-          & $clangPath.FullName --version
+          # Test the ROCm clang shipped in the installed wheel
+          & "${env:HIP_PATH}\lib\llvm\bin\clang.exe" --version

      - name: ccache
        uses: ggml-org/ccache-action@v1.2.21
@@ -134,28 +151,27 @@ jobs:
          # TODO: this build does not match the build in release.yml, so we use a different cache key
          #       ideally, the builds should match, similar to the CUDA build above so that we would be able
          #       to populate the ccache for the release with manual runs of this workflow
-          #key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
-          key: cuda-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
+          #key: release-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
+          key: cuda-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}

      - name: Build
        id: cmake_build
        run: |
-          $env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path)
-          $env:CMAKE_PREFIX_PATH="${env:HIP_PATH}"
          cmake -G "Unix Makefiles" -B build -S . `
-            -DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" `
-            -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" `
-            -DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.1/include/" `
+            -DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
+            -DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
+            -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
+            -DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
            -DCMAKE_BUILD_TYPE=Release `
            -DLLAMA_BUILD_BORINGSSL=ON `
-            -DROCM_DIR="${env:HIP_PATH}" `
+            -DHIP_PATH="${env:HIP_PATH}" `
            -DGGML_HIP=ON `
-            -DGPU_TARGETS="gfx1100"  `
+            -DGPU_TARGETS="gfx1100" `
            -DGGML_RPC=ON
          cmake --build build -j ${env:NUMBER_OF_PROCESSORS}

      - name: ccache-clear
        uses: ./.github/actions/ccache-clear
        with:
-          #key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
-          key: cuda-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
+          #key: release-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
+          key: cuda-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
@@ -39,9 +39,9 @@ jobs:
    strategy:
      matrix:
        include:
+        # thread and address doesn't run properly on some self hosted machines, so run it on Github instead
        - sanitizer: ADDRESS
-          machine: [self-hosted, X64, Linux]
-        # thread doesn't run properly on some self hosted machines, so run it on Github instead
+          machine: ubuntu-24.04
        - sanitizer: THREAD
          machine: ubuntu-24.04
        - sanitizer: UNDEFINED
@@ -54,14 +54,14 @@ jobs:
        id: checkout
        uses: actions/checkout@v6

-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.21
-        if: ${{ matrix.sanitizer == 'THREAD' }}
-        with:
-          key: ctest-thread-ubuntu-24.04
-          variant: ccache
-          evict-old-files: 1d
-          save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
+      # - name: ccache
+      #   uses: ggml-org/ccache-action@v1.2.21
+      #   if: ${{ matrix.sanitizer != 'UNDEFINED' }}
+      #   with:
+      #     key: ctest-${{ matrix.sanitizer }}-ubuntu-24.04
+      #     variant: ccache
+      #     evict-old-files: 1d
+      #     save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}

      # with UNDEFINED sanitizer, we have to build in Debug to avoid GCC 13 false-positive warnings
      - name: Build (undefined)
@@ -0,0 +1,46 @@
+name: Make Release
+
+on:
+  workflow_dispatch:
+    inputs:
+      dry_run:
+        description: 'Dry run - validate without creating the tag'
+        required: true
+        type: boolean
+        default: true
+
+env:
+  GH_TOKEN: ${{ github.token }}
+
+permissions:
+  contents: write
+
+jobs:
+  make-release:
+    runs-on: ubuntu-latest
+
+    steps:
+      - name: Checkout
+        uses: actions/checkout@v6
+
+      - name: Run release checks
+        id: checks
+        run: bash scripts/make-release-checks.sh ${{ github.event.inputs.dry_run == 'true' && '--dry-run' || '' }}
+        env:
+          GITHUB_REPOSITORY: ${{ github.repository }}
+
+      - name: Create release tag
+        if: ${{ github.event.inputs.dry_run == 'false' }}
+        run: |
+          VERSION="${{ steps.checks.outputs.version }}"
+          git config user.name "github-actions[bot]"
+          git config user.email "github-actions[bot]@users.noreply.github.com"
+          git tag -a "${VERSION}" -m "Release ${VERSION}"
+          git push origin "${VERSION}"
+          echo "Created and pushed tag ${VERSION}"
+
+      - name: Dry run summary
+        if: ${{ github.event.inputs.dry_run == 'true' }}
+        run: |
+          echo "Dry run complete - all checks passed."
+          echo "Would have created tag: ${{ steps.checks.outputs.version }}"
@@ -748,6 +748,135 @@ jobs:
          path: llama-bin-win-cpu-${{ matrix.arch }}.zip
          name: llama-bin-win-cpu-${{ matrix.arch }}.zip

+  windows-rocm:
+    needs: [check-release]
+    if: ${{ needs.check-release.outputs.should_release == 'true' }}
+
+    runs-on: windows-2022
+
+    strategy:
+      matrix:
+        include:
+          - ROCM_VERSION: "7.14.0"
+            gpu_targets: "gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
+            build: x64
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+        with:
+          fetch-depth: 0
+
+      - name: ccache
+        uses: ggml-org/ccache-action@v1.2.21
+        with:
+          key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
+          evict-old-files: 1d
+
+      # - name: Cache ROCm Installation
+      #   id: cache-rocm
+      #   uses: actions/cache@v5
+      #   with:
+      #     path: C:\TheRock\build
+      #     key: rocm-wheels-${{ matrix.ROCM_VERSION }}-multi-arch-${{ runner.os }}
+
+      - name: Setup ROCm
+        # if: steps.cache-rocm.outputs.cache-hit != 'true'
+        uses: ./.github/actions/windows-setup-rocm
+        with:
+          version: ${{ matrix.ROCM_VERSION }}
+
+      - name: Setup ROCm Environment
+        run: |
+          $ErrorActionPreference = "Stop"
+
+          # Activate venv from cache or fresh install
+          & C:\TheRock\build\.venv\Scripts\Activate.ps1
+
+          # Expand the devel tree (idempotent; no-op if already done during install)
+          rocm-sdk init
+          if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
+
+          # Get ROCm installation paths using the rocm-sdk CLI tool
+          $rocmPath = (rocm-sdk path --root)
+          if (-not $rocmPath) { throw "rocm-sdk path --root returned empty - devel package may not be installed" }
+          $rocmPath = $rocmPath.Trim()
+          $cmakePath = (rocm-sdk path --cmake).Trim()
+          $binPath = (rocm-sdk path --bin).Trim()
+          write-host "ROCm root: $rocmPath"
+          write-host "CMake path: $cmakePath"
+          write-host "Bin path: $binPath"
+
+          echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
+          echo "CMAKE_PREFIX_PATH=$cmakePath" >> $env:GITHUB_ENV
+          echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
+          echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
+          echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
+          echo "$binPath" >> $env:GITHUB_PATH
+
+          # Keep venv in PATH for subsequent steps
+          echo "C:\TheRock\build\.venv\Scripts" >> $env:GITHUB_PATH
+
+      - name: Build
+        run: |
+          mkdir build
+          cd build
+          cmake .. `
+            -G "Unix Makefiles" `
+            -DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
+            -DCMAKE_BUILD_TYPE=Release `
+            -DGGML_BACKEND_DL=ON `
+            -DGGML_NATIVE=OFF `
+            -DGGML_CPU=ON `
+            -DGGML_CPU_ALL_VARIANTS=ON `
+            -DGGML_HIP=ON `
+            -DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
+            -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
+            -DCMAKE_C_FLAGS="-Wno-error=incompatible-pointer-types" `
+            -DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
+            -DHIP_PATH="${env:HIP_PATH}" `
+            -DGGML_HIP_ROCWMMA_FATTN=ON `
+            -DAMDGPU_TARGETS="${{ matrix.gpu_targets }}"
+          cmake --build . --config Release --parallel ${env:NUMBER_OF_PROCESSORS}
+
+      - name: ccache-clear
+        uses: ./.github/actions/ccache-clear
+        with:
+          key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
+
+      - name: Verify HIP backend was built
+        run: |
+          $hipDll = Get-ChildItem -Path build\bin -Filter "ggml-hip*.dll" -ErrorAction SilentlyContinue
+          if (-not $hipDll) {
+            Write-Host "##[error]ggml-hip*.dll was NOT produced. The HIP backend silently failed to build."
+            Write-Host "Contents of build\bin:"
+            Get-ChildItem build\bin | Format-Table -AutoSize
+            exit 1
+          }
+          Write-Host "HIP backend artifact found:"
+          $hipDll | Format-Table FullName, Length -AutoSize
+
+      - name: Determine tag name
+        id: tag
+        uses: ./.github/actions/get-tag-name
+
+      - name: Get ROCm short version
+        run: |
+          $rocmVersionShort = ('${{ matrix.ROCM_VERSION }}'.Split('.')[0..1] -join '.')
+          echo "ROCM_VERSION_SHORT=$rocmVersionShort" >> $env:GITHUB_ENV
+
+      - name: Pack artifacts
+        run: |
+          cp "LICENSE" "build\bin\"
+          7z a -snl llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip .\build\bin\*
+
+      - name: Upload artifacts
+        uses: actions/upload-artifact@v6
+        with:
+          path: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip
+          name: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip
+
  windows:
    needs: [check-release]
    if: ${{ needs.check-release.outputs.should_release == 'true' }}
@@ -1156,248 +1285,123 @@ jobs:
          path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz
          name: llama-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz

-  ubuntu-22-rocm:
-    needs: [check-release, get-version]
-    if: ${{ needs.check-release.outputs.should_release == 'true' }}
+#   ubuntu-22-rocm:
+#     needs: [check-release, get-version]
+#     if: ${{ needs.check-release.outputs.should_release == 'true' }}

-    runs-on: ubuntu-22.04
+#     runs-on: ubuntu-22.04

-    permissions:
-      actions: write
+#     permissions:
+#       actions: write

-    strategy:
-      matrix:
-        include:
-          - ROCM_VERSION: "7.2.1"
-            gpu_targets: "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1150;gfx1200;gfx1201"
-            build: 'x64'
+#     strategy:
+#       matrix:
+#         include:
+#           - ROCM_VERSION: "7.14.0"
+#             gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
+#             build: 'x64'

-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-        with:
-          fetch-depth: 0
+#     steps:
+#       - name: Clone
+#         id: checkout
+#         uses: actions/checkout@v6
+#         with:
+#           fetch-depth: 0

-      - name: Setup Node.js
-        uses: actions/setup-node@v6
-        with:
-          node-version: "24"
-          cache: "npm"
-          cache-dependency-path: "tools/ui/package-lock.json"
+#       - name: Setup Node.js
+#         uses: actions/setup-node@v6
+#         with:
+#           node-version: "24"
+#           cache: "npm"
+#           cache-dependency-path: "tools/ui/package-lock.json"

-      - name: Free up disk space
-        uses: ggml-org/free-disk-space@v1.3.1
-        with:
-          tool-cache: true
+#       - name: Free up disk space
+#         uses: ggml-org/free-disk-space@v1.3.1
+#         with:
+#           tool-cache: true

-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.21
-        with:
-          key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+#       # - name: ccache
+#       #   uses: ggml-org/ccache-action@v1.2.21
+#       #   with:
+#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}

-      - name: Dependencies
-        id: depends
-        run: |
-          sudo apt install -y build-essential git cmake wget
+#       - name: Dependencies
+#         id: depends
+#         run: |
+#           sudo apt install -y build-essential git cmake wget

-      - name: Setup Legacy ROCm
-        if: matrix.ROCM_VERSION == '7.2.1'
-        id: legacy_env
-        run: |
-          sudo mkdir --parents --mode=0755 /etc/apt/keyrings
-          wget https://repo.radeon.com/rocm/rocm.gpg.key -O - | \
-            gpg --dearmor | sudo tee /etc/apt/keyrings/rocm.gpg > /dev/null
+#       - name: Setup TheRock with Wheels
+#         id: therock_env
+#         run: |
+#           # Create Python virtual environment
+#           python3 -m venv .venv
+#           source .venv/bin/activate

-          sudo tee /etc/apt/sources.list.d/rocm.list << EOF
-          deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/${{ matrix.ROCM_VERSION }} jammy main
-          EOF
+#           # Install ROCm wheels for build
+#           # libraries = HIP runtime and CMake configs needed for linking
+#           # devel = compilers, headers, static libs
+#           python -m pip install --upgrade pip
+#           python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"

-          sudo tee /etc/apt/preferences.d/rocm-pin-600 << EOF
-          Package: *
-          Pin: release o=repo.radeon.com
-          Pin-Priority: 600
-          EOF
+#           # Get ROCm installation paths using the rocm-sdk CLI tool
+#           ROCM_PATH=$(rocm-sdk path --root)
+#           CMAKE_PATH=$(rocm-sdk path --cmake)
+#           BIN_PATH=$(rocm-sdk path --bin)
+#           echo "ROCM_PATH=$ROCM_PATH"
+#           echo "CMAKE_PATH=$CMAKE_PATH"
+#           echo "BIN_PATH=$BIN_PATH"

-          sudo apt update
-          sudo apt-get install -y libssl-dev rocm-hip-sdk
+#           # Set environment variables
+#           echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
+#           echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
+#           echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
+#           echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
+#           echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV

-      - name: Setup TheRock
-        if: matrix.ROCM_VERSION != '7.2.1'
-        id: therock_env
-        run: |
-          wget https://repo.amd.com/rocm/tarball/therock-dist-linux-gfx1151-${{ matrix.ROCM_VERSION }}.tar.gz
-          mkdir install
-          tar -xf *.tar.gz -C install
-          export ROCM_PATH=$(pwd)/install
-          echo ROCM_PATH=$ROCM_PATH >> $GITHUB_ENV
-          echo PATH=$PATH:$ROCM_PATH/bin >> $GITHUB_ENV
-          echo LD_LIBRARY_PATH=$ROCM_PATH/lib:$ROCM_PATH/llvm/lib:$ROCM_PATH/lib/rocprofiler-systems >> $GITHUB_ENV
+#           # Keep venv activated for subsequent steps
+#           echo "$(pwd)/.venv/bin" >> $GITHUB_PATH

-      - name: Build with native CMake HIP support
-        id: cmake_build
-        run: |
-          cmake -B build -S . \
-            -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
-            -DCMAKE_BUILD_TYPE=Release \
-            -DGGML_BACKEND_DL=ON \
-            -DGGML_NATIVE=OFF \
-            -DCMAKE_INSTALL_RPATH='$ORIGIN' \
-            -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
-            -DGGML_CPU_ALL_VARIANTS=ON \
-            -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
-            -DGGML_HIP=ON \
-            -DHIP_PLATFORM=amd \
-            -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
-            ${{ env.CMAKE_ARGS }}
-          cmake --build build --config Release -j $(nproc)
+#       - name: Build with native CMake HIP support
+#         id: cmake_build
+#         run: |
+#           cmake -B build -S . \
+#             -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
+#             -DCMAKE_BUILD_TYPE=Release \
+#             -DGGML_BACKEND_DL=ON \
+#             -DGGML_NATIVE=OFF \
+#             -DCMAKE_INSTALL_RPATH='$ORIGIN' \
+#             -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+#             -DGGML_CPU_ALL_VARIANTS=ON \
+#             -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
+#             -DGGML_HIP=ON \
+#             -DHIP_PLATFORM=amd \
+#             -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
+#             ${{ env.CMAKE_ARGS }}
+#           cmake --build build --config Release -j $(nproc)

-      - name: ccache-clear
-        uses: ./.github/actions/ccache-clear
-        with:
-          key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+#       # - name: ccache-clear
+#       #   uses: ./.github/actions/ccache-clear
+#       #   with:
+#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}

-      - name: Determine tag name
-        id: tag
-        uses: ./.github/actions/get-tag-name
+#       - name: Determine tag name
+#         id: tag
+#         uses: ./.github/actions/get-tag-name

-      - name: Get ROCm short version
-        run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
+#       - name: Get ROCm short version
+#         run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV

-      - name: Pack artifacts
-        id: pack_artifacts
-        run: |
-          cp LICENSE ./build/bin/
-          tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
+#       - name: Pack artifacts
+#         id: pack_artifacts
+#         run: |
+#           cp LICENSE ./build/bin/
+#           tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .

-      - name: Upload artifacts
-        uses: actions/upload-artifact@v6
-        with:
-          path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
-          name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
-
-  windows-hip:
-    needs: [check-release, get-version]
-    if: ${{ needs.check-release.outputs.should_release == 'true' }}
-
-    runs-on: windows-2022
-
-    permissions:
-      actions: write
-
-    env:
-      HIPSDK_INSTALLER_VERSION: "26.Q1"
-
-    strategy:
-      matrix:
-        include:
-          - name: "radeon"
-            gpu_targets: "gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032"
-
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-
-      - name: Setup Node.js
-        uses: actions/setup-node@v6
-        with:
-          node-version: "24"
-          cache: "npm"
-          cache-dependency-path: "tools/ui/package-lock.json"
-
-      - name: Grab rocWMMA package
-        id: grab_rocwmma
-        run: |
-          curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2.1/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70201-81~24.04_amd64.deb"
-          7z x rocwmma.deb
-          7z x data.tar
-
-      - name: Cache ROCm Installation
-        id: cache-rocm
-        uses: actions/cache@v5
-        with:
-          path: C:\Program Files\AMD\ROCm
-          key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
-
-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.21
-        with:
-          key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
-
-      - name: Install ROCm
-        if: steps.cache-rocm.outputs.cache-hit != 'true'
-        id: depends
-        run: |
-          $ErrorActionPreference = "Stop"
-          write-host "Downloading AMD HIP SDK Installer"
-          Invoke-WebRequest -Uri "https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ env.HIPSDK_INSTALLER_VERSION }}-Win11-For-HIP.exe" -OutFile "${env:RUNNER_TEMP}\rocm-install.exe"
-          write-host "Installing AMD HIP SDK"
-          $proc = Start-Process "${env:RUNNER_TEMP}\rocm-install.exe" -ArgumentList '-install' -NoNewWindow -PassThru
-          $completed = $proc.WaitForExit(600000)
-          if (-not $completed) {
-              Write-Error "ROCm installation timed out after 10 minutes. Killing the process"
-              $proc.Kill()
-              exit 1
-          }
-          if ($proc.ExitCode -ne 0) {
-              Write-Error "ROCm installation failed with exit code $($proc.ExitCode)"
-              exit 1
-          }
-          write-host "Completed AMD HIP SDK installation"
-
-      - name: Verify ROCm
-        id: verify
-        run: |
-          # Find and test ROCm installation
-          $clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1
-          if (-not $clangPath) {
-            Write-Error "ROCm installation not found"
-            exit 1
-          }
-          & $clangPath.FullName --version
-
-      - name: Build
-        id: cmake_build
-        run: |
-          $env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path)
-          $env:CMAKE_PREFIX_PATH="${env:HIP_PATH}"
-          cmake -G "Unix Makefiles" -B build -S . `
-            -DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" `
-            -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" `
-            -DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.1/include/ -Wno-ignored-attributes -Wno-nested-anon-types" `
-            -DCMAKE_BUILD_TYPE=Release `
-            -DGGML_BACKEND_DL=ON `
-            -DGGML_NATIVE=OFF `
-            -DGGML_CPU=OFF `
-            -DGPU_TARGETS="${{ matrix.gpu_targets }}" `
-            -DGGML_HIP=ON `
-            -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} `
-            -DLLAMA_BUILD_BORINGSSL=ON
-          cmake --build build --target ggml-hip -j ${env:NUMBER_OF_PROCESSORS}
-          md "build\bin\rocblas\library\"
-          md "build\bin\hipblaslt\library"
-          cp "${env:HIP_PATH}\bin\libhipblas.dll" "build\bin\"
-          cp "${env:HIP_PATH}\bin\libhipblaslt.dll" "build\bin\"
-          cp "${env:HIP_PATH}\bin\rocblas.dll" "build\bin\"
-          cp "${env:HIP_PATH}\bin\rocblas\library\*" "build\bin\rocblas\library\"
-          cp "${env:HIP_PATH}\bin\hipblaslt\library\*" "build\bin\hipblaslt\library\"
-
-      - name: ccache-clear
-        uses: ./.github/actions/ccache-clear
-        with:
-          key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
-
-      - name: Pack artifacts
-        id: pack_artifacts
-        run: |
-          7z a -snl llama-bin-win-hip-${{ matrix.name }}-x64.zip .\build\bin\*
-
-      - name: Upload artifacts
-        uses: actions/upload-artifact@v6
-        with:
-          path: llama-bin-win-hip-${{ matrix.name }}-x64.zip
-          name: llama-bin-win-hip-${{ matrix.name }}-x64.zip
+#       - name: Upload artifacts
+#         uses: actions/upload-artifact@v6
+#         with:
+#           path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+#           name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz

  ios-xcode:
    needs: [check-release, get-version]
@@ -1572,9 +1576,9 @@ jobs:
      - windows-cpu
      - windows-cuda
      #- windows-sycl
-      - windows-hip
+      - windows-rocm
      - windows-openvino
-      - ubuntu-22-rocm
+      #- ubuntu-22-rocm
      - ubuntu-cpu
      - ubuntu-vulkan
      - ubuntu-24-openvino
@@ -1684,7 +1688,7 @@ jobs:
            - [Ubuntu s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-s390x.tar.gz)
            - [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-x64.tar.gz)
            - [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-arm64.tar.gz)
-            - [Ubuntu x64 (ROCm 7.2)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-7.2-x64.tar.gz)
+            - Ubuntu x64 (ROCm 7.14)[DISABLED](https://github.com/ggml-org/llama.cpp/pull/26969)
            - [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-openvino-${{ needs.ubuntu-24-openvino.outputs.openvino_version }}-x64.tar.gz)
            - [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp32-x64.tar.gz)
            - [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp16-x64.tar.gz)
@@ -1702,7 +1706,7 @@ jobs:
            - [Windows x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-vulkan-x64.zip)
            - [Windows x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-openvino-${{ needs.windows-openvino.outputs.openvino_version }}-x64.zip)
            - [Windows x64 (SYCL)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-sycl-x64.zip)
-            - [Windows x64 (HIP)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-hip-radeon-x64.zip)
+            - [Windows x64 (ROCm 7.14)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-rocm-7.14-x64.zip)

            **openEuler:**
            - [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23705)
@@ -110,7 +110,7 @@ jobs:
          source .venv/bin/activate
          cd tools/server/tests
          export ${{ matrix.extra_args }}
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Slow tests
        id: server_integration_tests_slow
@@ -119,4 +119,4 @@ jobs:
          source .venv/bin/activate
          cd tools/server/tests
          export ${{ matrix.extra_args }}
-          SLOW_TESTS=1 pytest -v -x
+          SLOW_TESTS=1 ./tests.sh
@@ -72,7 +72,7 @@ jobs:
        run: |
          cd tools/server/tests
          source venv/bin/activate
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx1, backend-sampling)
        id: server_integration_tests_backend_sampling
@@ -81,7 +81,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export LLAMA_ARG_BACKEND_SAMPLING=1
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx2)
        id: server_integration_tests_gpu2
@@ -90,7 +90,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export GGML_METAL_DEVICES=2
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx2, backend-sampling)
        id: server_integration_tests_gpu2_backend_sampling
@@ -99,7 +99,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export GGML_METAL_DEVICES=2 LLAMA_ARG_BACKEND_SAMPLING=1
-          pytest -v -x -m "not slow"
+          ./tests.sh

  server-cuda:
    runs-on: [self-hosted, llama-server, Linux, NVIDIA]
@@ -132,7 +132,7 @@ jobs:
        run: |
          cd tools/server/tests
          source venv/bin/activate
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx1, backend-sampling)
        id: server_integration_tests_backend_sampling
@@ -141,7 +141,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export LLAMA_ARG_BACKEND_SAMPLING=1
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx2)
        id: server_integration_tests_gpu2
@@ -150,7 +150,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export GGML_CUDA_DEVICES=2
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Tests (GPUx2, backend-sampling)
        id: server_integration_tests_gpu2_backend_sampling
@@ -159,7 +159,7 @@ jobs:
          cd tools/server/tests
          source venv/bin/activate
          export GGML_CUDA_DEVICES=2 LLAMA_ARG_BACKEND_SAMPLING=1
-          pytest -v -x -m "not slow"
+          ./tests.sh

  server-kleidiai:
    runs-on: ah-ubuntu_22_04-c8g_8x
@@ -219,4 +219,4 @@ jobs:
        run: |
          cd tools/server/tests
          source venv/bin/activate
-          pytest -v -x -m "not slow"
+          ./tests.sh
@@ -104,21 +104,21 @@ jobs:
        id: server_integration_tests
        run: |
          cd tools/server/tests
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Slow tests
        id: server_integration_tests_slow
        if: ${{ github.event.schedule || github.event.inputs.slow_tests == 'true' }}
        run: |
          cd tools/server/tests
-          SLOW_TESTS=1 pytest -v -x
+          SLOW_TESTS=1 ./tests.sh

      - name: Tests (Backend sampling)
        id: server_integration_tests_backend_sampling
        run: |
          cd tools/server/tests
          export LLAMA_ARG_BACKEND_SAMPLING=1
-          pytest -v -x -m "not slow"
+          ./tests.sh

      - name: Slow tests (Backend sampling)
        id: server_integration_tests_slow_backend_sampling
@@ -126,7 +126,7 @@ jobs:
        run: |
          cd tools/server/tests
          export LLAMA_ARG_BACKEND_SAMPLING=1
-          SLOW_TESTS=1 pytest -v -x
+          SLOW_TESTS=1 ./tests.sh

  windows:
    runs-on: windows-2025
@@ -167,15 +167,17 @@ jobs:

      - name: Tests
        id: server_integration_tests
+        shell: bash
        run: |
          cd tools/server/tests
-          $env:PYTHONIOENCODING = ":replace"
-          pytest -v -x -m "not slow"
+          export PYTHONIOENCODING=":replace"
+          ./tests.sh

      - name: Slow tests
        id: server_integration_tests_slow
        if: ${{ github.event.schedule || github.event.inputs.slow_tests == 'true' }}
+        shell: bash
        run: |
          cd tools/server/tests
-          $env:SLOW_TESTS = "1"
-          pytest -v -x
+          export SLOW_TESTS="1"
+          ./tests.sh
@@ -19,6 +19,8 @@ jobs:
        run: |
          cargo binstall komac@2.16.0 -y

+      # TODO: This should later be updated to publish releases instead of
+      # development release builds.
      - name: Find latest release
        id: find_latest_release
        uses: actions/github-script@v8
@@ -2,6 +2,26 @@ cmake_minimum_required(VERSION 3.14...3.28) # for add_link_options and implicit
 project("llama.cpp" C CXX)
 include(CheckIncludeFileCXX)

+### llama.cpp version
+set(LLAMA_VERSION_MAJOR 0)
+set(LLAMA_VERSION_MINOR 1)
+set(LLAMA_VERSION_PATCH 0)
+set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
+
+# whether this is a development/nightly build
+# set this to OFF when making a release from a release tag (vX.Y.Z)
+# ref: https://github.com/ggml-org/ggml/discussions/1579
+option(LLAMA_BUILD_IS_DEV "llama: dev build" ON)
+
+if (LLAMA_BUILD_IS_DEV)
+    set(LLAMA_VERSION "${LLAMA_VERSION_BASE}-dev")
+else()
+    # TODO: check that the current commit is tagged correctly according to the version specified above
+    set(LLAMA_VERSION "${LLAMA_VERSION_BASE}")
+endif()
+
+message(STATUS "llama.cpp version: ${LLAMA_VERSION}")
+
 #set(CMAKE_WARN_DEPRECATED YES)
 set(CMAKE_WARN_UNUSED_CLI YES)

@@ -24,9 +44,6 @@ if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
    set(LLAMA_STANDALONE ON)

    include(git-vars)
-
-    # configure project version
-    # TODO
 else()
    set(LLAMA_STANDALONE OFF)
 endif()
@@ -139,7 +156,6 @@ endif()
 if (NOT DEFINED LLAMA_BUILD_COMMIT)
    set(LLAMA_BUILD_COMMIT        ${BUILD_COMMIT})
 endif()
-set(LLAMA_INSTALL_VERSION 0.0.${LLAMA_BUILD_NUMBER})

 # override ggml options
 set(GGML_ALL_WARNINGS   ${LLAMA_ALL_WARNINGS})
@@ -275,12 +291,12 @@ configure_package_config_file(
              LLAMA_BIN_INSTALL_DIR )

 write_basic_package_version_file(
-        ${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
-    VERSION ${LLAMA_INSTALL_VERSION}
+        ${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
+        VERSION ${LLAMA_VERSION}
    COMPATIBILITY SameMajorVersion)

 install(FILES ${CMAKE_CURRENT_BINARY_DIR}/llama-config.cmake
-              ${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
+              ${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
        DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/llama)

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

 ## Contributing

@@ -1,5 +1,7 @@
 #include "build-info.h"

+#include "llama.h"
+
 #include <cstdio>
 #include <cstdlib>
 #include <string>
@@ -77,12 +79,12 @@ static const command cmds[] = {

 #undef UPDATE_HIDDEN

-static int version(int argc, char ** argv) {
-    printf("%s\n", llama_build_info());
+static int version(int /*argc*/, char ** /*argv*/) {
+    llama_print_build_info(llama_version());
    return 0;
 }

-static int licenses(int argc, char ** argv) {
+static int licenses(int /*argc*/, char ** /*argv*/) {
    for (int i = 0; LICENSES[i]; ++i) {
        printf("%s\n", LICENSES[i]);
    }
@@ -49,6 +49,14 @@ mkdir -p "$2"
 OUT=$(realpath "$1")
 MNT=$(realpath "$2")

+# gpu-rocm self-hosted runner can't upload logs to blob; keep each run's logs in
+# their own dir keyed by the GitHub run id so an Actions run URL maps to its logs.
+if [ -n "${GG_BUILD_ROCM}" ] && [ -n "${GITHUB_RUN_ID}" ]; then
+    OUT="$OUT/run-${GITHUB_RUN_ID}-${GITHUB_RUN_ATTEMPT:-1}"
+    mkdir -p "$OUT"
+    echo "ci results dir: $OUT"
+fi
+
 rm -f $OUT/*.log
 rm -f $OUT/*.exit
 rm -f $OUT/*.md
@@ -1,4 +1,4 @@
-set(LLAMA_VERSION      @LLAMA_INSTALL_VERSION@)
+set(LLAMA_VERSION      @LLAMA_VERSION@)
 set(LLAMA_BUILD_COMMIT @LLAMA_BUILD_COMMIT@)
 set(LLAMA_BUILD_NUMBER @LLAMA_BUILD_NUMBER@)
 set(LLAMA_SHARED_LIB   @BUILD_SHARED_LIBS@)
@@ -5,6 +5,6 @@ includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@

 Name: llama
 Description: Port of Facebook's LLaMA model in C/C++
-Version: @LLAMA_INSTALL_VERSION@
+Version: @LLAMA_VERSION@
 Libs: -L${libdir} -lggml -lggml-base -lllama
 Cflags: -I${includedir}
@@ -121,8 +121,8 @@ add_library(${TARGET}
    )

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

@@ -35,6 +35,7 @@
 #include <regex>
 #include <set>
 #include <string>
+#include <system_error>
 #include <thread> // for hardware_concurrency
 #include <vector>

@@ -704,12 +705,61 @@ void common_models_handler_apply(common_models_handler & handler, common_params
 // CLI argument parsing functions
 //

+// apply config files (if present), a later file overrides an earlier one:
+// 1. system-wide: /etc/llama.cpp/config.ini (%PROGRAMDATA%\llama.cpp\config.ini on windows)
+// 2. user-level: ${XDG_CONFIG_HOME:-~/.config}/llama.cpp/config.ini (%APPDATA%\llama.cpp\config.ini on windows)
+static void common_params_apply_system_config(common_params & params, llama_example ex) {
+    std::vector<std::string> paths;
+
+#if defined(_WIN32)
+    const std::string program_data = common_get_env("PROGRAMDATA");
+    if (!program_data.empty()) {
+        paths.push_back(program_data + "\\llama.cpp\\config.ini");
+    }
+#else
+    paths.push_back("/etc/llama.cpp/config.ini");
+#endif
+
+    try {
+        paths.push_back(fs_get_config_directory() + "config.ini");
+    } catch (const std::exception & e) {
+        LOG_DBG("cannot read user-level config file, skipping: %s\n", e.what());
+    }
+
+    std::vector<std::string> found;
+    for (const auto & path : paths) {
+        std::error_code ec;
+        if (std::filesystem::exists(path, ec)) {
+            found.push_back(path);
+        }
+    }
+    if (found.empty()) {
+        return;
+    }
+
+    common_preset_context ctx(ex);
+    ctx.ignore_unknown_keys = true; // the same config file is shared by all programs
+    for (const auto & path : found) {
+        LOG_INF("using config file: %s\n", path.c_str());
+        common_preset global;
+        common_presets presets = ctx.load_from_ini(path, global);
+        global.apply_to_params(params);
+        auto it = presets.find(COMMON_PRESET_DEFAULT_NAME);
+        if (it != presets.end()) {
+            it->second.apply_to_params(params);
+        }
+    }
+}
+
 static bool common_params_parse_ex(int argc, char ** argv, common_params_context & ctx_arg) {
    common_params & params = ctx_arg.params;

    // setup log directly from params.verbosity: see tools/cli/cli.cpp
    common_log_set_verbosity_thold(params.verbosity);

+    // config file applies first, so env variables and CLI arguments override it
+    common_params_apply_system_config(params, ctx_arg.ex);
+
    std::unordered_map<std::string, std::pair<common_arg *, bool>> arg_to_options;
    for (auto & opt : ctx_arg.options) {
        for (const auto & arg : opt.args) {
@@ -1390,8 +1440,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
        {"--version"},
        "show version and build info",
        [](common_params &) {
-            fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit());
-            fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
+            llama_print_build_info(llama_version());
            exit(0);
        }
    ));
@@ -2605,14 +2654,16 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
    ).set_env("LLAMA_ARG_DIO"));
    add_opt(common_arg(
        {"-lm", "--load-mode"}, "MODE",
-        "model loading mode (default: mmap)\n"
+        "model loading mode (default: auto)\n"
+        "- auto: mmap, unless a device does not support it\n"
        "- none: no special loading mode\n"
        "- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)\n"
        "- mlock: force system to keep model in RAM rather than swapping or compressing\n"
        "- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing\n"
        "- dio: use DirectIO if available\n",
        [](common_params & params, const std::string & value) {
-            /**/ if (value == "none")       { params.load_mode = LLAMA_LOAD_MODE_NONE;       }
+            /**/ if (value == "auto")       { params.load_mode = LLAMA_LOAD_MODE_AUTO;       }
+            else if (value == "none")       { params.load_mode = LLAMA_LOAD_MODE_NONE;       }
            else if (value == "mmap")       { params.load_mode = LLAMA_LOAD_MODE_MMAP;       }
            else if (value == "mlock")      { params.load_mode = LLAMA_LOAD_MODE_MLOCK;      }
            else if (value == "mmap+mlock") { params.load_mode = LLAMA_LOAD_MODE_MMAP_MLOCK; }
@@ -29,7 +29,7 @@ const char * llama_build_info(void) {
    return s.c_str();
 }

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

-void llama_print_build_info(void);
+void llama_print_build_info(const char *);
@@ -1166,6 +1166,16 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
        data.prompt += data.generation_prompt;
    }

+    std::vector<std::string> tool_call_starts = { "<tool_call>" };
+
+    // Match complete <function=name> opener for Qwen3-Coder models that occasionally omit the
+    // starting <tool_call>. The model may hallucinate a tool name, but it is preferable over
+    // constraining on <function which may occur in valid content generation, e.g. #include <functional>
+    foreach_function(inputs.tools, [&](const json & tool) {
+        const std::string name = tool.at("function").at("name");
+        tool_call_starts.push_back("<function=" + name + ">");
+    });
+
    auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) {
        auto generation_prompt = p.literal(GEN_PREFIX);

@@ -1238,7 +1248,7 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
            auto tool_calls = p.trigger_rule("tool-call-root", p.repeat(calls, min_calls, 1));

            return generation_prompt +
-                   (reasoning << p.content(p.until_one_of({ "<tool_call>", "<function=" })) << tool_calls);
+                   (reasoning << p.content(p.until_one_of(tool_call_starts)) << tool_calls);
        }

        // Content only parser
@@ -1264,12 +1274,9 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
        });

        if (data.grammar_lazy) {
-            data.grammar_triggers = {
-                { COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "<tool_call>" },
-                // Trigger on "<function" and not "<function=" because the trailing "=" is part of
-                // the token with the function name e.g. "=read"
-                { COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "<function"   },
-            };
+            for (const auto & start : tool_call_starts) {
+                data.grammar_triggers.push_back({ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, start });
+            }
        }
    }

@@ -3148,7 +3155,8 @@ static common_chat_params common_chat_params_init_muse_glimmer(const common_chat
        auto analysis = p.ref("analysis");

        auto recipient  = p.optional(p.literal(" to=user"));
-        auto final_msg  = p.rule("final", recipient + p.literal("<|message|>") + p.content(p.until("<|eot|>")));
+        auto final_msg  = p.rule("final", recipient + p.literal("<|message|>") +
+                                              p.content(p.until_one_of({ "<|eot|>", "<|eom|>" })));

        if (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE) {
            auto string_value = p.ac(
@@ -3204,7 +3212,8 @@ static common_chat_params common_chat_params_init_muse_glimmer(const common_chat
            if (inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED) {
                return p.zero_or_more(start + analysis) + start + tool_calls;
            }
-            return p.zero_or_more(start + analysis) + start + (tool_calls | final_msg);
+            auto trailing_calls = p.optional(p.literal("<|eom|>") + start + tool_calls);
+            return p.zero_or_more(start + analysis) + start + (tool_calls | (final_msg + trailing_calls));
        }

        return p.zero_or_more(start + analysis) + start + final_msg;
@@ -1019,20 +1019,21 @@ std::string fs_get_cache_directory() {
    std::string cache_directory = "";
    auto ensure_trailing_slash = [](std::string p) {
        // Make sure to add trailing slash
-        if (p.back() != DIRECTORY_SEPARATOR) {
+        if (p.empty() || p.back() != DIRECTORY_SEPARATOR) {
            p += DIRECTORY_SEPARATOR;
        }
        return p;
    };
-    if (getenv("LLAMA_CACHE")) {
-        cache_directory = std::getenv("LLAMA_CACHE");
-    } else {
+    cache_directory = common_get_env("LLAMA_CACHE");
+    if (cache_directory.empty()) {
 #if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \
        defined(__OpenBSD__) || defined(__NetBSD__)
-        if (std::getenv("XDG_CACHE_HOME")) {
-            cache_directory = std::getenv("XDG_CACHE_HOME");
-        } else if (std::getenv("HOME")) {
-            cache_directory = std::getenv("HOME") + std::string("/.cache/");
+        const std::string xdg_cache_home = common_get_env("XDG_CACHE_HOME");
+        const std::string home           = common_get_env("HOME");
+        if (!xdg_cache_home.empty()) {
+            cache_directory = xdg_cache_home;
+        } else if (!home.empty()) {
+            cache_directory = home + "/.cache/";
        } else {
 #if defined(__linux__)
            /* no $HOME is defined, fallback to getpwuid */
@@ -1047,9 +1048,16 @@ std::string fs_get_cache_directory() {
 #endif /* defined(__linux__) */
        }
 #elif defined(__APPLE__)
-        cache_directory = std::getenv("HOME") + std::string("/Library/Caches/");
+        cache_directory = common_get_env("HOME");
+        if (cache_directory.empty()) {
+            throw std::runtime_error("Failed to find $HOME directory");
+        }
+        cache_directory += "/Library/Caches/";
 #elif defined(_WIN32)
-        cache_directory = std::getenv("LOCALAPPDATA");
+        cache_directory = common_get_env("LOCALAPPDATA");
+        if (cache_directory.empty()) {
+            throw std::runtime_error("Failed to find %LOCALAPPDATA% directory");
+        }
 #elif defined(__EMSCRIPTEN__)
        GGML_ABORT("not implemented on this platform");
 #else
@@ -1061,6 +1069,51 @@ std::string fs_get_cache_directory() {
    return ensure_trailing_slash(cache_directory);
 }

+std::string fs_get_config_directory() {
+    std::string config_directory = "";
+    auto ensure_trailing_slash = [](std::string p) {
+        if (p.empty() || p.back() != DIRECTORY_SEPARATOR) {
+            p += DIRECTORY_SEPARATOR;
+        }
+        return p;
+    };
+#if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \
+        defined(__OpenBSD__) || defined(__NetBSD__) || defined(__APPLE__)
+    const std::string xdg_config_home = common_get_env("XDG_CONFIG_HOME");
+    const std::string home            = common_get_env("HOME");
+    if (!xdg_config_home.empty()) {
+        config_directory = xdg_config_home;
+    } else if (!home.empty()) {
+        config_directory = home + "/.config/";
+    } else {
+#if defined(__linux__)
+        /* no $HOME is defined, fallback to getpwuid */
+        struct passwd *pw = getpwuid(getuid());
+        if ((!pw) || (!pw->pw_dir)) {
+            throw std::runtime_error("Failed to find $HOME directory");
+        }
+
+        config_directory = std::string(pw->pw_dir) + std::string("/.config/");
+#else
+        throw std::runtime_error("Failed to find $HOME directory");
+#endif
+    }
+#elif defined(_WIN32)
+    config_directory = common_get_env("APPDATA");
+    if (config_directory.empty()) {
+        throw std::runtime_error("Failed to find %APPDATA% directory");
+    }
+#elif defined(__EMSCRIPTEN__)
+    // caller decides what to do when there is no config directory
+    throw std::runtime_error("not implemented on this platform");
+#else
+#  error Unknown architecture
+#endif
+    config_directory = ensure_trailing_slash(config_directory);
+    config_directory += "llama.cpp";
+    return ensure_trailing_slash(config_directory);
+}
+
 std::string fs_get_cache_file(const std::string & filename) {
    GGML_ASSERT(filename.find(DIRECTORY_SEPARATOR) == std::string::npos);
    std::string cache_directory = fs_get_cache_directory();
@@ -1639,6 +1692,7 @@ struct llama_context_params common_context_params_to_llama(const common_params &
    cparams.n_seq_max         = params.n_parallel;
    cparams.n_rs_seq          = params.speculative.need_n_rs_seq();
    cparams.n_outputs_max     = std::max(params.n_outputs_max, 0);
+    cparams.n_outputs_max_per_seq = std::max(params.n_outputs_max_per_seq, 0);
    cparams.n_batch           = params.n_batch;
    cparams.n_ubatch          = params.n_ubatch;
    cparams.n_threads         = params.cpuparams.n_threads;
@@ -447,6 +447,7 @@ struct common_params {
    int32_t n_parallel            =     1; // number of parallel sequences to decode
    int32_t n_sequences           =     1; // number of sequences to decode
    int32_t n_outputs_max         =     0; // max outputs in a batch (0 = n_batch)
+    int32_t n_outputs_max_per_seq =     1; // max outputs per sequence
    int32_t grp_attn_n            =     1; // group-attention factor
    int32_t grp_attn_w            =   512; // group-attention width
    int32_t n_print               =    -1; // print token count every n tokens (-1 = disabled)
@@ -472,7 +473,7 @@ struct common_params {
    std::vector<size_t> fit_params_target = std::vector<size_t>(llama_max_devices(), 1024 * 1024*1024);

    enum llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER; // how to split the model across GPUs
-    enum llama_load_mode  load_mode  = LLAMA_LOAD_MODE_MMAP; // how to load the model
+    enum llama_load_mode  load_mode  = LLAMA_LOAD_MODE_AUTO; // how to load the model

    common_cpu_params cpuparams;
    common_cpu_params cpuparams_batch;
@@ -880,6 +881,7 @@ bool fs_is_directory(const std::string & path);

 std::string fs_get_cache_directory();
 std::string fs_get_cache_file(const std::string & filename);
+std::string fs_get_config_directory();

 struct common_file_info {
    std::string path;
@@ -116,6 +116,8 @@ static llama_sampler_i llama_sampler_llg_i = {
    /* .backend_accept    = */ NULL,
    /* .backend_apply     = */ NULL,
    /* .backend_set_input = */ NULL,
+    /* .backend_reset     = */ NULL,
+    /* .copy_state        = */ NULL,
 };

 static size_t llama_sampler_llg_tokenize_fn(const void * user_data, const uint8_t * bytes, size_t bytes_len,
@@ -570,23 +570,34 @@ struct parser_executor {
    }

    static common_peg_parse_result handle_escape_sequence(common_peg_parse_context & ctx, size_t start, size_t & pos, const char delimiter) {
+        auto save = pos;
+
        ++pos; // consume '\'
        if (pos >= ctx.input.size()) {
            if (!ctx.is_lenient()) {
                return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
            }
+            pos = save; // suppress unmatched '\'
            return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_NEED_MORE_INPUT, start, pos);
        }

        char c = ctx.input[pos];
+
        if (c == delimiter || c == '\\' || c == '/' || c == 'b' || c == 'f' || c == 'n' || c == 'r' || c == 't') {
            ++pos;
            return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_SUCCESS, start, pos);
-        } else if (c == 'u') {
-            return handle_unicode_escape(ctx, start, pos);
-        } else {
-            return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
        }
+
+        if (c == 'u') {
+            auto result = handle_unicode_escape(ctx, start, pos);
+            if (result.need_more_input()) {
+                pos = save; // suppress incomplete sequence
+                return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_NEED_MORE_INPUT, start, pos);
+            }
+            return result;
+        }
+
+        return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
    }

    static common_peg_parse_result handle_unicode_escape(common_peg_parse_context & ctx, size_t start, size_t & pos) {
@@ -322,6 +322,8 @@ common_presets common_preset_context::load_from_ini(const std::string & path, co
                    preset.options[opt] = value;
                }
                LOG_DBG("accepted option: %s = %s\n", key.c_str(), preset.options[opt].c_str());
+            } else if (ignore_unknown_keys) {
+                LOG_WRN("ignoring option '%s' from %s: not supported by this program\n", key.c_str(), path.c_str());
            } else {
                throw std::runtime_error(string_format(
                    "option '%s' not recognized in preset '%s'",
@@ -59,6 +59,10 @@ struct common_preset_context {
    bool filter_allowed_keys = false;
    std::set<std::string> allowed_keys;

+    // if true, options unknown to the current example are skipped instead of being an error
+    // used for config files shared by all binaries, where each binary only knows a subset of options
+    bool ignore_unknown_keys = false;
+
    // if only_remote_allowed is true, only accept whitelisted keys
    common_preset_context(llama_example ex);

@@ -217,6 +217,8 @@ static struct llama_sampler_i common_reasoning_budget_i = {
    /* .backend_accept    = */ nullptr,
    /* .backend_apply     = */ nullptr,
    /* .backend_set_input = */ nullptr,
+    /* .backend_reset     = */ nullptr,
+    /* .copy_state        = */ nullptr,
 };

 static struct llama_sampler * common_reasoning_budget_clone(const struct llama_sampler * smpl) {
@@ -518,6 +518,26 @@ struct common_sampler * common_sampler_clone(common_sampler * gsmpl) {
    };
 }

+void common_sampler_copy(const common_sampler * src, common_sampler * dst) {
+    if (!src || !dst || src == dst) {
+        return;
+    }
+
+    GGML_ASSERT((src->grmr == nullptr) == (dst->grmr == nullptr));
+    GGML_ASSERT((src->rbudget == nullptr) == (dst->rbudget == nullptr));
+
+    llama_sampler_copy(src->grmr,    dst->grmr);
+    llama_sampler_copy(src->rbudget, dst->rbudget);
+    llama_sampler_copy(src->chain,   dst->chain);
+
+    dst->params     = src->params;
+    dst->prev       = src->prev;
+    dst->cur        = src->cur;
+    dst->cur_p      = src->cur_p;
+    dst->cur_p.data = src->cur_p.data ? dst->cur.data() : nullptr; // re-point to dst's buffer
+    dst->t_total_us = src->t_total_us;
+}
+
 void common_perf_print(const struct llama_context * ctx, const struct common_sampler * gsmpl) {
    // TODO: measure grammar performance

@@ -47,6 +47,7 @@ void common_sampler_free(struct common_sampler * gsmpl);
 void                    common_sampler_accept(struct common_sampler * gsmpl, llama_token token, bool is_generated);
 void                    common_sampler_reset (struct common_sampler * gsmpl);
 struct common_sampler * common_sampler_clone (struct common_sampler * gsmpl);
+void                    common_sampler_copy  (const struct common_sampler * src, struct common_sampler * dst);

 // arguments can be nullptr to skip printing
 void common_perf_print(const struct llama_context * ctx, const struct common_sampler * gsmpl);
@@ -171,12 +171,6 @@ struct common_speculative_impl {
    // (optional) serialize/restore per-seq internal state (e.g. eagle3's deferred boundary).
    virtual bool get_state(llama_seq_id /*seq_id*/, std::vector<uint8_t> & /*data*/) const { return false; }
    virtual void set_state(llama_seq_id /*seq_id*/, const std::vector<uint8_t> & /*data*/) {}
-
-    // true if this implementation requires the target context to extract post-norm embeddings
-    virtual bool need_embd() const = 0;
-
-    // true if this implementation requires the target context to extract pre-norm embeddings
-    virtual bool need_embd_nextn() const { return false; }
 };

 struct common_speculative_impl_draft_simple : public common_speculative_impl {
@@ -193,6 +187,10 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
        auto * ctx_dft = this->params.ctx_dft;
        auto * ctx_tgt = this->params.ctx_tgt;

+        if (!ctx_dft) {
+            throw std::runtime_error("draft-simple requires a draft context");
+        }
+
        SPC_TRC("%s", "adding speculative implementation 'draft-simple'\n");
        SPC_TRC("- n_max=%d, n_min=%d, p_min=%f\n", this->params.n_max, this->params.n_min, this->params.p_min);
        SPC_TRC("- gpu_layers=%d, cache_k=%s, cache_v=%s, ctx_tgt=%s, ctx_dft=%s, devices=[%s]\n",
@@ -385,10 +383,6 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
    void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
        // noop
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };


@@ -907,10 +901,6 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
        pending_g_last[seq_id].resize(n_embd_dec);
        std::memcpy(pending_g_last[seq_id].data(), data.data() + sizeof(llama_pos), (size_t) n_embd_dec * sizeof(float));
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 // DFlash: block-diffusion drafting with a draft-side KV cache injection
@@ -1247,10 +1237,6 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
    void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
        // noop
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 struct common_speculative_impl_draft_mtp : public common_speculative_impl {
@@ -1689,14 +1675,6 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
        const size_t row_bytes = (size_t) n_embd * sizeof(float);
        std::memcpy(pending_h[seq_id].data(), verify_h[seq_id].data() + (size_t) i_h * n_embd, row_bytes);
    }
-
-    bool need_embd() const override {
-        return false;
-    }
-
-    bool need_embd_nextn() const override {
-        return true;
-    }
 };

 // state of self-speculation (simple implementation, not ngram-map)
@@ -1743,10 +1721,6 @@ struct common_speculative_impl_ngram_simple : public common_speculative_impl {
    void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
        // noop
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 struct common_speculative_impl_ngram_map_k : public common_speculative_impl {
@@ -1801,10 +1775,6 @@ struct common_speculative_impl_ngram_map_k : public common_speculative_impl {

        common_ngram_map_accept(config[seq_id], n_accepted);
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 struct common_speculative_impl_ngram_mod : public common_speculative_impl {
@@ -1980,10 +1950,6 @@ struct common_speculative_impl_ngram_mod : public common_speculative_impl {
            }
        }
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 struct common_speculative_impl_ngram_cache : public common_speculative_impl {
@@ -2123,10 +2089,6 @@ struct common_speculative_impl_ngram_cache : public common_speculative_impl {
    void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
        // noop
    }
-
-    bool need_embd() const override {
-        return false;
-    }
 };

 struct common_speculative {
@@ -2299,6 +2261,7 @@ common_params common_base_params_to_speculative(const common_params & params) {
    result.cache_type_k  = params_spec.cache_type_k;
    result.cache_type_v  = params_spec.cache_type_v;
    result.n_outputs_max = params.n_parallel;
+    result.n_outputs_max_per_seq = 1;

    return result;
 }
@@ -2321,7 +2284,6 @@ common_speculative_init_result::common_speculative_init_result(
    const bool spec_mtp = std::find(params.speculative.types.begin(),
                                    params.speculative.types.end(),
                                    COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end();
-    GGML_ASSERT(has_draft || spec_mtp);

    auto mparams = common_model_params_to_llama(params);
    auto cparams = common_context_params_to_llama(params);
@@ -2384,6 +2346,17 @@ common_speculative_init_result_ptr common_speculative_init_from_params(common_pa
    return std::make_unique<common_speculative_init_result>(params, model_tgt, ctx_tgt);
 }

+common_speculative_output_limits common_speculative_get_output_limits(
+        int32_t n_batch, int32_t n_parallel, int32_t n_draft) {
+    const int64_t per_seq = 1 + (int64_t) std::max(0, n_draft);
+    const int64_t total   = (int64_t) n_parallel * per_seq;
+
+    return {
+        /* .total   = */ (int32_t) std::min<int64_t>(n_batch, total),
+        /* .per_seq = */ (int32_t) std::min<int64_t>(n_batch, per_seq),
+    };
+}
+
 // initialization of the speculative decoding system
 //
 common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq) {
@@ -2548,34 +2521,6 @@ bool common_speculative_process(common_speculative * spec, const llama_batch & b
    return result;
 }

-bool common_speculative_need_embd(common_speculative * spec) {
-    if (spec == nullptr) {
-        return false;
-    }
-
-    for (auto & impl : spec->impls) {
-        if (impl->need_embd()) {
-            return true;
-        }
-    }
-
-    return false;
-}
-
-bool common_speculative_need_embd_nextn(common_speculative * spec) {
-    if (spec == nullptr) {
-        return false;
-    }
-
-    for (auto & impl : spec->impls) {
-        if (impl->need_embd_nextn()) {
-            return true;
-        }
-    }
-
-    return false;
-}
-
 void common_speculative_draft(common_speculative * spec) {
    if (spec == nullptr) {
        return;
@@ -2660,7 +2605,10 @@ void common_speculative_draft(common_speculative * spec) {
 void common_speculative_accept(common_speculative * spec, llama_seq_id seq_id, uint16_t n_accepted) {
    common_speculative_impl * impl = spec->impl_last[seq_id];

-    GGML_ASSERT(impl);
+    if (impl == nullptr) {
+        GGML_ASSERT(n_accepted == 0);
+        return;
+    }

    {
        common_time_meas tm(impl->t_accept_us, !impl->gen_perf);
@@ -25,6 +25,15 @@ int32_t common_speculative_n_max(const common_params_speculative * spec);

 common_params common_base_params_to_speculative(const common_params & params);

+struct common_speculative_output_limits {
+    int32_t total;
+    int32_t per_seq;
+};
+
+// return the output limits needed for speculative decoding
+common_speculative_output_limits common_speculative_get_output_limits(
+        int32_t n_batch, int32_t n_parallel, int32_t n_draft);
+
 common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq);

 void common_speculative_free(common_speculative * spec);
@@ -58,12 +67,6 @@ void common_speculative_begin(common_speculative * spec, llama_seq_id seq_id, co
 // process the batch and update the internal state of the speculative context
 bool common_speculative_process(common_speculative * spec, const llama_batch & batch);

-// true if any implementation requires target post-norm embeddings to be extracted
-bool common_speculative_need_embd(common_speculative * spec);
-
-// true if any implementation requires target nextn embeddings to be extracted
-bool common_speculative_need_embd_nextn(common_speculative * spec);
-
 // generate drafts for the sequences specified with `common_speculative_get_draft_params`
 void common_speculative_draft(common_speculative * spec);

@@ -214,6 +214,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
    "Qwen3MoeForCausalLM": "qwen",
    "Qwen3NextForCausalLM": "qwen",
    "Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
+    "PocketTTSModel": "pockettts",
    "Qwen3TTSForConditionalGeneration": "qwen3tts",
    "Qwen3VLForConditionalGeneration": "qwen3vl",
    "Qwen3VLMoeForConditionalGeneration": "qwen3vl",
@@ -310,6 +311,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
    "Qwen2_5_VLForConditionalGeneration": "qwenvl",
    "Qwen3ASRForConditionalGeneration": "qwen3vl",
    "Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
+    "PocketTTSModel": "pockettts",
    "Qwen3TTSForConditionalGeneration": "qwen3tts",
    "Qwen3VLForConditionalGeneration": "qwen3vl",
    "Qwen3VLMoeForConditionalGeneration": "qwen3vl",
@@ -58,6 +58,11 @@ logger = logging.getLogger("hf-to-gguf")
 AnyModel = TypeVar("AnyModel", bound="type[ModelBase]")


+# for checkpoints that ship no config.json, we will try to provide a synthetic one
+HparamsMatcher = Callable[[Path], bool]
+HparamsLoader = Callable[[Path], dict[str, Any]]
+
+
 class SentencePieceTokenTypes(IntEnum):
    NORMAL = 1
    UNKNOWN = 2
@@ -77,6 +82,7 @@ class ModelBase:
        ModelType.TEXT: {},
        ModelType.MMPROJ: {},
    }
+    _hparams_loaders: list[tuple[HparamsMatcher, HparamsLoader]] = []

    dir_model: Path
    ftype: gguf.LlamaFileType
@@ -823,7 +829,7 @@ class ModelBase:
            elif any(str(v.get("quant_algo")).endswith("NVFP4") for v in quant_layers.values() if isinstance(v, dict)):
                quant_algo = "NVFP4"

-        self._is_nvfp4 = quant_algo == "NVFP4"
+        self._is_nvfp4 = quant_algo in ("NVFP4", "W4A16_NVFP4")
        self._is_mxfp4 = quant_method == "mxfp4"

        # NVFP4 weights are repacked and written directly to gguf_writer.
@@ -1040,6 +1046,24 @@ class ModelBase:

        return part_names

+    @staticmethod
+    def load_hparams_guess(dir_model: Path) -> dict[str, Any] | None:
+        # some models ship no config.json, will try to guess them
+        from conversion import load_all_models
+        load_all_models()
+
+        for matcher, loader in ModelBase._hparams_loaders:
+            if matcher(dir_model):
+                return loader(dir_model)
+        return None
+
+    @classmethod
+    def register_hparams_loader(cls, matcher: HparamsMatcher) -> Callable[[HparamsLoader], HparamsLoader]:
+        def inner(loader: HparamsLoader) -> HparamsLoader:
+            cls._hparams_loaders.append((matcher, loader))
+            return loader
+        return inner
+
    @staticmethod
    def load_hparams(dir_model: Path, is_mistral_format: bool):
        if is_mistral_format:
@@ -1053,6 +1077,10 @@ class ModelBase:
            config = AutoConfig.from_pretrained(dir_model, trust_remote_code=False).to_dict()
        except Exception as e:
            logger.warning(f"Failed to load model config from {dir_model}: {e}")
+            if not (dir_model / "config.json").is_file():
+                config = ModelBase.load_hparams_guess(dir_model)
+                if config is not None:
+                    return config
            logger.warning("Trying to load config.json instead")
            with open(dir_model / "config.json", "r", encoding="utf-8") as f:
                config = json.load(f)
@@ -665,7 +665,18 @@ class Gemma4Model(Gemma3Model):
        swa_layers = [t == "sliding_attention" for t in self.hparams["layer_types"]]
        self.gguf_writer.add_sliding_window_pattern(swa_layers)

-        head_dim_full = self.hparams["global_head_dim"]
+        per_layer_config = self.hparams.get("per_layer_config")
+        layer_types = self.hparams.get("layer_types", [])
+        if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
+                        head_dim_full = layer_config["head_dim"]
+                        break
+
+        assert head_dim_full is not None
+
        head_dim_swa = self.hparams["head_dim"]
        # correct the head dim for global/swa layers
        self.gguf_writer.add_key_length(head_dim_full)
@@ -685,8 +696,14 @@ class Gemma4Model(Gemma3Model):
            n_ff_arr = [n_ff if il < first_kv_shared_layer_idx else n_ff * 2 for il in range(self.block_count)]
            self.gguf_writer.add_feed_forward_length(n_ff_arr)

-        # handle num_global_key_value_heads
-        num_key_value_heads_full = self.hparams.get("num_global_key_value_heads")
+        if (num_key_value_heads_full := self.hparams.get("num_global_key_value_heads")) is None and per_layer_config is not None:
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "num_key_value_heads" in layer_config:
+                        num_key_value_heads_full = layer_config["num_key_value_heads"]
+                        break
+
        num_key_value_heads_swa = self.hparams.get("num_key_value_heads")
        if num_key_value_heads_full is not None and num_key_value_heads_swa is not None:
            value_arr = [num_key_value_heads_swa if is_swa else num_key_value_heads_full for is_swa in swa_layers]
@@ -708,7 +725,19 @@ class Gemma4Model(Gemma3Model):
        # IMPORTANT: this ROPE_FREQS tensor is ONLY used by the full_attention layers
        rope_params_full = self.hparams["rope_parameters"]["full_attention"]
        assert rope_params_full["rope_type"] == "proportional"
-        head_dim_full = (self.hparams["global_head_dim"])
+
+        per_layer_config = self.hparams.get("per_layer_config")
+        if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
+            layer_types = self.hparams.get("layer_types", [])
+            for layer_idx, layer_config in per_layer_config.items():
+                layer_idx = int(layer_idx)
+                if layer_idx < len(layer_types):
+                    if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
+                        head_dim_full = layer_config["head_dim"]
+                        break
+
+        assert head_dim_full is not None
+
        partial_rotary_factor_full = rope_params_full["partial_rotary_factor"]
        n_rot_full = int(head_dim_full * partial_rotary_factor_full / 2)
        n_unrot_full = int(head_dim_full / 2) - n_rot_full
@@ -275,10 +275,18 @@ class NemotronHModel(GraniteHybridModel):
                return None
        elif cls.mtp_only:
            # --mtp: export the MTP head plus the tensors it shares with the target model
+            # Include lm_head scale sidecars so NVFP4 packing sees them.
            keep = name in (
                "backbone.embeddings.weight",
                "backbone.norm_f.weight",
                "lm_head.weight",
+                "lm_head.weight_scale",
+                "lm_head.weight_scale_2",
+                "lm_head.weight_scale_inv",
+                "lm_head.input_scale",
+                "lm_head.input_global_scale",
+                "lm_head.weight_global_scale",
+                "lm_head.weight_packed",
            )
            if not keep:
                return None
@@ -0,0 +1,378 @@
+from __future__ import annotations
+
+import re
+from pathlib import Path
+from typing import Any, Iterable, TYPE_CHECKING
+
+import torch
+
+if TYPE_CHECKING:
+    from torch import Tensor
+
+from .base import ModelBase, MmprojModel, SentencePieceTokenTypes, TextModel, gguf, logger
+
+# Pocket TTS is a CALM: the backbone conditions a flow-matching decoder that generates one
+# continuous 32-d latent per frame. There is no codebook in this model.
+# The checkpoint ships no config.json, hparams come from _load_hparams() below.
+#
+# Tricks being used to support this model via existing llama.cpp code paths:
+# - bos_before_voice and bos_emb are learned input vectors, not tokens
+#   they are appended to the embedding table as extra tokens, to be looked up like any other row
+# - bos_emb lives in latent space, so input_linear is folded into it here
+# - the backbone has no lm_head, the embedding table is reused as output for the unused logits
+#
+# pipeline stage mapping:
+#   mimi encoder + speaker_proj --> mapped to normal mtmd audio encoder
+#   flow_lm.transformer         --> mapped to normal libllama text model (autoregressive)
+#   flow_lm.flow_net + out_eos  --> MTMD_GEN_PROCESS_TYPE_GEN_CODE
+#   mimi decoder                --> MTMD_GEN_PROCESS_TYPE_GEN_WAV
+
+# indices into mimi.encoder.model / mimi.decoder.model for stage i, see SEANetEncoder/SEANetDecoder
+_ENC_RES_IDX   = lambda i: 1 + 3 * i  # noqa: E731
+_ENC_SCALE_IDX = lambda i: 3 + 3 * i  # noqa: E731
+_DEC_SCALE_IDX = lambda i: 2 + 3 * i  # noqa: E731
+_DEC_RES_IDX   = lambda i: 3 + 3 * i  # noqa: E731
+
+_N_SEANET_STAGES = 3
+_SAMPLE_RATE = 24000
+
+
+def _tensor_shapes(dir_model: Path) -> dict[str, tuple[int, ...]]:
+    part_names = ModelBase.get_model_part_names(dir_model, "model", ".safetensors")
+    if len(part_names) != 1:
+        return {}
+    with gguf.utility.SafetensorsLocal(dir_model / part_names[0]) as part:
+        return {name: tuple(part[name].shape) for name in part.keys()}
+
+
+@ModelBase.register_hparams_loader(lambda dir_model: "flow_lm.bos_emb" in _tensor_shapes(dir_model))
+def _load_hparams(dir_model: Path) -> dict[str, Any]:
+    logger.info("gguf: detected pocket-tts checkpoint, deriving hparams from tensor shapes")
+    shapes = _tensor_shapes(dir_model)
+    n_vocab, n_embd = shapes["flow_lm.conditioner.embed.weight"]
+    n_layer = sum(1 for name in shapes if re.fullmatch(r"flow_lm\.transformer\.layers\.\d+\.norm1\.weight", name))
+    n_layer_a = sum(1 for name in shapes if re.fullmatch(r"mimi\.encoder_transformer\.transformer\.layers\.\d+\.norm1\.weight", name))
+    n_embd_a = shapes["mimi.encoder_transformer.transformer.layers.0.norm1.weight"][0]
+    return {
+        "architectures": ["PocketTTSModel"],
+        "model_type": "pockettts",
+        "num_hidden_layers": n_layer,
+        "hidden_size": n_embd,
+        "intermediate_size": shapes["flow_lm.transformer.layers.0.linear1.weight"][0],
+        # the transformer is fully causal with no context limit, this only bounds the KV cache
+        "max_position_embeddings": 4096,
+        # not in the checkpoint, but every released variant uses head_dim 64
+        "num_attention_heads": n_embd // 64,
+        # extra rows for the learned input vectors, see _embd_table()
+        "vocab_size": n_vocab + (2 if "flow_lm.bos_before_voice" in shapes else 1),
+        "rope_theta": 10000.0,
+        "layer_norm_eps": 1e-5,
+        "audio_config": {
+            "num_hidden_layers": n_layer_a,
+            "hidden_size": n_embd_a,
+            "intermediate_size": shapes["mimi.encoder_transformer.transformer.layers.0.linear1.weight"][0],
+            "num_attention_heads": n_embd_a // 64,
+        },
+    }
+
+
+@ModelBase.register("PocketTTSModel")
+class PocketTTSModel(TextModel):
+    model_arch = gguf.MODEL_ARCH.POCKETTTS
+
+    _LAYER_TENSOR_MAP = {
+        "norm1":               gguf.MODEL_TENSOR.ATTN_NORM,
+        "norm2":               gguf.MODEL_TENSOR.FFN_NORM,
+        "self_attn.out_proj":  gguf.MODEL_TENSOR.ATTN_OUT,
+        "linear1":             gguf.MODEL_TENSOR.FFN_UP,
+        "linear2":             gguf.MODEL_TENSOR.FFN_DOWN,
+    }
+
+    def set_vocab(self):
+        # this is a unigram sentencepiece model, llama.cpp's SPM tokenizer cannot do
+        # unigram segmentation, so use the UGM tokenizer instead
+        from sentencepiece import sentencepiece_model_pb2 as model
+
+        proto = model.ModelProto()  # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
+        proto.ParseFromString(open(self.dir_model / "tokenizer.model", "rb").read())
+        assert proto.trainer_spec.model_type == 1, "expected a unigram tokenizer"
+
+        tokens, scores, toktypes = self._create_vocab_sentencepiece()
+
+        # the last rows of the embedding table are not sentencepiece pieces
+        extra = self._extra_tokens()
+        for i, name in enumerate(extra):
+            tokens[len(tokens) - len(extra) + i] = name.encode("utf-8")
+            toktypes[len(tokens) - len(extra) + i] = SentencePieceTokenTypes.CONTROL
+            scores[len(tokens) - len(extra) + i] = -1000.0
+
+        self.gguf_writer.add_tokenizer_model("t5")
+        self.gguf_writer.add_tokenizer_pre("default")
+        self.gguf_writer.add_token_list(tokens)
+        self.gguf_writer.add_token_scores(scores)
+        self.gguf_writer.add_token_types(toktypes)
+        self.gguf_writer.add_add_space_prefix(proto.normalizer_spec.add_dummy_prefix)
+        self.gguf_writer.add_remove_extra_whitespaces(proto.normalizer_spec.remove_extra_whitespaces)
+        if proto.normalizer_spec.precompiled_charsmap:
+            self.gguf_writer.add_precompiled_charsmap(proto.normalizer_spec.precompiled_charsmap)
+        self.gguf_writer.add_add_bos_token(False)
+        self.gguf_writer.add_add_eos_token(False)
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        if not name.startswith("flow_lm."):
+            return  # mimi and the flow net go to the mmproj
+
+        if name == "flow_lm.conditioner.embed.weight":
+            yield (self.format_tensor_name(gguf.MODEL_TENSOR.TOKEN_EMBD), self._embd_table(data_torch))
+            return
+
+        if name.startswith("flow_lm.out_norm."):
+            suffix = "." + name.rsplit(".", 1)[1]
+            yield (self.format_tensor_name(gguf.MODEL_TENSOR.OUTPUT_NORM, suffix=suffix), data_torch)
+            return
+
+        if name.startswith("flow_lm.transformer.layers."):
+            assert bid is not None
+            key_with_suffix = name.split(f"layers.{bid}.", 1)[1]
+            key, suffix = key_with_suffix.rsplit(".", 1)
+
+            if key == "self_attn.in_proj":
+                q, k, v = data_torch.chunk(3, dim=0)
+                yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_Q, bid), q)
+                yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_K, bid), k)
+                yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_V, bid), v)
+                return
+
+            tensor = self._LAYER_TENSOR_MAP.get(key)
+            if tensor is not None:
+                yield (self.format_tensor_name(tensor, bid, suffix="." + suffix), data_torch)
+                return
+
+        return
+
+    def _extra_tokens(self) -> list[str]:
+        # the conditioner's padding row, then the learned vectors appended by _embd_table().
+        # bos_before_voice only exists when the pack sets insert_bos_before_voice
+        names = ["<|pad|>"]
+        if "flow_lm.bos_before_voice" in self.model_tensors:
+            names.append("<|bos_before_voice|>")
+        names.append("<|audio_bos|>")
+        return names
+
+    def _embd_table(self, embed: Tensor) -> Tensor:
+        rows = [embed]
+        if "flow_lm.bos_before_voice" in self.model_tensors:
+            rows.append(self.model_tensors["flow_lm.bos_before_voice"]().reshape(1, -1).to(embed.dtype))
+
+        # bos_emb is a latent, it only enters the backbone through input_linear
+        bos_emb = self.model_tensors["flow_lm.bos_emb"]()
+        input_linear = self.model_tensors["flow_lm.input_linear.weight"]()
+        audio_bos = torch.nn.functional.linear(bos_emb.float(), input_linear.float()).reshape(1, -1)
+        rows.append(audio_bos.to(embed.dtype))
+
+        return torch.cat(rows, dim=0)
+
+
+@ModelBase.register("PocketTTSModel")
+class PocketTTSMmprojModel(MmprojModel):
+    has_audio_encoder = True
+    has_vision_encoder = False
+
+    _MIMI_TFM_MAP = {
+        "norm1":              (gguf.MODEL_TENSOR.A_ENC_INPUT_NORM,  gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_NORM),
+        "norm2":              (gguf.MODEL_TENSOR.A_ENC_OUTPUT_NORM, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_NORM),
+        "self_attn.out_proj": (gguf.MODEL_TENSOR.A_ENC_OUTPUT,      gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_OUT),
+        "linear1":            (gguf.MODEL_TENSOR.A_ENC_FFN_UP,      gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_UP),
+        "linear2":            (gguf.MODEL_TENSOR.A_ENC_FFN_DOWN,    gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_DOWN),
+        "layer_scale_1.scale": (gguf.MODEL_TENSOR.A_ENC_ATTN_SCALE, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_SCALE),
+        "layer_scale_2.scale": (gguf.MODEL_TENSOR.A_ENC_FFN_SCALE_LS, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_SCALE),
+    }
+    _MIMI_TFM_QKV = (
+        (gguf.MODEL_TENSOR.A_ENC_ATTN_Q, gguf.MODEL_TENSOR.A_ENC_ATTN_K, gguf.MODEL_TENSOR.A_ENC_ATTN_V),
+        (gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_Q, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_K, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_V),
+    )
+
+    def set_gguf_parameters(self):
+        self.gguf_writer.add_file_type(self.ftype)
+        assert self.hparams_audio is not None
+
+        # voice-prompt encoder: mimi encoder + speaker_proj
+        self.gguf_writer.add_clip_has_audio_encoder(True)
+        # note: the 24kHz sample rate is hardcoded on the clip.cpp side, like the other audio models
+        self.gguf_writer.add_clip_audio_projector_type(gguf.VisionProjectorType.POCKETTTS_SPKENC)
+        self.gguf_writer.add_audio_projection_dim(self.n_embd_text)
+        self.gguf_writer.add_audio_block_count(self.hparams_audio["num_hidden_layers"])
+        self.gguf_writer.add_audio_embedding_length(self.hparams_audio["hidden_size"])
+        self.gguf_writer.add_audio_feed_forward_length(self.hparams_audio["intermediate_size"])
+        self.gguf_writer.add_audio_head_count(self.hparams_audio["num_attention_heads"])
+        self.gguf_writer.add_audio_attention_layernorm_eps(1e-5)
+        # mimi convolves the waveform directly, it is passed around as a 1-row "mel"
+        self.gguf_writer.add_audio_num_mel_bins(1)
+
+        # generation: flow-matching decoder + mimi decoder
+        # the SEANet and flow net hparams are constant across the family, clip.cpp holds them
+        self.gguf_writer.add_clip_has_gen_audio_encoder(True)
+        self.gguf_writer.add_clip_gen_audio_projector_type(gguf.VisionProjectorType.POCKETTTS_GEN)
+        self.gguf_writer.add_gen_audio_projection_dim(self.n_embd_text)
+        self.gguf_writer.add_gen_audio_embedding_length(self.hparams_audio["hidden_size"])
+        self.gguf_writer.add_gen_audio_feed_forward_length(self.hparams_audio["intermediate_size"])
+        self.gguf_writer.add_gen_audio_block_count(self.hparams_audio["num_hidden_layers"])
+        self.gguf_writer.add_gen_audio_head_count(self.hparams_audio["num_attention_heads"])
+        self.gguf_writer.add_gen_audio_attention_layernorm_eps(1e-5)
+
+        self.gguf_writer.add_gen_audio_model_variant(self.dir_model.name)
+
+    def tensor_force_quant(self, name, new_name, bid, n_dims):
+        del name, bid, n_dims
+        # conv1d/conv1d_dw kernels must be F16, ggml_conv_1d(_dw) has no BF16 path
+        if ".seanet." in new_name or new_name in ("a.downsample.conv.weight", "a.gen.wav.upsample.weight"):
+            return gguf.GGMLQuantizationType.F16
+        return False
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        del bid  # the block index of the mimi transformers is parsed here, not by the base class
+        T = gguf.MODEL_TENSOR
+
+        if name in ("flow_lm.bos_emb", "flow_lm.bos_before_voice", "flow_lm.conditioner.embed.weight"):
+            return  # folded into the backbone embedding table
+        if name.startswith("flow_lm.transformer.") or name.startswith("flow_lm.out_norm."):
+            return  # backbone
+
+        if name == "flow_lm.speaker_proj_weight":
+            yield (self.format_tensor_name(T.A_ENC_SPEAKER_PROJ), data_torch)
+            return
+        if name == "flow_lm.input_linear.weight":
+            yield (self.format_tensor_name(T.A_GEN_INPUT_LINEAR), data_torch)
+            return
+        if name == "flow_lm.emb_mean":
+            yield (self.format_tensor_name(T.A_GEN_EMB_MEAN, suffix=""), data_torch)
+            return
+        if name == "flow_lm.emb_std":
+            yield (self.format_tensor_name(T.A_GEN_EMB_STD, suffix=""), data_torch)
+            return
+        if name.startswith("flow_lm.out_eos."):
+            suffix = "." + name.rsplit(".", 1)[1]
+            yield (self.format_tensor_name(T.A_GEN_OUT_EOS, suffix=suffix), data_torch)
+            return
+
+        if name.startswith("flow_lm.flow_net."):
+            yield from self._flow_net_tensor(name, data_torch)
+            return
+
+        if name == "mimi.downsample.conv.conv.weight":
+            yield (self.format_tensor_name(T.A_ENC_DOWNSAMPLE_CONV), data_torch)
+            return
+        if name == "mimi.upsample.convtr.convtr.weight":
+            yield (self.format_tensor_name(T.A_GEN_WAV_UPSAMPLE), data_torch)
+            return
+        if name == "mimi.quantizer.output_proj.weight":
+            yield (self.format_tensor_name(T.A_GEN_WAV_QUANT_OUT), data_torch.squeeze(-1))
+            return
+
+        if "_transformer.transformer.layers." in name:
+            yield from self._mimi_tfm_tensor(name, data_torch)
+            return
+
+        if name.startswith("mimi.encoder.model.") or name.startswith("mimi.decoder.model."):
+            yield from self._seanet_tensor(name, data_torch)
+            return
+
+        return
+
+    def _flow_net_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
+        T = gguf.MODEL_TENSOR
+        key = name.split("flow_lm.flow_net.", 1)[1]
+        suffix = "." + key.rsplit(".", 1)[1]
+
+        simple = {
+            "input_proj":                 T.A_GEN_FLOW_INPUT_PROJ,
+            "cond_embed":                 T.A_GEN_FLOW_COND_EMBD,
+            "final_layer.linear":         T.A_GEN_FLOW_FINAL_PROJ,
+            "final_layer.adaLN_modulation.1": T.A_GEN_FLOW_FINAL_ADA,
+        }
+        tensor = simple.get(key.rsplit(".", 1)[0])
+        if tensor is not None:
+            yield (self.format_tensor_name(tensor, suffix=suffix), data_torch)
+            return
+
+        if key.startswith("time_embed."):
+            bid = int(key.split(".")[1])
+            rest = key.split(f"time_embed.{bid}.", 1)[1]
+            time_map = {
+                "freqs":       (T.A_GEN_FLOW_TIME_FREQS, ""),
+                "mlp.0":       (T.A_GEN_FLOW_TIME_UP,    suffix),
+                "mlp.2":       (T.A_GEN_FLOW_TIME_DOWN,  suffix),
+                "mlp.3.alpha": (T.A_GEN_FLOW_TIME_NORM,  ""),
+            }
+            entry = time_map.get(rest) or time_map.get(rest.rsplit(".", 1)[0])
+            if entry is not None:
+                yield (self.format_tensor_name(entry[0], bid, suffix=entry[1]), data_torch)
+            return
+
+        if key.startswith("res_blocks."):
+            bid = int(key.split(".")[1])
+            rest = key.split(f"res_blocks.{bid}.", 1)[1].rsplit(".", 1)[0]
+            blk_map = {
+                "in_ln":                T.A_GEN_FLOW_BLK_NORM,
+                "mlp.0":                T.A_GEN_FLOW_BLK_UP,
+                "mlp.2":                T.A_GEN_FLOW_BLK_DOWN,
+                "adaLN_modulation.1":   T.A_GEN_FLOW_BLK_ADA,
+            }
+            tensor = blk_map.get(rest)
+            if tensor is not None:
+                yield (self.format_tensor_name(tensor, bid, suffix=suffix), data_torch)
+            return
+
+    def _mimi_tfm_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
+        is_decoder = name.startswith("mimi.decoder_transformer.")
+        bid = int(name.split("_transformer.transformer.layers.", 1)[1].split(".")[0])
+        key_with_suffix = name.split(f".layers.{bid}.", 1)[1]
+
+        if key_with_suffix == "self_attn.in_proj.weight":
+            q, k, v = data_torch.chunk(3, dim=0)
+            names = self._MIMI_TFM_QKV[1 if is_decoder else 0]
+            for tensor, part in zip(names, (q, k, v)):
+                yield (self.format_tensor_name(tensor, bid), part)
+            return
+
+        key, suffix = key_with_suffix.rsplit(".", 1)
+        entry = self._MIMI_TFM_MAP.get(key) or self._MIMI_TFM_MAP.get(key_with_suffix)
+        if entry is None:
+            return
+        tensor = entry[1 if is_decoder else 0]
+        suffix = ".weight" if key_with_suffix.endswith(".scale") else "." + suffix
+        yield (self.format_tensor_name(tensor, bid, suffix=suffix), data_torch)
+
+    def _seanet_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
+        T = gguf.MODEL_TENSOR
+        is_decoder = name.startswith("mimi.decoder.")
+        idx = int(name.split(".model.", 1)[1].split(".")[0])
+        suffix = "." + name.rsplit(".", 1)[1]
+
+        conv_in, conv_out, res1, res2, scale = (
+            (T.A_GEN_WAV_SEANET_CONV_IN, T.A_GEN_WAV_SEANET_CONV_OUT, T.A_GEN_WAV_SEANET_RES_CONV1,
+             T.A_GEN_WAV_SEANET_RES_CONV2, T.A_GEN_WAV_SEANET_SCALE_CONV)
+            if is_decoder else
+            (T.A_ENC_SEANET_CONV_IN, T.A_ENC_SEANET_CONV_OUT, T.A_ENC_SEANET_RES_CONV1,
+             T.A_ENC_SEANET_RES_CONV2, T.A_ENC_SEANET_SCALE_CONV)
+        )
+
+        if idx == 0:
+            yield (self.format_tensor_name(conv_in, suffix=suffix), data_torch)
+            return
+        if idx == 3 * _N_SEANET_STAGES + 2:
+            yield (self.format_tensor_name(conv_out, suffix=suffix), data_torch)
+            return
+
+        for stage in range(_N_SEANET_STAGES):
+            res_idx = _DEC_RES_IDX(stage) if is_decoder else _ENC_RES_IDX(stage)
+            scale_idx = _DEC_SCALE_IDX(stage) if is_decoder else _ENC_SCALE_IDX(stage)
+            if idx == scale_idx:
+                yield (self.format_tensor_name(scale, stage, suffix=suffix), data_torch)
+                return
+            if idx == res_idx:
+                # block.1 is the dilated conv, block.3 the pointwise one (0 and 2 are ELU)
+                inner = int(name.split(".block.", 1)[1].split(".")[0])
+                tensor = res1 if inner == 1 else res2
+                yield (self.format_tensor_name(tensor, stage, suffix=suffix), data_torch)
+                return
@@ -647,10 +647,13 @@ class DFlashModel(Qwen3Model):
        # own tokenizer logic, not the Qwen default).
        from . import get_model_class
        with open(self.target_model_dir / "config.json", "r", encoding="utf-8") as f:
-            target_arch = json.load(f)["architectures"][0]
+            target_hparams = json.load(f)
+            target_arch = target_hparams["architectures"][0]
        target_cls = get_model_class(target_arch)

        if target_cls is not type(self):
+            if target_arch == "NemotronHForCausalLM":
+                setattr(self, "is_moe", "num_experts_per_tok" in target_hparams)
            target_cls.set_vocab(self)  # ty: ignore[unresolved-attribute]
        else:
            super().set_vocab()
@@ -688,6 +691,12 @@ class DFlashModel(Qwen3Model):
            name = "model." + name
        return super().filter_tensors((name, gen))

+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        if name == "model.embed_tokens.weight" and not self.hparams.get("has_embed_tokens", True):
+            return
+
+        yield from super().modify_tensors(data_torch, name, bid)
+

@ModelBase.register("Qwen3DSparkModel")
 class DSparkModel(DFlashModel):
@@ -4,7 +4,7 @@

 The INI preset feature, introduced in [PR#17859](https://github.com/ggml-org/llama.cpp/pull/17859), allows users to create reusable and shareable parameter configurations for llama.cpp.

-### Using Presets with the Server
+## Using Presets with the Server

 When running multiple models on the server (router mode), INI preset files can be used to configure model-specific parameters. Please refer to the [server documentation](../tools/server/README.md) for more details.

@@ -93,3 +93,18 @@ llama-server -hf user/repo:gpt-oss-120b-hf
 ```

 Please make sure to provide the correct `hf-repo` for each child preset. Otherwise, you may get error: `The specified tag is not a valid quantization scheme.`
+
+## System-level config
+
+The system-level config, added in PR [#26118](https://github.com/ggml-org/llama.cpp/pull/26118), allows sharing the same set of options among multiple tools and examples. Unlike the sections above, it is not limited to the server.
+
+These files are loaded on startup if present. A later file overrides an earlier one:
+1. System-wide: `/etc/llama.cpp/config.ini` (or `%PROGRAMDATA%\llama.cpp\config.ini` on Windows)
+2. User-level: `$XDG_CONFIG_HOME/llama.cpp/config.ini`, `~/.config/llama.cpp/config.ini` by default (or `%APPDATA%\llama.cpp\config.ini` on Windows)
+
+The config file is applied first, then its options are overridden by ENV variables, CLI arguments and model presets (in router mode).
+
+Note:
+- Only the `[*]` and default sections are used; options written before any section header belong to "default. Named sections are ignored
+- Tool-specific options can be specified, but will be ignored (with a warning) if the example doesn't support it<br/>Example: if you specify `port = 1234`, only `llama-server` will use it, other examples will ignore it
+- `model` or `hf-repo` are not recommended to be configured system-level, because it may introduce conflicts<br/>Example: a `hf-repo` in the config file still takes effect when you pass `-m` on the command line, so you may load a different model than expected
@@ -0,0 +1,49 @@
+# Release process
+
+llama.cpp uses [semantic versioning](https://semver.org) (`MAJOR.MINOR.PATCH`).
+
+## Version bump guidelines
+
+| Change type | Version component |
+|---|---|
+| Breaking change to the public C API (`include/llama.h`)         | `MAJOR` |
+| Backward-compatible features, model support, or API addition    | `MINOR` |
+| Bug fix with no API change                                      | `PATCH` |
+
+The version is set in the three variables at the top of the root `CMakeLists.txt`:
+
+```cmake
+set(LLAMA_VERSION_MAJOR 0)
+set(LLAMA_VERSION_MINOR 1)
+set(LLAMA_VERSION_PATCH 0)
+```
+
+_A version bump should be included in the PR that introduces the change, or in a
+dedicated bump commit merged before the release is cut._
+
+_TODO: add PR labels (`semver: patch`, `semver: minor`, `semver: major`) to help
+identify which PRs require a version bump before cutting a release._
+
+## Making a release
+
+Releases are created by running the [make-release](.github/workflows/make-release.yml)
+which is a manual workflow.
+
+The workflow creates an annotated git tag (e.g. `v0.1.0`) and pushes it to the
+remote. No GitHub Release object is created, the tag is the release artifact.
+
+## Building a release
+
+By default, `LLAMA_BUILD_IS_DEV=ON` which appends a `-dev` suffix to `LLAMA_VERSION`,
+marking the build as a nightly/development build. Distributors building from a
+release tag must pass `-DLLAMA_BUILD_IS_DEV=OFF` to produce a clean version string
+(e.g. `0.1.0` instead of `0.1.0-dev`).
+
+## How releases reach users
+Currently releases are not published to github releases, only nightly/development
+builds are available there. The way users can access releases are using the following
+channels:
+
+- **llama-install.sh**  — downloads pre-built binaries built from the release tag.
+- **Package managers**  — consume the git tag directly.
+- **Build from source** — users clone the repo and check out the tag.
@@ -202,6 +202,12 @@ Example Video:

 If a draft model is combined with a draftless decoding the draftless decoding has higher precedence.

+### Backend Sampling
+
+Use `--backend-sampling` to run supported target-model samplers on the model backend. Draft-model sampling uses the backend by default and can be controlled with `--spec-draft-backend-sampling` and `--no-spec-draft-backend-sampling`.
+
+Unsupported samplers and device layouts fall back to CPU sampling. Tensor split mode does not support backend sampling. A fixed seed produces repeatable random draws, but stochastic CPU and backend sampling can still select different tokens because floating-point operations can differ between implementations and devices. Use greedy sampling when exact output matching is required.
+
 ### General Speculative Parameters

 ```
@@ -3,9 +3,11 @@
 #include "common.h"
 #include "ngram-cache.h"
 #include "sampling.h"
+#include "speculative.h"
 #include "log.h"
 #include "llama.h"

+#include <algorithm>
 #include <clocale>
 #include <cstdint>
 #include <cstdio>
@@ -27,6 +29,10 @@ int main(int argc, char ** argv){
    // max. number of additional tokens to draft if match is found
    const int n_draft = params.speculative.draft.n_max;

+    const auto output_limits = common_speculative_get_output_limits(params.n_batch, params.n_parallel, n_draft);
+    params.n_outputs_max = output_limits.total;
+    params.n_outputs_max_per_seq = output_limits.per_seq;
+
    // init llama.cpp
    llama_backend_init();
    llama_numa_init(params.numa);
@@ -1,6 +1,6 @@
 --extra-index-url https://download.pytorch.org/whl/cpu
 torch
-torchvision
+torchvision; platform_machine != "s390x"
 transformers
 huggingface-hub
 accelerate
@@ -2,12 +2,15 @@

 import argparse
 import os
+import sys
 import importlib
 import torch
 import numpy as np

 from transformers import AutoTokenizer, AutoConfig, AutoModelForCausalLM
-from pathlib import Path
+
+sys.path.insert(0, os.path.join(os.path.dirname(__file__), '..'))
+from utils.common import save_output_data

 unreleased_model_name = os.getenv('UNRELEASED_MODEL_NAME')

@@ -54,6 +57,7 @@ print(f"Model name: {model_name}")

 prompt = "Hello world today"
 input_ids = tokenizer(prompt, return_tensors="pt").input_ids  # ty: ignore[call-non-callable]
+token_ids = input_ids[0].cpu().tolist()
 print(f"Input tokens: {input_ids}")
 print(f"Input text: {repr(prompt)}")
 print(f"Tokenized: {tokenizer.convert_ids_to_tokens(input_ids[0])}")  # ty: ignore[unresolved-attribute]
@@ -74,21 +78,8 @@ with torch.no_grad():
    print(f"Hidden dimension: {token_embeddings.shape[-1]}")
    print(f"Number of tokens: {token_embeddings.shape[0]}")

-    # Save raw token embeddings
-    data_dir = Path("data")
-    data_dir.mkdir(exist_ok=True)
-    bin_filename = data_dir / f"pytorch-{model_name}-embeddings.bin"
-    txt_filename = data_dir / f"pytorch-{model_name}-embeddings.txt"
-
-    # Save all token embeddings as binary
    print(token_embeddings)
-    token_embeddings.astype(np.float32).tofile(bin_filename)
-
-    # Save as text for inspection
-    with open(txt_filename, "w") as f:
-        for i, embedding in enumerate(token_embeddings):
-            for j, val in enumerate(embedding):
-                f.write(f"{i} {j} {val:.6f}\n")
+    save_output_data(token_embeddings, token_ids, prompt, model_name, type_suffix="-embeddings")

    # Print embeddings per token in the requested format
    print("\nToken embeddings:")
@@ -110,5 +101,3 @@ with torch.no_grad():
    for i, token in enumerate(tokens):
        print(f"  Token {i}: {repr(token)}")

-    print(f"Saved bin logits to: {bin_filename}")
-    print(f"Saved txt logist to: {txt_filename}")
@@ -3,10 +3,47 @@
 Demonstration of basic greedy speculative decoding

 ```bash
+# spec-type draft-simple
 ./bin/llama-speculative-simple \
-    -m  ../models/qwen2.5-32b-coder-instruct/ggml-model-q8_0.gguf \
-    -md ../models/qwen2.5-1.5b-coder-instruct/ggml-model-q4_0.gguf \
-    -f test.txt -c 0 -ngl 99 --color on \
-    --sampling-seq k --top-k 1 -fa on --temp 0.0 \
-    -ngld 99 --spec-draft-n-max 16 --spec-draft-n-draft-min 5 --draft-p-min 0.9
+        -hf  ggml-org/Qwen3-8B-Base-GGUF:Q8_0 \
+        -hfd ggml-org/Qwen3-0.6B-Base-GGUF \
+        -p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-simple --spec-draft-n-max 7 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
+
+# spec-type draft-mtp
+./bin/llama-speculative-simple \
+        -hf ggml-org/Qwen3.6-27B-GGUF:Q8_0 \
+        -p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-mtp --spec-draft-n-max 3 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
+
+# spec-type draft-mtp (with shared KV cache)
+# note: this model needs a <s> token at the start to somewhat work without the chat template
+./bin/llama-speculative-simple \
+        -hf ggml-org/Gemma-4-31B-it-GGUF:Q8_0 \
+        -p "<s>Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-mtp --spec-draft-n-max 3 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
+
+# spec-type draft-eagle3
+./bin/llama-speculative-simple \
+        -hf ggml-org/gpt-oss-20b-GGUF \
+        -p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-eagle3 --spec-draft-n-max 3 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
+
+# spec-type draft-dflash
+./bin/llama-speculative-simple \
+        -hf ggml-org/Qwen3-8B-GGUF \
+        -p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-dflash --spec-draft-n-max 7 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
+
+# spec-type draft-dspark
+./bin/llama-speculative-simple \
+        -hf ggml-org/Qwen3-8B-GGUF \
+        -p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
+        --spec-type draft-dspark --spec-draft-n-max 7 -ngld 99 --color on \
+        -n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
 ```
@@ -5,6 +5,7 @@
 #include "log.h"
 #include "llama.h"

+#include <algorithm>
 #include <clocale>
 #include <cstdio>
 #include <cstring>
@@ -29,6 +30,11 @@ int main(int argc, char ** argv) {
        return 1;
    }

+    const auto output_limits = common_speculative_get_output_limits(
+            params.n_batch, params.n_parallel, common_speculative_n_max(&params.speculative));
+    params.n_outputs_max = output_limits.total;
+    params.n_outputs_max_per_seq = output_limits.per_seq;
+
    // init llama.cpp
    llama_backend_init();
    llama_numa_init(params.numa);
@@ -45,45 +51,23 @@ int main(int argc, char ** argv) {

    const llama_vocab * vocab = llama_model_get_vocab(model_tgt);

-    // load the draft model
-    llama_model_ptr model_dft;
-    llama_context_ptr ctx_dft;
+    // load the draft model (if any) - this also creates the MTP draft context when MTP speculation is enabled
+    common_speculative_init_result_ptr spec_init;

-    // TODO: simplify this logic
    {
-        const auto & params_spec = params.speculative.draft;
+        common_params params_dft = common_base_params_to_speculative(params);

-        auto params_dft = params;
-
-        params_dft.devices      = params_spec.devices;
-        params_dft.model        = params_spec.mparams;
-        params_dft.n_gpu_layers = params_spec.n_gpu_layers;
-
-        if (params_spec.cpuparams.n_threads > 0) {
-            params_dft.cpuparams.n_threads       = params.speculative.draft.cpuparams.n_threads;
-            params_dft.cpuparams_batch.n_threads = params.speculative.draft.cpuparams_batch.n_threads;
-        }
-
-        params_dft.tensor_buft_overrides = params.speculative.draft.tensor_buft_overrides;
-
-        auto mparams_dft = common_model_params_to_llama(params_dft);
-
-        model_dft.reset(llama_model_load_from_file(params_dft.model.path.c_str(), mparams_dft));
-        if (model_dft == nullptr) {
-            LOG_ERR("failed to load draft model, '%s'\n", params_dft.model.path.c_str());
-            return 1;
-        }
-
-        auto cparams = common_context_params_to_llama(params_dft);
-        ctx_dft.reset(llama_init_from_model(model_dft.get(), cparams));
+        spec_init = common_speculative_init_from_params(params_dft, model_tgt, ctx_tgt);

        params.speculative.draft.ctx_tgt = ctx_tgt;
-        params.speculative.draft.ctx_dft = ctx_dft.get();
+        params.speculative.draft.ctx_dft = spec_init->context();
    }

+    llama_context * ctx_dft = params.speculative.draft.ctx_dft;
+
    // check if the context supports partial sequence removal
-    const bool use_ckpt_tgt = (common_context_can_seq_rm(ctx_tgt)       == COMMON_CONTEXT_SEQ_RM_TYPE_FULL);
-    const bool use_ckpt_dft = (common_context_can_seq_rm(ctx_dft.get()) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL);
+    const bool use_ckpt_tgt = common_context_can_seq_rm(ctx_tgt) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL;
+    const bool use_ckpt_dft = common_context_can_seq_rm(ctx_dft) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL;

    if (use_ckpt_tgt) {
        LOG_INF("speculative decoding will use checkpoints (context does not support partial sequence removal)\n");
@@ -129,9 +113,30 @@ int main(int argc, char ** argv) {
    // target model sampling context
    common_sampler_ptr smpl(common_sampler_init(model_tgt, params.sampling));

-    // eval the prompt
-    llama_decode(ctx_tgt,       llama_batch_get_one(inp.data(), inp.size() - 1));
-    llama_decode(ctx_dft.get(), llama_batch_get_one(inp.data(), inp.size() - 1));
+    // init the speculator
+    const auto & params_spec = params.speculative;
+
+    struct common_speculative * spec = common_speculative_init(params.speculative, 1);
+
+    if (spec == nullptr) {
+        LOG_ERR("%s", "failed to initialize speculative decoding\n");
+        return 1;
+    }
+
+    // eval the prompt on the target and feed it to the speculative implementation(s)
+    {
+        llama_batch batch_prompt = llama_batch_init(inp.size(), 0, 1);
+        for (size_t i = 0; i < inp.size() - 1; ++i) {
+            common_batch_add(batch_prompt, inp[i], i, { seq_id }, false);
+        }
+
+        llama_decode(ctx_tgt, batch_prompt);
+
+        if (!common_speculative_process(spec, batch_prompt)) {
+            LOG_ERR("%s", "failed to process speculative prompt\n");
+            return 1;
+        }
+    }

    // note: keep the last token separate!
    llama_token id_last = inp.back();
@@ -142,18 +147,12 @@ int main(int argc, char ** argv) {

    int n_past = inp.size() - 1;

-    // init the speculator
-    const auto & params_spec = params.speculative;
-
-    struct common_speculative * spec = common_speculative_init(params.speculative, 1);
-
    common_speculative_begin(spec, seq_id, prompt_tgt);

    llama_batch batch_tgt = llama_batch_init(llama_n_batch(ctx_tgt), 0, 1);

-    size_t n_draft = 0;
-
    llama_tokens draft;
+
    common_prompt_checkpoint ckpt;

    const auto t_enc_end = ggml_time_us();
@@ -175,13 +174,20 @@ int main(int argc, char ** argv) {
                    llama_memory_seq_pos_max(llama_get_memory(ctx_tgt), seq_id));

            if (use_ckpt_dft) {
-                ckpt.update_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+                ckpt.update_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
            }

+            // determine the max draft that fits the remaining context and generation budget
+            int n_draft_max = (int) llama_n_ctx(ctx_tgt) - n_past - 2;
+            if (params.n_predict >= 0) {
+                n_draft_max = std::min(n_draft_max, params.n_predict - n_predict - 1);
+            }
+            n_draft_max = std::max(n_draft_max, 0);
+
            // generate a new draft
            common_speculative_get_draft_params(spec, seq_id) = {
                /* .drafting   = */ true,
-                /* .n_max      = */ -1,
+                /* .n_max      = */ n_draft_max,
                /* .n_past     = */ n_past,
                /* .id_last    = */ id_last,
                /* .prompt     = */ &prompt_tgt,
@@ -189,9 +195,6 @@ int main(int argc, char ** argv) {
            };
            common_speculative_draft(spec);

-            // save the original draft size
-            n_draft = draft.size();
-
            // save a checkpoint of the target context before evaluating the draft
            // this allows us to restore the state if partial draft acceptance occurs
            if (!draft.empty()) {
@@ -200,10 +203,13 @@ int main(int argc, char ** argv) {
                }
            }

-            {
-                ckpt.load_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+            // reset the draft context to the checkpoint before verification
+            if (ctx_dft) {
+                if (use_ckpt_dft) {
+                    ckpt.load_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+                }

-                llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, ckpt.pos_max + 1, -1);
+                llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, ckpt.pos_max + 1, -1);
            }
        } else {
            // we have a previous (partial) draft to reuse from checkpoint restoration
@@ -227,10 +233,10 @@ int main(int argc, char ** argv) {
            llama_decode(ctx_tgt, batch_tgt);
        }

-        // evaluate the same batch with the draft model
-        {
-            // TODO: extend to support MTP, Eagle, etc. See server code for reference
-            llama_decode(ctx_dft.get(), batch_tgt);
+        // feed the batch to the speculative implementation(s) - this drives the draft model, MTP, Eagle3, etc.
+        if (!common_speculative_process(spec, batch_tgt)) {
+            LOG_ERR("%s", "failed to process speculative batch\n");
+            break;
        }

        // only save the sampler sampler state if we use checkpoints
@@ -239,6 +245,9 @@ int main(int argc, char ** argv) {
            smpl_save.reset(common_sampler_clone(smpl.get()));
        }

+        // save the size of the draft being verified
+        const size_t n_draft = draft.size();
+
        // sample from the full target batch and return the accepted tokens based on the target sampler
        //
        // for each token to be accepted, the sampler would have to sample that same token
@@ -255,8 +264,8 @@ int main(int argc, char ** argv) {
        // check for partial draft acceptance:
        // if the context doesn't support partial sequence removal, restore the checkpoint
        // and make the accepted tokens the new partial draft for the next iteration
-        if (use_ckpt_tgt && ids.size() - 1 < draft.size()) {
-            LOG_DBG("partial acceptance: %zu < %zu, restoring checkpoint\n", ids.size() - 1, draft.size());
+        if (use_ckpt_tgt && ids.size() - 1 < n_draft) {
+            LOG_DBG("partial acceptance: %zu < %zu, restoring checkpoint\n", ids.size() - 1, n_draft);

            draft = std::move(ids);

@@ -266,10 +275,10 @@ int main(int argc, char ** argv) {
                llama_memory_seq_rm(llama_get_memory(ctx_tgt), seq_id, ckpt.pos_max + 1, -1);
            }

-            {
-                ckpt.load_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+            if (ctx_dft) {
+                ckpt.load_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);

-                llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, ckpt.pos_max + 1, -1);
+                llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, ckpt.pos_max + 1, -1);
            }

            prompt_tgt.resize(ckpt.n_tokens);
@@ -320,8 +329,11 @@ int main(int argc, char ** argv) {
        {
            LOG_DBG("clear kv cache from any extra tokens, n_past = %d\n", n_past);

-            llama_memory_seq_rm(llama_get_memory(ctx_tgt),       seq_id, n_past, -1);
-            llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, n_past, -1);
+            llama_memory_seq_rm(llama_get_memory(ctx_tgt), seq_id, n_past, -1);
+
+            if (ctx_dft) {
+                llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, n_past, -1);
+            }
        }

        if ((params.n_predict >= 0 && n_predict > params.n_predict) || has_eos) {
@@ -347,6 +359,7 @@ int main(int argc, char ** argv) {

    LOG_INF("\n");
    LOG_INF("draft:\n\n");
+    common_speculative_print_stats(spec);

    LOG_INF("\n");
    LOG_INF("target:\n\n");
@@ -1,6 +1,7 @@
 #include "arg.h"
 #include "common.h"
 #include "sampling.h"
+#include "speculative.h"
 #include "log.h"
 #include "llama.h"

@@ -57,6 +58,11 @@ int main(int argc, char ** argv) {
    // max number of parallel drafting sequences (i.e. tree branches)
    const int n_seq_dft = params.n_parallel;

+    const auto output_limits = common_speculative_get_output_limits(
+            params.n_batch, params.n_parallel, params.speculative.draft.n_max);
+    params.n_outputs_max = output_limits.total;
+    params.n_outputs_max_per_seq = output_limits.per_seq;
+
    // probability threshold for splitting a draft branch (only for n_seq_dft > 1)
    const float p_draft_split = params.speculative.draft.p_split;

@@ -83,6 +89,8 @@ int main(int argc, char ** argv) {
    params.devices = params.speculative.draft.devices;
    params.model = params.speculative.draft.mparams;
    params.n_gpu_layers = params.speculative.draft.n_gpu_layers;
+    params.n_outputs_max = params.n_parallel;
+    params.n_outputs_max_per_seq = 1;
    if (params.speculative.draft.cpuparams.n_threads > 0) {
        params.cpuparams.n_threads = params.speculative.draft.cpuparams.n_threads;
    }
@@ -0,0 +1,3 @@
+llama-build-install
+install
+build
@@ -0,0 +1,13 @@
+cmake_minimum_required(VERSION 3.14)
+project(llama-simple)
+
+set(CMAKE_CXX_STANDARD 17)
+
+find_package(llama 0.1.0 REQUIRED)
+
+add_executable(test-cmake test-cmake.cpp)
+target_link_libraries(test-cmake PRIVATE llama)
+target_compile_definitions(test-cmake PRIVATE
+    LLAMA_BUILD_NUMBER=${LLAMA_BUILD_NUMBER}
+    LLAMA_BUILD_COMMIT="${LLAMA_BUILD_COMMIT}"
+)
@@ -0,0 +1,36 @@
+## cmake-test
+
+This is just for manually testing/developing of a llama.cpp installation to
+enable troubleshooting issues and exploration. The idea is that this can be used
+after making changes to llama.cpp installation cmake configuration and then
+verify it locally.
+
+### Usage
+The following will configure, build, and install llama.cpp
+
+Configuring/build/install:
+```console
+./build-install.sh
+```
+The above command will create a directory named `install` in the current directory
+which will have the follwing files in its lib directory:
+```console
+(venv) $ ls install/lib/
+cmake                   libggml.so          libllama-common.so.0      libllama.so.0.1.0  llama.cpp
+libggml-base.so         libggml.so.0        libllama-common.so.0.1.0  libmtmd.so         pkgconfig
+libggml-base.so.0       libggml.so.0.19.0   libllama.so               libmtmd.so.0
+libggml-base.so.0.19.0  libllama-common.so  libllama.so.0             libmtmd.so.0.1.0
+```
+
+Build/run this project using the installation created above:
+```console
+(venv) $ ./build.sh
+-- Configuring done (0.0s)
+-- Generating done (0.0s)
+-- Build files have been written to: /home/danbev/work/ai/llama.cpp/examples/test-cmake/build
+[100%] Built target test-cmake
+[test-cmake] Using llama.cpp version 0.1.0-dev-b10335
+[test-cmake] Initializing backend...
+load_backend: loaded CPU backend from /home/danbev/work/ai/llama.cpp/examples/test-cmake/install/lib/llama.cpp/libggml-cpu-alderlake.so
+[test-cmake] Backend initialized.
+```
@@ -0,0 +1,19 @@
+#!/bin/bash
+
+set -e
+
+rm -rf llama-build-install install
+
+cmake --fresh -S ../../. -B llama-build-install -DCMAKE_BUILD_TYPE=Release \
+  -DBUILD_SHARED_LIBS=ON \
+  -DGGML_BACKEND_DL=ON \
+  -DGGML_CPU_ALL_VARIANTS=ON \
+  -DLLAMA_TESTS_INSTALL=OFF \
+  -DCMAKE_INSTALL_PREFIX="${PWD}/install" \
+  -DGGML_BACKEND_DIR="${PWD}/install/lib/llama.cpp" \
+  -DGGML_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
+  -DLLAMA_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
+  -DLLAMA_TOOLS_INSTALL=OFF
+
+cmake --build llama-build-install --parallel 12
+cmake --install llama-build-install
@@ -0,0 +1,7 @@
+#!/bin/bash
+
+set -e
+
+cmake -S . -B build -DCMAKE_PREFIX_PATH="${PWD}/install"
+cmake --build build
+LD_LIBRARY_PATH="${PWD}/install/lib/llama.cpp:${PWD}/install/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" ./build/test-cmake
@@ -0,0 +1,12 @@
+#include "llama.h"
+#include <cstdio>
+
+int main(void) {
+    printf("[test-cmake] version: %s, build: %d (%s)\n",
+           llama_version(), LLAMA_BUILD_NUMBER, LLAMA_BUILD_COMMIT);
+    printf("[test-cmake] Initializing backend...\n");
+    llama_backend_init();
+    printf("[test-cmake] Backend initialized.\n");
+    llama_backend_free();
+    return 0;
+}
@@ -402,7 +402,7 @@ configure_package_config_file(
              GGML_BIN_INSTALL_DIR)

 write_basic_package_version_file(
-        ${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
+        ${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
    VERSION ${GGML_INSTALL_VERSION}
    COMPATIBILITY SameMajorVersion)

@@ -414,7 +414,7 @@ message(STATUS "ggml version: ${GGML_INSTALL_VERSION}")
 message(STATUS "ggml commit:  ${GGML_BUILD_COMMIT}")

 install(FILES ${CMAKE_CURRENT_BINARY_DIR}/ggml-config.cmake
-              ${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
+              ${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
        DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/ggml)

 if (MSVC)
@@ -113,6 +113,7 @@ set_and_check(GGML_LIB_DIR "@PACKAGE_GGML_LIB_INSTALL_DIR@")
 if(NOT TARGET ggml::ggml)
    find_package(Threads REQUIRED)

+    unset(GGML_LIBRARY CACHE)
    find_library(GGML_LIBRARY ggml
        REQUIRED
        HINTS ${GGML_LIB_DIR}
@@ -121,8 +122,10 @@ if(NOT TARGET ggml::ggml)
    add_library(ggml::ggml UNKNOWN IMPORTED)
    set_target_properties(ggml::ggml
        PROPERTIES
-            IMPORTED_LOCATION "${GGML_LIBRARY}")
+            IMPORTED_LOCATION "${GGML_LIBRARY}"
+            INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}")

+    unset(GGML_BASE_LIBRARY CACHE)
    find_library(GGML_BASE_LIBRARY ggml-base
        REQUIRED
        HINTS ${GGML_LIB_DIR}
@@ -132,6 +135,7 @@ if(NOT TARGET ggml::ggml)
    set_target_properties(ggml::ggml-base
        PROPERTIES
            IMPORTED_LOCATION "${GGML_BASE_LIBRARY}"
+            INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}"
            INTERFACE_LINK_LIBRARIES "${GGML_BASE_INTERFACE_LINK_LIBRARIES}")

    set(_ggml_all_targets "")
@@ -140,6 +144,7 @@ if(NOT TARGET ggml::ggml)
            string(REPLACE "-" "_" _ggml_backend_pfx "${_ggml_backend}")
            string(TOUPPER "${_ggml_backend_pfx}" _ggml_backend_pfx)

+            unset(${_ggml_backend_pfx}_LIBRARY CACHE)
            find_library(${_ggml_backend_pfx}_LIBRARY ${_ggml_backend}
                REQUIRED
                HINTS ${GGML_LIB_DIR}
@@ -154,6 +154,8 @@ extern "C" {
        bool buffer_from_host_ptr;
        // event synchronization
        bool events;
+        // mmap is supported for loading
+        bool mmap_support;
    };

    // all the device properties
@@ -132,6 +132,7 @@ static void ggml_backend_meta_device_get_props(ggml_backend_dev_t dev, ggml_back
        /* .host_buffer           = */ false, // Not implemented.
        /* .buffer_from_host_ptr  = */ false, // Not implemented.
        /* .events                = */ false, // Not implemented.
+        /* .mmap_support          = */ true,
    };
    for (ggml_backend_dev_t simple_dev : meta_dev_ctx->simple_devs) {
        ggml_backend_dev_props tmp_props;
@@ -140,6 +141,7 @@ static void ggml_backend_meta_device_get_props(ggml_backend_dev_t dev, ggml_back
        props->caps.host_buffer          = props->caps.host_buffer          && tmp_props.caps.host_buffer;
        props->caps.buffer_from_host_ptr = props->caps.buffer_from_host_ptr && tmp_props.caps.buffer_from_host_ptr;
        props->caps.events               = props->caps.events               && tmp_props.caps.events;
+        props->caps.mmap_support         = props->caps.mmap_support         && tmp_props.caps.mmap_support;
    }
 }

@@ -367,6 +367,7 @@ static void ggml_backend_blas_device_get_props(ggml_backend_dev_t dev, struct gg
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ true,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

@@ -2815,6 +2815,7 @@ static void ggml_backend_cann_device_get_props(ggml_backend_dev_t dev, ggml_back
        /* .host_buffer           = */ host_buffer,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ true,
+        /* .mmap_support          = */ true,
    };
 }

@@ -1,90 +1,12 @@
 #include "ggml-backend-impl.h"
+#include "ggml-feats.h"

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

 static int ggml_backend_cpu_aarch64_score() {
    int score = 1;
-    aarch64_features af;
+    const ggml_feats_arch64_runtime_t af = ggml_feats_get_arch64_runtime();
+    GGML_UNUSED(af);

 #ifdef GGML_USE_DOTPROD
    if (!af.has_dotprod) { return 0; }
@@ -116,4 +38,4 @@ static int ggml_backend_cpu_aarch64_score() {

 GGML_BACKEND_DL_SCORE_IMPL(ggml_backend_cpu_aarch64_score)

-# endif // defined(__aarch64__)
+# endif // defined(__aarch64__) || defined(_M_ARM64)
@@ -397,6 +397,7 @@ static void ggml_backend_cpu_device_get_props(ggml_backend_dev_t dev, struct ggm
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ true,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

@@ -2,10 +2,12 @@
 // SPDX-License-Identifier: MIT
 //
 #include <arm_neon.h>
-#include <assert.h>
-#include <stdio.h>
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
 #include <atomic>
 #include <cfloat>
+#include <cctype>
 #include <algorithm>
 #include <cmath>
 #include <stdexcept>
@@ -17,25 +19,21 @@
 #include <cstddef>
 #include <cstdint>
 #include <fstream>
-#include <set>
+#include <map>
 #include <iostream>
 #include <climits>
+#include <charconv>
+#include <system_error>
 #if defined(__linux__)
 #include <asm/hwcap.h>
+#include <dirent.h>
 #include <sys/auxv.h>
 #include <sys/types.h>
 #include <sys/stat.h>
 #include <unistd.h>
-#ifndef HWCAP2_SME2
-#define HWCAP2_SME2 (1UL << 37)
-#endif
 #elif defined(__APPLE__)
-#include <string_view>
 #include <sys/sysctl.h>
 #include <sys/types.h>
-#elif defined(_WIN32)
-#include <windows.h>
-#include <excpt.h>
 #endif

 #include "kleidiai.h"
@@ -43,6 +41,7 @@
 #include "ggml-cpu.h"
 #include "ggml-cpu-impl.h"
 #include "ggml-impl.h"
+#include "ggml-feats.h"
 #include "ggml-backend-impl.h"
 #include "ggml-threading.h"
 #include "traits.h"
@@ -64,8 +63,8 @@ struct ggml_kleidiai_context {
    ggml_kleidiai_kernels * kernels_q4;
    ggml_kleidiai_kernels * kernels_q8;
    ggml_kleidiai_kernels * kernels_f32;
-    int sme_thread_cap; // <= 0 means “SME disabled/unknown”;
-    int thread_hint;    // <= 0 means “no hint”
+    int sme_thread_cap; // <= 0 means "SME disabled/unknown"
+    int thread_hint;    // <= 0 means "no hint"
    int chunk_multiplier;
 } static ctx = { CPU_FEATURE_NONE, nullptr, nullptr, nullptr, 0, -1, 4 };

@@ -93,24 +92,117 @@ static const char* cpu_feature_to_string(cpu_feature f) {
    }
 }

+#if defined(__linux__) && defined(__aarch64__)
+static bool parse_cpu_dir_name(const char* name, size_t* cpu) {
+    if (strncmp(name, "cpu", 3) != 0 ||
+        name[3] < '0' || name[3] > '9') {
+        return false;
+    }
+
+    const char* first = name + 3;
+    const char* last = name + strlen(name);
+
+    size_t value = 0;
+    const auto [end, ec] = std::from_chars(first, last, value, 10);
+
+    if (ec != std::errc{} || end != last) {
+        return false;
+    }
+
+    *cpu = value;
+    return true;
+}
+
+static std::vector<size_t> detect_cpu_ids() {
+    std::vector<size_t> cpus;
+
+    DIR * dir = opendir("/sys/devices/system/cpu");
+    if (dir == nullptr) {
+        return cpus;
+    }
+
+    while (dirent * entry = readdir(dir)) {
+        size_t cpu = 0;
+        if (parse_cpu_dir_name(entry->d_name, &cpu)) {
+            cpus.push_back(cpu);
+        }
+    }
+    closedir(dir);
+
+    std::sort(cpus.begin(), cpus.end());
+    cpus.erase(std::unique(cpus.begin(), cpus.end()), cpus.end());
+    return cpus;
+}
+#endif
+
+#if defined(__APPLE__) && defined(__aarch64__)
+static bool apple_sme_counted_perf_level(std::string name) {
+    for (std::string::size_type i = 0; i < name.size(); ++i) {
+        name[i] = (char) std::tolower((unsigned char) name[i]);
+    }
+
+    // Conservative ceiling: only count perf-level names observed to provide full SME throughput.
+    // Future names should be calibrated here before they raise the automatic SME thread cap.
+    return name.find("super") != std::string::npos ||
+           name.find("performance") != std::string::npos;
+}
+#endif
+
+static void add_smcus_from_smidr(uint64_t smidr, size_t & num_private, std::map<uint32_t, size_t> & shared_counts) {
+    // Arm ARM: SMIDR_EL1. SH==0 is implementation-defined; keep the existing
+    // conservative policy and only treat zero affinity as private.
+    const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
+    const uint32_t nsmc = (uint32_t)((smidr >> 56) & 0xF);
+    const size_t shared_count = nsmc == 0xF ? 1 : (size_t)nsmc + 1;
+    const uint32_t affinity = (uint32_t)(smidr & 0xFFFu);
+    const uint32_t affinity2 = (uint32_t)((smidr >> 32) & 0xFFFFFu);
+    const uint32_t id = (affinity2 << 12) | affinity;
+
+    if (nsmc == 0xF) {
+        GGML_LOG_WARN("kleidiai: NSMC detected as 0xF indicating reseved value, setting min safe shared SMCU count to 1");
+    }
+
+    switch (sh) {
+        case 2: // private SMCU
+            ++num_private;
+            break;
+        case 3: // shared SMCU
+            if (shared_counts[id] < shared_count) {
+                shared_counts[id] = shared_count;
+            }
+            break;
+        case 0:
+            if (id == 0) {
+                ++num_private;
+            } else if (shared_counts[id] < shared_count) {
+                shared_counts[id] = shared_count;
+            }
+            break;
+        default:
+            break;
+    }
+}
+
 static size_t detect_num_smcus() {
-    if (!ggml_cpu_has_sme()) {
+    const auto runtime_feat = ggml_feats_get_arch64_runtime();
+    if (!runtime_feat.has_sme) {
        return 0;
    }

 #if defined(__linux__) && defined(__aarch64__)
    // Linux/aarch64: Best-effort count of Streaming Mode Compute Units (SMCUs) via SMIDR_EL1 sysfs.
    size_t num_private = 0;
-    std::set<uint32_t> shared_ids;
+    std::map<uint32_t, size_t> shared_counts;

-    for (size_t cpu = 0;; ++cpu) {
+    const std::vector<size_t> cpus = detect_cpu_ids();
+    for (const size_t cpu : cpus) {
        const std::string path =
            "/sys/devices/system/cpu/cpu" + std::to_string(cpu) +
            "/regs/identification/smidr_el1";

        std::ifstream file(path);
        if (!file.is_open()) {
-            break;
+            continue;
        }

        uint64_t smidr = 0;
@@ -118,54 +210,69 @@ static size_t detect_num_smcus() {
            continue;
        }

-        // Arm ARM: SMIDR_EL1
-        const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
-        // Build an "affinity-like" identifier for shared SMCUs.
-        // Keep the original packing logic, but isolate it here.
-        const uint32_t id = (uint32_t)((smidr & 0xFFFu) | ((smidr >> 20) & 0xFFFFF000u));
-
-        switch (sh) {
-            case 0b10: // private SMCU
-                ++num_private;
-                break;
-            case 0b11: // shared SMCU
-                shared_ids.emplace(id);
-                break;
-            case 0b00:
-                // Ambiguous / implementation-defined. Be conservative:
-                // treat id==0 as private, otherwise as shared.
-                if (id == 0) ++num_private;
-                else shared_ids.emplace(id);
-                break;
-            default:
-                break;
-        }
+        add_smcus_from_smidr(smidr, num_private, shared_counts);
    }

-    return num_private + shared_ids.size();
+    size_t total = num_private;
+    for (const auto & entry : shared_counts) {
+        total += entry.second;
+    }
+    return total;

 #elif defined(__APPLE__) && defined(__aarch64__)
-    // table for known M4 variants. Users can override via GGML_KLEIDIAI_SME=<n>.
-    char chip_name[256] = {};
-    size_t size = sizeof(chip_name);
+    int perf_levels = 0;
+    size_t size = sizeof(perf_levels);
+    if (sysctlbyname("hw.nperflevels", &perf_levels, &size, nullptr, 0) != 0 ||
+        size != sizeof(perf_levels) || perf_levels <= 0) {
+        return 0;
+    }

-    if (sysctlbyname("machdep.cpu.brand_string", chip_name, &size, nullptr, 0) == 0) {
-        const std::string brand(chip_name);
+    size_t units = 0;
+    for (int i = 0; i < perf_levels; ++i) {
+        char key[64] = {};
+        int physical_cpus = 0;
+        int cpus_per_l2 = 0;

-        struct ModelSMCU { const char *match; size_t smcus; };
-        static const ModelSMCU table[] = {
-            { "M4 Ultra", 2 },
-            { "M4 Max",   2 },
-            { "M4 Pro",   2 },
-            { "M4",       1 },
-        };
+        snprintf(key, sizeof(key), "hw.perflevel%d.physicalcpu", i);
+        size = sizeof(physical_cpus);
+        if (sysctlbyname(key, &physical_cpus, &size, nullptr, 0) != 0 ||
+            size != sizeof(physical_cpus) || physical_cpus <= 0) {
+            continue;
+        }

-        for (const auto &e : table) {
-            if (brand.find(e.match) != std::string::npos) {
-                return e.smcus;
-            }
+        snprintf(key, sizeof(key), "hw.perflevel%d.cpusperl2", i);
+        size = sizeof(cpus_per_l2);
+        if (sysctlbyname(key, &cpus_per_l2, &size, nullptr, 0) != 0 ||
+            size != sizeof(cpus_per_l2) || cpus_per_l2 <= 0) {
+            continue;
+        }
+
+        snprintf(key, sizeof(key), "hw.perflevel%d.name", i);
+        size = 0;
+        if (sysctlbyname(key, nullptr, &size, nullptr, 0) != 0 || size == 0) {
+            continue;
+        }
+
+        std::string name(size, '\0');
+        if (sysctlbyname(key, &name[0], &size, nullptr, 0) != 0) {
+            continue;
+        }
+        name.resize(size);
+        while (!name.empty() && name.back() == '\0') {
+            name.pop_back();
+        }
+
+        if (apple_sme_counted_perf_level(name)) {
+            units += (size_t) ((physical_cpus + cpus_per_l2 - 1) / cpus_per_l2);
        }
    }
+
+    return units;
+
+#elif defined(_WIN32) && (defined(_M_ARM64) || defined(__aarch64__))
+    // No verified Windows arm64 SMCU detection path yet. Return unknown and use
+    // GGML_KLEIDIAI_SME=N as a diagnostics/debug override for SME thread cap
+    // calibration until a detection mechanism is verified on real hardware.
    return 0;

 #else
@@ -198,15 +305,18 @@ static void init_kleidiai_context(void) {
    if (!initialized) {
        initialized = true;

+        // Optional diagnostics/debug overrides; production defaults come from runtime detection.
        const char *env_sme         = getenv("GGML_KLEIDIAI_SME");
        const char *env_threads     = getenv("GGML_TOTAL_THREADS");
        const char *env_chunk_mult  = getenv("GGML_KLEIDIAI_CHUNK_MULTIPLIER");

+        const auto runtime_feat = ggml_feats_get_arch64_runtime();
+
        size_t detected_smcus = 0;

-        ctx.features  = (ggml_cpu_has_dotprod()     ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
-                        (ggml_cpu_has_matmul_int8() ? CPU_FEATURE_I8MM    : CPU_FEATURE_NONE) |
-                        ((ggml_cpu_has_sve() && ggml_cpu_get_sve_cnt() == QK8_0) ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);
+        ctx.features  = (runtime_feat.has_dotprod  ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
+                        (runtime_feat.has_i8mm     ? CPU_FEATURE_I8MM    : CPU_FEATURE_NONE) |
+                        (runtime_feat.sve_cnt == QK8_0 ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);

        if (env_threads) {
            bool ok = false;
@@ -224,54 +334,54 @@ static void init_kleidiai_context(void) {
            }
        }

-        // SME policy:
-        // - env unset => auto-detect SMCUs; enable SME only if detected > 0.
-        // - env=0     => force off.
-        // - env>0     => force N cores, if the binary was built with SME.
        int sme_cores = 0;
        bool sme_env_ok = false;
        bool sme_env_set = (env_sme != nullptr);

+        const bool has_supported_sme_family = runtime_feat.has_sme;
+        bool sme_cap_detected = false;
+
+        if (has_supported_sme_family) {
+            detected_smcus = detect_num_smcus();
+            sme_cap_detected = detected_smcus > 0;
+            // Some platforms expose SME without exposing a calibrated SMCU count.
+            // Use one SME thread as the conservative default; add platform SMCU detection to raise it.
+            sme_cores = sme_cap_detected ? (int)detected_smcus : 1;
+
+            if (!sme_env_set && !sme_cap_detected) {
+                GGML_LOG_INFO("kleidiai: SME detected; SMCU count unavailable, using conservative SME thread cap=1\n");
+            }
+        }
+
+        // Runtime-detect SME support and available SMCUs first. The detected SMCU
+        // count is used as the SME thread cap, and GGML_KLEIDIAI_SME can debug-override that:
+        //   - unset: use runtime detection.
+        //   - 0:     disable SME-family kernels.
+        //   - N > 0: use N as the SME thread cap, if an SME-family kernel is selectable.
        if (sme_env_set) {
            bool ok = false;
            int v = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
            sme_env_ok = ok;

-            if (!ok) {
-                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; falling back to runtime SME-core detection\n");
-                detected_smcus = detect_num_smcus();
-                sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
-            } else if (v == 0) {
-                sme_cores = 0;
-            } else if (!ggml_cpu_has_sme()) {
-                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but the binary was not built with SME; disabling SME\n", v);
-                sme_cores = 0;
+            if (ok) {
+                if (has_supported_sme_family) {
+                    sme_cores = v;
+                } else {
+                    if (v > 0) {
+                        GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but SME is not supported on this CPU; disabling SME-family kernels\n", v);
+                    }
+                    sme_cores = 0;
+                }
            } else {
-                sme_cores = v;
+                GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; using automatic SME thread cap\n");
            }
-        } else {
-            detected_smcus = detect_num_smcus();
-            sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
        }

-        if (!sme_env_set && ggml_cpu_has_sme() && sme_cores == 0) {
-            GGML_LOG_WARN("kleidiai: runtime SME-core detection returned 0; falling back to NEON\n");
-        }
-
-        if (sme_cores > 0) {
+        if (sme_cores > 0 && has_supported_sme_family) {
            ctx.features |= CPU_FEATURE_SME;
-#if defined(__aarch64__) && defined(__linux__)
-            // ARM guarantees SME2 implies SME, so only check SME2 when SME is enabled.
-            if (getauxval(AT_HWCAP2) & HWCAP2_SME2) {
+            if (runtime_feat.has_sme2) {
                ctx.features |= CPU_FEATURE_SME2;
            }
-#elif defined(__aarch64__) && defined(__APPLE__)
-            int feat_sme2 = 0;
-            size_t size = sizeof(feat_sme2);
-            if (sysctlbyname("hw.optional.arm.FEAT_SME2", &feat_sme2, &size, NULL, 0) == 0 && feat_sme2) {
-                ctx.features |= CPU_FEATURE_SME2;
-            }
-#endif
        }

        // Kernel selection
@@ -297,16 +407,19 @@ static void init_kleidiai_context(void) {
            GGML_LOG_INFO("kleidiai: primary f32 kernel feature %s\n", cpu_feature_to_string(ctx.kernels_f32->required_cpu));
        }

-        ctx.sme_thread_cap = (ctx.features & CPU_FEATURE_SME) ? sme_cores : 0;
+        const bool has_selected_sme_family_kernel =
+            (ctx.kernels_q4  && is_sme_family(ctx.kernels_q4->required_cpu)) ||
+            (ctx.kernels_q8  && is_sme_family(ctx.kernels_q8->required_cpu)) ||
+            (ctx.kernels_f32 && is_sme_family(ctx.kernels_f32->required_cpu));
+        ctx.sme_thread_cap = has_selected_sme_family_kernel ? sme_cores : 0;

-        if (ctx.features & CPU_FEATURE_SME) {
-            const bool has_sme2 = (ctx.features & CPU_FEATURE_SME2) != CPU_FEATURE_NONE;
+        if (has_selected_sme_family_kernel) {
            if (sme_env_set && sme_env_ok && sme_cores > 0) {
-                GGML_LOG_INFO("kleidiai: SME%s enabled (GGML_KLEIDIAI_SME=%d override)\n",
-                    has_sme2 ? "2" : "", sme_cores);
+                GGML_LOG_INFO("kleidiai: SME enabled (GGML_KLEIDIAI_SME=%d debug override)\n", sme_cores);
+            } else if (sme_cap_detected) {
+                GGML_LOG_INFO("kleidiai: SME enabled (runtime-detected SME thread cap=%d)\n", sme_cores);
            } else {
-                GGML_LOG_INFO("kleidiai: SME%s enabled (runtime-detected SME cores=%d)\n",
-                    has_sme2 ? "2" : "", sme_cores);
+                GGML_LOG_INFO("kleidiai: SME enabled (runtime SME detected, conservative thread cap=%d)\n", sme_cores);
            }
        } else {
            GGML_LOG_INFO("kleidiai: SME disabled\n");
@@ -467,7 +580,7 @@ static int kleidiai_collect_kernel_chain_common(
    }

    if (is_sme_family(primary->required_cpu)) {
-        const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~CPU_FEATURE_SME & ~CPU_FEATURE_SME2);
+        const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~(CPU_FEATURE_SME | CPU_FEATURE_SME2));
        if (fallback_mask != CPU_FEATURE_NONE) {
            ggml_kleidiai_kernels * fallback = select_fallback(fallback_mask);
            if (fallback && fallback != primary &&
@@ -1077,13 +1190,14 @@ class tensor_traits : public ggml::cpu::tensor_traits {
        const int ith_total = params->ith;

        int sme_slot = -1;
+        int non_sme_slot = -1;
        for (int i = 0; i < runtime_count; ++i) {
            if (is_sme_family(runtime[i].kernels->required_cpu)) {
                sme_slot = i;
                break;
            }
        }
-        int non_sme_slot = -1;
+
        for (int i = 0; i < runtime_count; ++i) {
            if (!is_sme_family(runtime[i].kernels->required_cpu)) {
                non_sme_slot = i;
@@ -1865,6 +1865,37 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
    ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
 }

+// returns true when ggml_cuda_mul_mat_id takes the fallback path that requires stream synchronization
+// [TAG_MUL_MAT_ID_CUDA_GRAPHS]
+static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int cc) {
+    const ggml_tensor * src0 = dst->src[0];
+    const ggml_tensor * src1 = dst->src[1];
+
+    if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
+        return true;
+    }
+
+    if (dst->ne[2] <= MMVQ_MAX_BATCH_SIZE) {
+        if (ggml_is_quantized(src0->type)) {
+            if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc)) {
+                return false;
+            }
+        } else if (GGML_CUDA_CC_IS_AMD(cc)) {
+            return false;
+        }
+    }
+
+    if (ggml_cuda_should_use_mmq(src0->type, cc, src1->ne[2], /*n_experts=*/src0->ne[2])) {
+        return false;
+    }
+
+    if (ggml_cuda_should_use_mmf(src0->type, cc, WARP_SIZE, src0->ne, src0->nb, src1->ne[2], /*mul_mat_id=*/true)) {
+        return false;
+    }
+
+    return true;
+}
+
 static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
    const ggml_tensor * src0 = dst->src[0];
    const ggml_tensor * src1 = dst->src[1];
@@ -1907,7 +1938,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
    }

    // note: this path should not be reached when recording CUDA graphs, because it requires stream synchronization
-    // TODO: add asserts to verify this. should work with CUDA, HIP, etc.
+    GGML_ASSERT(ggml_cuda_mul_mat_id_needs_sync(dst, cc));
    cudaStream_t stream = ctx.stream();

    GGML_ASSERT(nb12 % nb11 == 0);
@@ -2522,10 +2553,8 @@ static bool ggml_cuda_graph_check_compability(ggml_cgraph * cgraph) {
        // [TAG_MUL_MAT_ID_CUDA_GRAPHS]
        if (node->op == GGML_OP_MUL_MAT_ID) {
            const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
-            const int mmvq_mmid_max = get_mmvq_mmid_max_batch(node->src[0]->type, cc);
-            if (!ggml_is_quantized(node->src[0]->type) || node->ne[2] > mmvq_mmid_max) {
-                // under these conditions, the mul_mat_id operation will need to synchronize the stream, so we cannot use CUDA graphs
-                // TODO: figure out a way to enable for larger batch sizes, without hurting performance
+            if (ggml_cuda_mul_mat_id_needs_sync(node, cc)) {
+                // the mul_mat_id fallback path synchronizes the stream, so we cannot use CUDA graphs
                // ref: https://github.com/ggml-org/llama.cpp/pull/18958
                use_cuda_graph = false;
 #ifndef NDEBUG
@@ -4801,6 +4830,7 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back
        /* .host_buffer           = */ host_buffer,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ events,
+        /* .mmap_support          = */ props->type != GGML_BACKEND_DEVICE_TYPE_IGPU,
    };
 }

@@ -141,6 +141,57 @@ static __global__ void rwkv_wkv7_f32(const int B, const int T, const int C, cons
    }
 }

+template <int rows_per_block>
+static __global__ void __launch_bounds__(WARP_SIZE * rows_per_block, 2)
+rwkv_wkv7_f32_t1_warp_row(const int T, const int C, const int H, const float * r, const float * w, const float * k, const float * v, const float * a, const float * b, const float * s, float * dst) {
+    constexpr int head_size = CUDA_WKV_BLOCK_SIZE;
+    constexpr int half_head = head_size / 2;
+
+    const int lane = threadIdx.x;
+    const int row  = blockIdx.y * rows_per_block + threadIdx.y;
+    const int bid  = blockIdx.x;
+
+    const int batch_i = bid / H;
+    const int head_i  = bid % H;
+    const int state_size = C * head_size;
+    const int head_off = head_i * head_size;
+    const int t = batch_i * C + head_off + row;
+
+    __shared__ float _r[head_size], _w[head_size], _k[head_size], _a[head_size], _b[head_size];
+
+    if (threadIdx.y == 0) {
+        _r[lane] = r[batch_i * C + head_off + lane];
+        _w[lane] = w[batch_i * C + head_off + lane];
+        _k[lane] = k[batch_i * C + head_off + lane];
+        _a[lane] = a[batch_i * C + head_off + lane];
+        _b[lane] = b[batch_i * C + head_off + lane];
+
+        _r[lane + half_head] = r[batch_i * C + head_off + lane + half_head];
+        _w[lane + half_head] = w[batch_i * C + head_off + lane + half_head];
+        _k[lane + half_head] = k[batch_i * C + head_off + lane + half_head];
+        _a[lane + half_head] = a[batch_i * C + head_off + lane + half_head];
+        _b[lane + half_head] = b[batch_i * C + head_off + lane + half_head];
+    }
+    __syncthreads();
+
+    const int64_t state_base = batch_i * state_size + head_i * head_size * head_size + row * head_size;
+    const float s0 = s[state_base + lane];
+    const float s1 = s[state_base + lane + half_head];
+    const float sa = warp_reduce_sum(_a[lane] * s0 + _a[lane + half_head] * s1);
+
+    const float vt  = v[t];
+    const float st0 = s0 * _w[lane]             + _k[lane]             * vt + sa * _b[lane];
+    const float st1 = s1 * _w[lane + half_head] + _k[lane + half_head] * vt + sa * _b[lane + half_head];
+    const float y   = warp_reduce_sum(st0 * _r[lane] + st1 * _r[lane + half_head]);
+
+    dst[T * C + state_base + lane]             = st0;
+    dst[T * C + state_base + lane + half_head] = st1;
+
+    if (lane == 0) {
+        dst[t] = y;
+    }
+}
+
 void ggml_cuda_op_rwkv_wkv6(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
    const float * k_d  = (const float *)dst->src[0]->data;
    const float * v_d  = (const float *)dst->src[1]->data;
@@ -191,7 +242,10 @@ void ggml_cuda_op_rwkv_wkv7(ggml_backend_cuda_context & ctx, ggml_tensor * dst)
    GGML_ASSERT(C % H == 0);
    GGML_ASSERT(C / H == CUDA_WKV_BLOCK_SIZE || C / H == CUDA_WKV_BLOCK_SIZE * 2);

-    if (C / H == CUDA_WKV_BLOCK_SIZE) {
+    if (T / B == 1 && C / H == CUDA_WKV_BLOCK_SIZE) {
+        constexpr int rows_per_block = 4;
+        rwkv_wkv7_f32_t1_warp_row<rows_per_block><<<dim3(B * H, CUDA_WKV_BLOCK_SIZE / rows_per_block), dim3(WARP_SIZE, rows_per_block), 0, stream>>>(T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
+    } else if (C / H == CUDA_WKV_BLOCK_SIZE) {
        rwkv_wkv7_f32<CUDA_WKV_BLOCK_SIZE><<<B * H, C / H, 0, stream>>>(B, T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
    } else {
        rwkv_wkv7_f32<CUDA_WKV_BLOCK_SIZE * 2><<<B * H, C / H, 0, stream>>>(B, T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
@@ -1646,6 +1646,7 @@ static void ggml_backend_et_device_get_props(ggml_backend_dev_t dev, struct ggml
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

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

@@ -681,6 +681,7 @@ static void ggml_backend_metal_device_get_props(ggml_backend_dev_t dev, ggml_bac
        /* .host_buffer          = */ false,
        /* .buffer_from_host_ptr = */ true,
        /* .events               = */ true,
+        /* .mmap_support         = */ true,
    };
 }

@@ -73,6 +73,7 @@ typedef const void * (*get_adreno_bin_kernel_func_t)(
 //------------------------------------------------------------------------------

 bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor);
+
 static bool ggml_cl_is_q4_0_soa(const ggml_tensor * tensor);
 static bool ggml_cl_is_q8_0_soa(const ggml_tensor * tensor);
 static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
@@ -4629,6 +4630,23 @@ static std::string ggml_opencl_fa_compile_opts(ggml_backend_opencl_context * bac
    if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E) {
        opts += " -D FA_C8_NO_SG_PIN";
    }
+    // Transposed K tile in local memory: the KV rows the QK loop walks together become
+    // adjacent, so a group of them is ONE 128-bit local read instead of several narrow
+    // ones. The QK loop is LDS-read-issue-bound (a wrong-math probe that kept every FMA/dp4a
+    // but removed the LDS reads ran the kernel ~40% faster), so this is worth up to +26% on
+    // fa=1 prefill. Output is bit-identical -- only the layout moves.
+    //
+    // DK <= 128 only. At DK=256 (gemma-3-4b) it measures 1-2% NEGATIVE and reproduces across
+    // rounds; padding the row stride does not recover it, so the cause is not a simple bank
+    // conflict and the wider tile does not want this layout.
+    //
+    // Default on within that gate; GGML_OPENCL_FA_K_LDS_T=0 restores the row-major tile.
+    {
+        const char * e = getenv("GGML_OPENCL_FA_K_LDS_T");
+        if ((e == nullptr || e[0] != '0') && cfg->dk <= 128) {
+            opts += " -D FA_K_LDS_T";
+        }
+    }
    return opts;
 }

@@ -4911,8 +4929,13 @@ static bool ggml_opencl_ensure_fa_variant(ggml_backend_opencl_context * backend_
        const int x = (e && e[0]) ? atoi(e) : 0;
        return (x == 8 || x == 16 || x == 32) ? x : 0;   // 0 = per-gen default
    }();
+    // X2E needs 16 to keep per-lane o_acc at 128B (the compiler spills the
+    // kernel-default width); X1E does not spill, but C=16 is still a measured
+    // +28-30% DK128-GQA4 decode win there (X1-85, kv 4096/8192), neutral on
+    // DK64 / GQA1 / quant-KV.
    const int fa_cl_c_gqa4 = fa_cl_c_env ? fa_cl_c_env
-        : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ? 16 : 0);
+        : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ||
+           backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E ? 16 : 0);
    const std::string opts_cl_c_gqa4 = fa_cl_c_gqa4
        ? " -D FA_CL_C=" + std::to_string(fa_cl_c_gqa4) : std::string();
    const std::string fa_cl_c_g8_val = std::to_string(fa_cl_c_gqa4 ? fa_cl_c_gqa4 * 2 : 16);
@@ -7058,6 +7081,19 @@ inline bool enable_adreno_trans_weight(const ggml_backend_opencl_context *backen
    return ((elem_num < 128 * 1024 * 1024) && adreno_kernel && shape_ok);  // max element num: 2**27
 }

+inline bool enable_adreno_trans_weight_q5_K(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
+    if (!use_adreno_kernels(backend_ctx, tensor)) {
+        return false;
+    }
+
+    const size_t elem_num = ggml_nelements(tensor);
+    const size_t q_img_width = elem_num / 8;
+    const size_t qh_img_width = elem_num / 16;
+
+    return q_img_width <= backend_ctx->image_max_buffer_size &&
+           qh_img_width <= backend_ctx->image_max_buffer_size;
+}
+
 static inline bool use_flat_gemv_for_large_m_q4_K(const ggml_tensor *tensor) {
    // gemv_noshuffle variant perf drops for large M, use flat variant for large M.
    // threshold is well above typical hidden/FFN dims, but below typical vocab sizes.
@@ -9237,7 +9273,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,

 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
        cl_kernel kernel = backend_ctx->kernel_convert_block_q5_K;
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {
            kernel = backend_ctx->kernel_convert_block_q5_K_noshuffle;
        }
 #else
@@ -9272,7 +9308,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,

        tensor->extra = extra;
 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {

            int M = tensor->ne[1];
            int K = tensor->ne[0];
@@ -10370,7 +10406,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
            CL_CHECK(clReleaseMemObject(data_device));
            return;
        }
-        if (use_adreno_kernels(backend_ctx, tensor)) {
+        if (enable_adreno_trans_weight_q5_K(backend_ctx, tensor)) {
            int M = tensor->ne[1];
            int K = tensor->ne[0];

@@ -10777,6 +10813,7 @@ static void ggml_backend_opencl_device_get_props(ggml_backend_dev_t dev, struct
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ false,
+        /* .mmap_support          = */ false,
    };
 }

@@ -18909,7 +18946,8 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
        }

        // q5_K x fp32
-        if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32) {
+        if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32 &&
+            enable_adreno_trans_weight_q5_K(backend_ctx, src0)) {
            ggml_cl_mul_mat_q5_K_f32_adreno(backend, src0, src1, dst);
            return;
        }
@@ -211,7 +211,30 @@ __kernel void FA_TILE_NAME(

    float slope = get_alibi_slope(max_bias, head_idx, n_head_log2, m0, m1);

+#ifdef FA_K_LDS_T
+    // K tile transposed: [dk vec][kv row] instead of [kv row][dk vec].
+    //
+    // The QK loop walks 2 or 4 KV rows at a time against the same dk element. Row-major
+    // those are DK_VEC half4s apart, so each is its own 64-bit local read. Transposed they
+    // are adjacent, so a pair is one 128-bit read -- half the LDS issues for the same bytes,
+    // no extra registers, arithmetic untouched.
+    //
+    // This kernel looked like it should be FMA-bound (a half4 mad does ~4 ALU ops per LDS
+    // read, unlike the 1:1 of the dp4a loop), but it is NOT: a wrong-math probe that kept
+    // every FMA and removed the LDS reads ran it 38.6% faster (18.92 -> 11.62 ms/op).
+    // Explicitly 16-byte aligned: FA_LK_PAIR below reads two adjacent half4 as one float4,
+    // and the element type only obliges the compiler to align this array to 8. The indices
+    // are even so the offset is a multiple of 16, but the base has to be too, and relying
+    // on the compiler to over-align it is relying on luck.
+    __local KV_DATA_TYPE4 l_k[DK_VEC][BLOCK_N] __attribute__((aligned(16)));
+#define FA_LK(ROW, C) l_k[C][ROW]
+    // Two adjacent KV rows as one 128-bit local read (half4 pair == 16 B). j is even and
+    // BLOCK_N is even, so &l_k[c][j] is 16 B past a 16 B-aligned base.
+#define FA_LK_PAIR(C, J) as_half8(*(__local const float4 *)(&l_k[C][J]))
+#else
    __local KV_DATA_TYPE4 l_k[BLOCK_N][DK_VEC];
+#define FA_LK(ROW, C) l_k[ROW][C]
+#endif
    __local KV_DATA_TYPE4 l_v[BLOCK_N][DV_VEC];

 #if N_SPLIT > 1 && !defined(HAS_SUBGROUP_SHUFFLE)
@@ -254,17 +277,17 @@ __kernel void FA_TILE_NAME(
 #ifdef FA_K_IMG
                if (use_kv_pad) {
                    const ulong k_row_offset = batch_idx * k_tile_nb3 + head_kv_idx * k_tile_nb2 + k_row_idx * k_nb1;
-                    l_k[row][col] = ((__global KV_DATA_TYPE4*)(k_tile_base + k_row_offset))[col];
+                    FA_LK(row, col) = ((__global KV_DATA_TYPE4*)(k_tile_base + k_row_offset))[col];
                } else {
                    const int k_row_px = batch_idx * k_pitch_px_batch + head_kv_idx * k_pitch_px_head + k_row_idx * k_pitch_px_row;
-                    l_k[row][col] = read_imageh(k_img, k_row_px + col);
+                    FA_LK(row, col) = read_imageh(k_img, k_row_px + col);
                }
 #else
                const ulong k_row_offset = batch_idx * k_tile_nb3 + head_kv_idx * k_tile_nb2 + k_row_idx * k_nb1;
-                l_k[row][col] = ((__global KV_DATA_TYPE4*)(k_tile_base + k_row_offset))[col];
+                FA_LK(row, col) = ((__global KV_DATA_TYPE4*)(k_tile_base + k_row_offset))[col];
 #endif
            } else {
-                l_k[row][col] = (KV_DATA_TYPE4)(0.0h);
+                FA_LK(row, col) = (KV_DATA_TYPE4)(0.0h);
            }
        }
        for (int i = tid; i < BLOCK_N * DV_VEC; i += WG_SIZE) {
@@ -292,8 +315,15 @@ __kernel void FA_TILE_NAME(
                FA_UNROLL
                for (int k = 0; k < SPLIT_DK_VEC; k++) {
                    const ACC_TYPE4 qk = q_priv[k];
+#if defined(FA_K_LDS_T)
+                    // 2 KV rows adjacent in the transposed tile: one 128-bit local read.
+                    const half8 kk = FA_LK_PAIR(dk_off + k, j);
+                    ACC_TYPE4 dot0 = qk * CONVERT_KV_ACC4(kk.lo);
+                    ACC_TYPE4 dot1 = qk * CONVERT_KV_ACC4(kk.hi);
+#else
                    ACC_TYPE4 dot0 = qk * CONVERT_KV_ACC4(l_k[j  ][dk_off + k]);
                    ACC_TYPE4 dot1 = qk * CONVERT_KV_ACC4(l_k[j+1][dk_off + k]);
+#endif
                    partial0 += dot0.s0 + dot0.s1 + dot0.s2 + dot0.s3;
                    partial1 += dot1.s0 + dot1.s1 + dot1.s2 + dot1.s3;
                }
@@ -359,7 +389,7 @@ __kernel void FA_TILE_NAME(
            ACC_TYPE4 dot_acc = (ACC_TYPE4)(0.0f);
            FA_UNROLL
            for (int k = 0; k < SPLIT_DK_VEC; k++) {
-                dot_acc = mad(q_priv[k], CONVERT_KV_ACC4(l_k[j][dk_off + k]), dot_acc);
+                dot_acc = mad(q_priv[k], CONVERT_KV_ACC4(FA_LK(j, dk_off + k)), dot_acc);
            }
            local_partial[j][tid] =
                dot_acc.s0 + dot_acc.s1 + dot_acc.s2 + dot_acc.s3;
@@ -452,10 +482,21 @@ __kernel void FA_TILE_NAME(
                FA_UNROLL
                for (int k = 0; k < DK_VEC; k++) {
                    const ACC_TYPE4 qk = q_priv[k];
+#if defined(FA_K_LDS_T)
+                    // 4 KV rows adjacent in the transposed tile: two 128-bit local reads
+                    // instead of four 64-bit ones.
+                    const half8 kk01 = FA_LK_PAIR(k, j);
+                    const half8 kk23 = FA_LK_PAIR(k, j + 2);
+                    dot_acc0 = mad(qk, CONVERT_KV_ACC4(kk01.lo), dot_acc0);
+                    dot_acc1 = mad(qk, CONVERT_KV_ACC4(kk01.hi), dot_acc1);
+                    dot_acc2 = mad(qk, CONVERT_KV_ACC4(kk23.lo), dot_acc2);
+                    dot_acc3 = mad(qk, CONVERT_KV_ACC4(kk23.hi), dot_acc3);
+#else
                    dot_acc0 = mad(qk, CONVERT_KV_ACC4(l_k[j][k]),   dot_acc0);
                    dot_acc1 = mad(qk, CONVERT_KV_ACC4(l_k[j+1][k]), dot_acc1);
                    dot_acc2 = mad(qk, CONVERT_KV_ACC4(l_k[j+2][k]), dot_acc2);
                    dot_acc3 = mad(qk, CONVERT_KV_ACC4(l_k[j+3][k]), dot_acc3);
+#endif
                }
                ACC_TYPE s0 = (dot_acc0.s0 + dot_acc0.s1 + dot_acc0.s2 + dot_acc0.s3) * scale;
                ACC_TYPE s1 = (dot_acc1.s0 + dot_acc1.s1 + dot_acc1.s2 + dot_acc1.s3) * scale;
@@ -1631,8 +1631,25 @@ __kernel void flash_attn_f32_q4_0(
    float slope = get_alibi_slope(max_bias, head_idx, n_head_log2, m0, m1);

 #ifdef FA_HAVE_INT_DOT
+// Accessors so the staging code is layout-agnostic.
+#ifdef FA_K_LDS_T
+#define FA_K_PACKED(ROW, IDX) l_k_packed[IDX][ROW]
+#define FA_K_SCALE(ROW, BLK)  l_k_scale[BLK][ROW]
+#else
+#define FA_K_PACKED(ROW, IDX) l_k_packed[ROW][IDX]
+#define FA_K_SCALE(ROW, BLK)  l_k_scale[ROW][BLK]
+#endif
+
+#ifdef FA_K_LDS_T
+    // K tile transposed: the 4 KV rows the QK loop walks together become adjacent, so each
+    // (block, group) step is ONE 128-bit local read instead of four 32-bit ones. The QK
+    // loop is LDS-read-issue-bound.
+    __local uint  l_k_packed[DK_Q4_BLOCKS_PREFILL * 8][BLOCK_N];
+    __local float l_k_scale [DK_Q4_BLOCKS_PREFILL][BLOCK_N];
+#else
    __local uint  l_k_packed[BLOCK_N][DK_Q4_BLOCKS_PREFILL * 8];
    __local float l_k_scale [BLOCK_N][DK_Q4_BLOCKS_PREFILL];
+#endif
 #else
    __local half4 l_k[BLOCK_N][DK_VEC];
 #endif
@@ -1660,17 +1677,17 @@ __kernel void flash_attn_f32_q4_0(
                    const global char * blk_ptr = k_base + k_row_off + blk * Q4_0_BLOCK_SIZE;
                    const float df = (float) vload_half(0, (const global half *) blk_ptr);
                    const global uchar * qs = (const global uchar *)(blk_ptr + 2);
-                    l_k_scale[row][blk] = df;
+                    FA_K_SCALE(row, blk) = df;
                    uint k_packed[8];
                    pack_q4_0_nibbles(qs, k_packed);
                    #pragma unroll
                    for (int j = 0; j < 8; ++j) {
-                        l_k_packed[row][blk * 8 + j] = k_packed[j];
+                        FA_K_PACKED(row, blk * 8 + j) = k_packed[j];
                    }
                } else {
-                    l_k_scale[row][blk] = 0.0f;
+                    FA_K_SCALE(row, blk) = 0.0f;
                    #pragma unroll
-                    for (int j = 0; j < 8; ++j) l_k_packed[row][blk * 8 + j] = 0u;
+                    for (int j = 0; j < 8; ++j) FA_K_PACKED(row, blk * 8 + j) = 0u;
                }
            }
 #else
@@ -1760,6 +1777,19 @@ __kernel void flash_attn_f32_q4_0(
                for (int b_local = 0; b_local < SPLIT_DK_Q4_BLOCKS; ++b_local) {
                    const int b = k_blk_base + b_local;
                    int sum0 = 0, sum1 = 0, sum2 = 0, sum3 = 0;
+#ifdef FA_K_LDS_T
+                    // 4 KV rows are adjacent in the transposed tile: one 128-bit local
+                    // read per (block, group) instead of four 32-bit ones.
+                    #pragma unroll
+                    for (int g = 0; g < 8; ++g) {
+                        const uint  qp  = q_packed_pf[b_local * 8 + g];
+                        const uint4 kq4 = vload4(0, &l_k_packed[b * 8 + g][j]);
+                        sum0 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s0, sum0);
+                        sum1 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s1, sum1);
+                        sum2 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s2, sum2);
+                        sum3 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s3, sum3);
+                    }
+#else
                    #pragma unroll
                    for (int g = 0; g < 8; ++g) {
                        const uint qp = q_packed_pf[b_local * 8 + g];
@@ -1768,12 +1798,21 @@ __kernel void flash_attn_f32_q4_0(
                        sum2 = dot_acc_sat_4x8packed_ss_int(qp, l_k_packed[j+2][b * 8 + g], sum2);
                        sum3 = dot_acc_sat_4x8packed_ss_int(qp, l_k_packed[j+3][b * 8 + g], sum3);
                    }
+#endif
                    const float qd    = q_d_pf[b_local];
                    const int   q_sum = q_sum_pf[b_local];
+#ifdef FA_K_LDS_T
+                    const float4 ks4 = vload4(0, &l_k_scale[b][j]);
+                    s0 += (float)(sum0 - 8 * q_sum) * qd * ks4.s0;
+                    s1 += (float)(sum1 - 8 * q_sum) * qd * ks4.s1;
+                    s2 += (float)(sum2 - 8 * q_sum) * qd * ks4.s2;
+                    s3 += (float)(sum3 - 8 * q_sum) * qd * ks4.s3;
+#else
                    s0 += (float)(sum0 - 8 * q_sum) * qd * l_k_scale[j  ][b];
                    s1 += (float)(sum1 - 8 * q_sum) * qd * l_k_scale[j+1][b];
                    s2 += (float)(sum2 - 8 * q_sum) * qd * l_k_scale[j+2][b];
                    s3 += (float)(sum3 - 8 * q_sum) * qd * l_k_scale[j+3][b];
+#endif
                }
 #else
                ACC_TYPE4 dot_acc0 = (ACC_TYPE4)(0.0f);
@@ -1393,8 +1393,31 @@ __kernel void flash_attn_f32_q8_0(
    float slope = get_alibi_slope(max_bias, head_idx, n_head_log2, m0, m1);

 #ifdef FA_HAVE_INT_DOT
+// Accessors so the staging code is layout-agnostic.
+#ifdef FA_K_LDS_T
+#define FA_K_PACKED(ROW, IDX) l_k_packed[IDX][ROW]
+#define FA_K_SCALE(ROW, BLK)  l_k_scale[BLK][ROW]
+#else
+#define FA_K_PACKED(ROW, IDX) l_k_packed[ROW][IDX]
+#define FA_K_SCALE(ROW, BLK)  l_k_scale[ROW][BLK]
+#endif
+
+#ifdef FA_K_LDS_T
+    // K tile transposed: [block*8 + g][kv row] instead of [kv row][block*8 + g].
+    //
+    // The QK loop walks 4 KV rows at a time against the same (b, g), so in the original
+    // layout those 4 values are BLOCK_N*8 uints apart and cost 4 separate 32-bit local
+    // reads. Transposed they are adjacent, so they are one 128-bit read -- 4x fewer LDS
+    // issues for the same bytes and no extra registers. That matters because the QK loop
+    // is LDS-read-issue-bound: a wrong-math probe that kept every dp4a but cut the LDS
+    // reads ran the whole kernel 41% faster (18.51 -> 10.91 ms/op), and deleting QK
+    // outright only reached 10.88 -- i.e. essentially ALL of QK's cost is these reads.
+    __local uint  l_k_packed[DK_Q8_BLOCKS_PREFILL * 8][BLOCK_N];
+    __local float l_k_scale [DK_Q8_BLOCKS_PREFILL][BLOCK_N];
+#else
    __local uint  l_k_packed[BLOCK_N][DK_Q8_BLOCKS_PREFILL * 8];
    __local float l_k_scale [BLOCK_N][DK_Q8_BLOCKS_PREFILL];
+#endif
 #else
    __local half4 l_k[BLOCK_N][DK_VEC];
 #endif
@@ -1427,7 +1450,7 @@ __kernel void flash_attn_f32_q8_0(
                    const global char * blk_ptr = k_base + k_row_off + blk * Q8_0_BLOCK_SIZE;
                    const float df = (float) vload_half(0, (const global half *) blk_ptr);
                    const global uchar * qs = (const global uchar *)(blk_ptr + 2);
-                    l_k_scale[row][blk] = df;
+                    FA_K_SCALE(row, blk) = df;
                    #pragma unroll
                    for (int j = 0; j < 8; ++j) {
                        uint k_packed =
@@ -1435,12 +1458,12 @@ __kernel void flash_attn_f32_q8_0(
                             ((uint) qs[j*4 + 1]) <<  8 |
                             ((uint) qs[j*4 + 2]) << 16 |
                             ((uint) qs[j*4 + 3]) << 24;
-                        l_k_packed[row][blk * 8 + j] = k_packed;
+                        FA_K_PACKED(row, blk * 8 + j) = k_packed;
                    }
                } else {
-                    l_k_scale[row][blk] = 0.0f;
+                    FA_K_SCALE(row, blk) = 0.0f;
                    #pragma unroll
-                    for (int j = 0; j < 8; ++j) l_k_packed[row][blk * 8 + j] = 0u;
+                    for (int j = 0; j < 8; ++j) FA_K_PACKED(row, blk * 8 + j) = 0u;
                }
            }
 #else
@@ -1556,6 +1579,19 @@ __kernel void flash_attn_f32_q8_0(
                for (int b_local = 0; b_local < SPLIT_DK_Q8_BLOCKS; ++b_local) {
                    const int b = k_blk_base + b_local;
                    int sum0 = 0, sum1 = 0, sum2 = 0, sum3 = 0;
+#if defined(FA_K_LDS_T)
+                    // The 4 KV rows are adjacent in the transposed tile, so each (b, g)
+                    // step is ONE 128-bit local read instead of four 32-bit ones.
+                    #pragma unroll
+                    for (int g = 0; g < 8; ++g) {
+                        const uint qp = q_packed_pf[b_local * 8 + g];
+                        const uint4 kq4 = vload4(0, &l_k_packed[b * 8 + g][j]);
+                        sum0 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s0, sum0);
+                        sum1 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s1, sum1);
+                        sum2 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s2, sum2);
+                        sum3 = dot_acc_sat_4x8packed_ss_int(qp, kq4.s3, sum3);
+                    }
+#else
                    #pragma unroll
                    for (int g = 0; g < 8; ++g) {
                        const uint qp = q_packed_pf[b_local * 8 + g];
@@ -1564,11 +1600,20 @@ __kernel void flash_attn_f32_q8_0(
                        sum2 = dot_acc_sat_4x8packed_ss_int(qp, l_k_packed[j+2][b * 8 + g], sum2);
                        sum3 = dot_acc_sat_4x8packed_ss_int(qp, l_k_packed[j+3][b * 8 + g], sum3);
                    }
+#endif
                    const float qd = q_d_pf[b_local];
+#ifdef FA_K_LDS_T
+                    const float4 ks4 = vload4(0, &l_k_scale[b][j]);
+                    s0 += (float)sum0 * qd * ks4.s0;
+                    s1 += (float)sum1 * qd * ks4.s1;
+                    s2 += (float)sum2 * qd * ks4.s2;
+                    s3 += (float)sum3 * qd * ks4.s3;
+#else
                    s0 += (float)sum0 * qd * l_k_scale[j  ][b];
                    s1 += (float)sum1 * qd * l_k_scale[j+1][b];
                    s2 += (float)sum2 * qd * l_k_scale[j+2][b];
                    s3 += (float)sum3 * qd * l_k_scale[j+3][b];
+#endif
                }
 #else
                ACC_TYPE4 dot_acc0 = (ACC_TYPE4)(0.0f);
@@ -763,6 +763,7 @@ static void ggml_backend_openvino_device_get_props(ggml_backend_dev_t dev, ggml_
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

@@ -1881,6 +1881,7 @@ static void ggml_backend_rpc_device_get_props(ggml_backend_dev_t dev, struct ggm
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

@@ -5649,6 +5649,7 @@ static void ggml_backend_sycl_device_get_props(ggml_backend_dev_t dev, ggml_back
        /* .host_buffer           = */ host_buffer,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ events,
+        /* .mmap_support          = */ true,
    };
 }

@@ -111,6 +111,7 @@ uint32_t backend_device_get_props(apir_encoder * enc, apir_decoder * dec, virgl_
    apir_encode_bool_t(enc, &props.caps.host_buffer);
    apir_encode_bool_t(enc, &props.caps.buffer_from_host_ptr);
    apir_encode_bool_t(enc, &props.caps.events);
+    apir_encode_bool_t(enc, &props.caps.mmap_support);

    return 0;
 }
@@ -7,7 +7,7 @@
 #include <cstdint>

 #define APIR_PROTOCOL_MAJOR 0
-#define APIR_PROTOCOL_MINOR 1
+#define APIR_PROTOCOL_MINOR 2

 #define APIR_HANDSHAKE_MAGIC 0xab1e

@@ -11,9 +11,9 @@ static ggml_backend_buffer_t ggml_backend_remoting_buffer_type_alloc_buffer(ggml

    context->gpu = gpu;

-    bool async__unused, host_buffer__unused, events__unused;
+    bool async__unused, host_buffer__unused, events__unused, mmap_support__unused;
    bool buffer_from_host_ptr;
-    apir_device_get_props(gpu, &async__unused, &host_buffer__unused, &buffer_from_host_ptr, &events__unused);
+    apir_device_get_props(gpu, &async__unused, &host_buffer__unused, &buffer_from_host_ptr, &events__unused, &mmap_support__unused);

    if (buffer_from_host_ptr) {
        context->apir_context = apir_device_buffer_from_ptr(gpu, size, size);
@@ -65,7 +65,7 @@ static void ggml_backend_remoting_device_get_props(ggml_backend_dev_t dev, ggml_

    virtgpu * gpu = DEV_TO_GPU(dev);
    apir_device_get_props(gpu, &props->caps.async, &props->caps.host_buffer, &props->caps.buffer_from_host_ptr,
-                          &props->caps.events);
+                          &props->caps.events, &props->caps.mmap_support);

    props->caps.buffer_from_host_ptr = false;
    props->caps.async                = false;
@@ -144,7 +144,8 @@ void apir_device_get_props(virtgpu * gpu,
                           bool *    async,
                           bool *    host_buffer,
                           bool *    buffer_from_host_ptr,
-                           bool *    events) {
+                           bool *    events,
+                           bool *    mmap_support) {
    apir_encoder *        encoder;
    apir_decoder *        decoder;
    ApirForwardReturnCode ret;
@@ -157,6 +158,7 @@ void apir_device_get_props(virtgpu * gpu,
    apir_decode_bool_t(decoder, host_buffer);
    apir_decode_bool_t(decoder, buffer_from_host_ptr);
    apir_decode_bool_t(decoder, events);
+    apir_decode_bool_t(decoder, mmap_support);

    remote_call_finish(gpu, encoder, decoder);

@@ -13,7 +13,8 @@ void                           apir_device_get_props(struct virtgpu * gpu,
                                                     bool *           async,
                                                     bool *           host_buffer,
                                                     bool *           buffer_from_host_ptr,
-                                                     bool *           events);
+                                                     bool *           events,
+                                                     bool *           mmap_support);
 apir_buffer_context_t          apir_device_buffer_from_ptr(struct virtgpu * gpu, size_t size, size_t max_tensor_size);

 /* buffer-type */
@@ -4627,6 +4627,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q5_1], matmul_q5_1_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q8_0], matmul_q8_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q2_K], matmul_q2_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
+        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_TQ2_0], matmul_tq2_0_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q3_K], matmul_q3_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_K], matmul_q4_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q5_K], matmul_q5_k_f16, mmq_wg_denoms_k, warptile_mmq_k, vk_mat_mat_push_constants, 3)
@@ -4667,6 +4668,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
@@ -4739,6 +4741,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0], matmul_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );

        CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_K], matmul_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_TQ2_0], matmul_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
        CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q3_K], matmul_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
        CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_K], matmul_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
        CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_K], matmul_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
@@ -4783,6 +4786,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
        CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
        CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
        CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
        CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
        CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
@@ -4873,6 +4877,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_1], matmul_q5_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
        CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0], matmul_q8_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
        CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_K], matmul_q2_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+        CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_TQ2_0], matmul_tq2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
        CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q3_K], matmul_q3_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
        CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_K], matmul_q4_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
        CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q5_K], matmul_q5_k_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
@@ -4921,6 +4926,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
            CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_subgroup_q5_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
            CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_subgroup_q8_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
            CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_subgroup_q2_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+            CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_subgroup_tq2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
            CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_subgroup_q3_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
            CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_subgroup_q4_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
            CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_subgroup_q5_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
@@ -4968,6 +4974,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
            CREATE_MM2(GGML_TYPE_Q5_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_1], matmul_id_q5_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
            CREATE_MM2(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q8_0], matmul_id_q8_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
            CREATE_MM2(GGML_TYPE_Q2_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_K], matmul_id_q2_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+            CREATE_MM2(GGML_TYPE_TQ2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_TQ2_0], matmul_id_tq2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
            CREATE_MM2(GGML_TYPE_Q3_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q3_K], matmul_id_q3_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
            CREATE_MM2(GGML_TYPE_Q4_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_K], matmul_id_q4_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
            CREATE_MM2(GGML_TYPE_Q5_K, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q5_K], matmul_id_q5_k_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
@@ -5047,6 +5054,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        CREATE_MM(GGML_TYPE_Q8_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q8_0].f32acc, matmul_q8_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);

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

@@ -18013,6 +18035,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                    case GGML_TYPE_IQ4_NL:
                    case GGML_TYPE_MXFP4:
                    case GGML_TYPE_NVFP4:
+                    case GGML_TYPE_TQ2_0:
                        break;
                    default:
                        return false;
@@ -18118,6 +18141,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                    case GGML_TYPE_IQ4_NL:
                    case GGML_TYPE_MXFP4:
                    case GGML_TYPE_NVFP4:
+                    case GGML_TYPE_TQ2_0:
                    case GGML_TYPE_I32:
                        return true;
                    default:
@@ -608,6 +608,20 @@ vec2 get_dm(uint ib, uint a_offset) {
 }
 #endif

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

+layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0 {
+   block_tq2_0 block;
+};
+
+layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0_packed16 {
+   block_tq2_0_packed16 block;
+};
+
+float16_t dequantFuncTQ2_0(const in decodeBufTQ2_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
+{
+    decodeBufTQ2_0_packed16 bl16 = decodeBufTQ2_0_packed16(bl);
+    const uint idx = coordInBlock[1];
+
+    const uint qsshift = (idx & 0x60) >> 4;             // 0,2,4,6
+
+    uint qs = uint32_t(bl16.block.qs[((idx & 0x80) >> 3) + ((idx & 0x1E) >> 1)]);
+    qs = (qs >> qsshift) & 0x0303;
+    qs = unpack8(qs)[idx & 1];
+
+    return bl.block.d * (float16_t(int(qs)) - float16_t(1.0));
+}
+
+f16vec4 dequantFuncTQ2_0_v(const in decodeBufTQ2_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
+{
+    const uint idx = coordInBlock[1];
+
+    const uint qsshift = (idx & 0x60) >> 4;             // 0,2,4,6
+    const uint qsi = ((idx & 0x80) >> 2) + (idx & 0x1C); // byte index of 4-aligned group
+
+    const uint qsw = (uint(bl.block.qs[qsi]))
+                   | (uint(bl.block.qs[qsi + 1]) << 8)
+                   | (uint(bl.block.qs[qsi + 2]) << 16)
+                   | (uint(bl.block.qs[qsi + 3]) << 24);
+    const u8vec4 q = unpack8((qsw >> qsshift) & 0x03030303);
+
+    return bl.block.d * (f16vec4(q) - f16vec4(1.0));
+}
+
 layout(buffer_reference, std430, buffer_reference_align = 4) buffer decodeBufQ2_K {
   block_q2_K block;
 };
@@ -1368,6 +1406,9 @@ f16vec4 dequantFuncNVFP4_v(const in decodeBufNVFP4 bl, const in uint blockCoords
 #elif defined(DATA_A_Q8_0)
 #define dequantFuncA dequantFuncQ8_0
 #define dequantFuncA_v dequantFuncQ8_0_v
+#elif defined(DATA_A_TQ2_0)
+#define dequantFuncA dequantFuncTQ2_0
+#define dequantFuncA_v dequantFuncTQ2_0_v
 #elif defined(DATA_A_Q2_K)
 #define dequantFuncA dequantFuncQ2_K
 #define dequantFuncA_v dequantFuncQ2_K_v
@@ -0,0 +1,31 @@
+#version 450
+
+#include "dequant_head.glsl"
+
+layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
+
+layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
+layout (binding = 1) writeonly buffer D {D_TYPE data_b[];};
+
+void main() {
+    [[unroll]] for (uint wgy = 0; wgy < 256; wgy++) {
+        const uint i = gl_WorkGroupID.x * 256 + wgy;
+        if (i >= p.nel / QUANT_K) {
+            return;
+        }
+
+        const uint tid = gl_LocalInvocationID.x;
+        const uint ip = tid / 32;         // group 0,1 (128 elems each)
+        const uint il = tid - 32 * ip;    // byte in group 0..31
+
+        const uint y_idx = i * QUANT_K + 128 * ip + il;
+
+        const uint8_t qs = data_a[i].qs[32 * ip + il];
+
+        const FLOAT_TYPE d = FLOAT_TYPE(data_a[i].d);
+        data_b[y_idx +  0] = D_TYPE(d * FLOAT_TYPE(int((qs >> 0) & 3) - 1));
+        data_b[y_idx + 32] = D_TYPE(d * FLOAT_TYPE(int((qs >> 2) & 3) - 1));
+        data_b[y_idx + 64] = D_TYPE(d * FLOAT_TYPE(int((qs >> 4) & 3) - 1));
+        data_b[y_idx + 96] = D_TYPE(d * FLOAT_TYPE(int((qs >> 6) & 3) - 1));
+    }
+}
@@ -0,0 +1,102 @@
+#version 450
+#extension GL_EXT_shader_explicit_arithmetic_types_int32 : require
+
+#include "mul_mat_vec_base.glsl"
+
+layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
+
+FLOAT_TYPE temp[NUM_COLS][NUM_ROWS];
+
+// ternary TQ2_0: w = (q - 1) * d. Same qs group/level layout as q2_K, but a
+// single f16 scale per 256-block and no mins:
+//   sum_e b_e * (q_e - 1) * d  =  d * (sum_e b_e * q_e  -  sum_e b_e)
+void calc_superblock(const uint a_offset, const uint b_offset, const uint v_im, const uint q_offset, const uint y_offset, const uint i, const uint num_blocks_per_row, const uint first_row, const uint num_rows) {
+    const uint y_idx = i * QUANT_K + y_offset;
+
+    [[unroll]] for (uint n = 0; n < num_rows; ++n) {
+        const uint ib0 = a_offset + (first_row+n)*num_blocks_per_row;
+        if (i >= num_blocks_per_row) {
+            continue;
+        }
+
+        const uint32_t qs_u32 = uint32_t(data_a_packed16[ib0 + i].qs[q_offset / 2]) | (uint32_t(data_a_packed16[ib0 + i].qs[q_offset / 2 + 8]) << 16);
+        const vec4 qs_u32_0 = vec4(unpack8(qs_u32 & 0x03030303));
+        const vec4 qs_u32_2 = vec4(unpack8((qs_u32 >> 2) & 0x03030303));
+        const vec4 qs_u32_4 = vec4(unpack8((qs_u32 >> 4) & 0x03030303));
+        const vec4 qs_u32_6 = vec4(unpack8((qs_u32 >> 6) & 0x03030303));
+
+        const FLOAT_TYPE d = FLOAT_TYPE(data_a[ib0 + i].d);
+
+        [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
+            vec2 b0 =   vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 +  0]);
+            vec2 b16 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 +  8]);
+            vec2 b32 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 16]);
+            vec2 b48 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 24]);
+            vec2 b64 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 32]);
+            vec2 b80 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 40]);
+            vec2 b96 =  vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 48]);
+            vec2 b112 = vec2(data_b_v2[(j*p.batch_stride_b + b_offset + y_idx) / 2 + 56]);
+
+            FLOAT_TYPE sumq = FLOAT_TYPE(0.0);
+            FLOAT_TYPE sumb = FLOAT_TYPE(0.0);
+            [[unroll]] for (int l = 0; l < 2; ++l) {
+                sumq = fma(FLOAT_TYPE(b0[l]),   FLOAT_TYPE(qs_u32_0[l  ]),
+                       fma(FLOAT_TYPE(b16[l]),  FLOAT_TYPE(qs_u32_0[l+2]),
+                       fma(FLOAT_TYPE(b32[l]),  FLOAT_TYPE(qs_u32_2[l  ]),
+                       fma(FLOAT_TYPE(b48[l]),  FLOAT_TYPE(qs_u32_2[l+2]),
+                       fma(FLOAT_TYPE(b64[l]),  FLOAT_TYPE(qs_u32_4[l  ]),
+                       fma(FLOAT_TYPE(b80[l]),  FLOAT_TYPE(qs_u32_4[l+2]),
+                       fma(FLOAT_TYPE(b96[l]),  FLOAT_TYPE(qs_u32_6[l  ]),
+                       fma(FLOAT_TYPE(b112[l]), FLOAT_TYPE(qs_u32_6[l+2]), sumq))))))));
+                sumb += FLOAT_TYPE(b0[l]) + FLOAT_TYPE(b16[l]) + FLOAT_TYPE(b32[l]) + FLOAT_TYPE(b48[l])
+                      + FLOAT_TYPE(b64[l]) + FLOAT_TYPE(b80[l]) + FLOAT_TYPE(b96[l]) + FLOAT_TYPE(b112[l]);
+            }
+            temp[j][n] = fma(d, sumq - sumb, temp[j][n]);
+        }
+    }
+}
+
+void compute_outputs(const uint32_t first_row, const uint32_t num_rows) {
+    uint a_offset, b_offset, d_offset;
+    get_offsets(a_offset, b_offset, d_offset);
+
+    const uint num_blocks_per_row = p.ncols / QUANT_K;
+
+    // 16 threads are used to process each block
+    const uint it_size = gl_WorkGroupSize.x/16;
+    const uint tid = gl_LocalInvocationID.x;
+    const uint itid = tid%16;  // 0...15
+    const uint ix = tid/16;
+
+    const uint v_im = itid/8;                                // 0 or 1. 0 computes 0..., 1 computes 128...
+    const uint v_in = itid - 8*v_im;                         // 0...7
+
+    const uint l0 = 2*v_in;                                  // 0...15
+    const uint q_offset = 32*v_im + l0;
+    const uint y_offset = 128*v_im + l0;
+
+    [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
+        [[unroll]] for (uint i = 0; i < NUM_ROWS; ++i) {
+            temp[j][i] = FLOAT_TYPE(0);
+        }
+    }
+
+    for (uint i0 = 0; i0 < num_blocks_per_row; i0 += it_size)
+        calc_superblock(a_offset, b_offset, v_im, q_offset, y_offset, i0 + ix, num_blocks_per_row, first_row, num_rows);
+
+    reduce_result(temp, d_offset, first_row, num_rows, tid);
+}
+
+void main() {
+    const uint first_row = NUM_ROWS * (gl_WorkGroupID.x + gl_NumWorkGroups.x * gl_WorkGroupID.z);
+
+    // do NUM_ROWS at a time, unless there aren't enough remaining rows
+    if (first_row + NUM_ROWS <= p.stride_d) {
+        compute_outputs(first_row, NUM_ROWS);
+    } else {
+        if (first_row >= p.stride_d) {
+            return;
+        }
+        compute_outputs(first_row, p.stride_d - first_row);
+    }
+}
@@ -182,6 +182,22 @@ void load_a_to_shmem(const uint pos_a, const uint row, const uint col, const uin

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

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

 struct block_q3_K
@@ -72,6 +72,7 @@ const std::vector<std::string> type_names = {
    "iq4_nl",
    "mxfp4",
    "nvfp4",
+    "tq2_0",
    "bf16",
 };

@@ -733,7 +734,7 @@ void process_shaders() {
    for (const auto& tname : type_names) {
        // mul mat vec
        std::string data_a_key = "DATA_A_" + to_uppercase(tname);
-        std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_")) ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";
+        std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_") || tname == "tq2_0") ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";

        string_to_spv("mul_mat_vec_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}));
        string_to_spv("mul_mat_vec_" + tname + "_f16_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float16_t"}, {"B_TYPEV2", "f16vec2"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}));
@@ -2815,11 +2815,25 @@ class ggml_webgpu_shader_lib {
        key.common.v_direct &= decisions.use_sg_matrix && key.common.v_type == GGML_TYPE_F16;
        key.use_sg_matrix = decisions.use_sg_matrix;

-        const uint32_t max_kv_tile = ggml_webgpu_flash_attn_max_kv_tile(
+        uint32_t max_kv_tile = ggml_webgpu_flash_attn_max_kv_tile(
            context.wg_mem_limit_bytes, decisions.q_tile, decisions.use_sg_matrix ? context.sg_mat_n : 1u,
            key.common.head_dim_qk, key.common.head_dim_v, key.common.has_mask,
            key.common.k_direct || key.common.v_direct);
-        GGML_ASSERT(max_kv_tile > 0);
+
+        // WorkGroup storage size isn't enough for some params with subgroup matrices path (ref. https://github.com/ggml-org/llama.cpp/pull/26566)
+        if (max_kv_tile == 0) {
+            GGML_ASSERT(decisions.use_sg_matrix);
+            // switch to flash_attn_reg_tile path
+            decisions.use_sg_matrix = false;
+            decisions.q_tile        = GGML_WEBGPU_FLASH_ATTN_TILE_Q_TILE;
+            key.common.k_direct     = false;
+            key.common.v_direct     = false;
+            key.use_sg_matrix       = false;
+            max_kv_tile             = ggml_webgpu_flash_attn_max_kv_tile(
+                context.wg_mem_limit_bytes, decisions.q_tile, 1u, key.common.head_dim_qk, key.common.head_dim_v,
+                key.common.has_mask, key.common.k_direct || key.common.v_direct);
+            GGML_ASSERT(max_kv_tile > 0);
+        }

        decisions.kv_tile = decisions.use_sg_matrix ?
                                std::min(max_kv_tile, context.sg_mat_n * GGML_WEBGPU_FLASH_ATTN_PREFERRED_KV_SG_TILES) :
@@ -2993,6 +3007,10 @@ class ggml_webgpu_shader_lib {
                defines.push_back("SRC_F16");
                variant += "_f16";
                break;
+            case GGML_TYPE_I32:
+                defines.push_back("SRC_I32");
+                variant += "_i32";
+                break;
            default:
                GGML_ABORT("Unsupported src type for cpy shader");
        }
@@ -3942,6 +3942,7 @@ static void ggml_backend_webgpu_device_get_props(ggml_backend_dev_t dev, struct
        /* .host_buffer           = */ false,
        /* .buffer_from_host_ptr  = */ false,
        /* .events                = */ false,
+        /* .mmap_support          = */ true,
    };
 }

@@ -4283,9 +4284,8 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
            break;
        case GGML_OP_CPY:
        case GGML_OP_CONT:
-            supports_op = ((op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) &&
-                           (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16)) ||
-                          (op->type == GGML_TYPE_I32 && src0->type == GGML_TYPE_F32);
+            supports_op = (op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_I32) &&
+                          (src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_I32);
            break;
        case GGML_OP_SET:
            supports_op = src0->type == src1->type && src0->type == op->type &&
@@ -4,6 +4,8 @@ enable f16;
 #define SRC_TYPE f32
 #elif defined(SRC_F16)
 #define SRC_TYPE f16
+#elif defined(SRC_I32)
+#define SRC_TYPE i32
 #endif

 #ifdef DST_F32
@@ -487,7 +487,8 @@ static void ggml_backend_zdnn_device_get_props(ggml_backend_dev_t dev, ggml_back
        /* .async                = */ false,
        /* .host_buffer          = */ false,
        /* .buffer_from_host_ptr = */ false,
-        /* .events               = */ false
+        /* .events               = */ false,
+        /* .mmap_support         = */ true,
    };
 }

@@ -654,7 +654,8 @@ static void ggml_backend_zendnn_device_get_props(ggml_backend_dev_t dev, struct
        /* .async                = */ false,
        /* .host_buffer          = */ false,
        /* .buffer_from_host_ptr = */ true,
-        /* .events               = */ false
+        /* .events               = */ false,
+        /* .mmap_support         = */ true,
    };
 }

@@ -611,6 +611,13 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr
        GGML_ASSERT(int64_t(ctx->kv.size()) == n_kv);

        const int alignment_idx = gguf_find_key(ctx, GGUF_KEY_GENERAL_ALIGNMENT);
+        if (alignment_idx != -1 && gguf_get_kv_type(ctx, alignment_idx) != GGUF_TYPE_UINT32) {
+            GGML_LOG_ERROR("%s: key '%s' must be of type %s but is %s\n",
+                __func__, GGUF_KEY_GENERAL_ALIGNMENT, gguf_type_name(GGUF_TYPE_UINT32),
+                gguf_type_name(gguf_get_kv_type(ctx, alignment_idx)));
+            gguf_free(ctx);
+            return nullptr;
+        }
        ctx->alignment = alignment_idx == -1 ? GGUF_DEFAULT_ALIGNMENT : gguf_get_val_u32(ctx, alignment_idx);

        if (ctx->alignment == 0 || (ctx->alignment & (ctx->alignment - 1)) != 0) {
@@ -682,9 +689,11 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr
            }

            // check that the total number of elements is representable
-            if (ok && ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) ||
-                       (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) ||
-                       (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) {
+            // (a zero-element tensor is trivially representable; the guard also avoids a division by zero below)
+            if (ok && ggml_nelements(&info.t) > 0 &&
+                ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) ||
+                 (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) ||
+                 (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) {

                GGML_LOG_ERROR("%s: total number of elements in tensor '%s' with shape "
                    "(%" PRIi64 ", %" PRIi64 ", %" PRIi64 ", %" PRIi64 ") is >= %" PRIi64 "\n",
@@ -407,6 +407,8 @@ class Keys:

    class ClipGenAudio:
        PROJECTOR_TYPE      = "clip.gen.audio.projector_type" # for mixed modality models
+        # name of the weight variant, for settings that are not in the checkpoint
+        MODEL_VARIANT       = "clip.gen.audio.model_variant"
        EMBEDDING_LENGTH    = "clip.gen.audio.embedding_length"
        FEED_FORWARD_LENGTH = "clip.gen.audio.feed_forward_length"
        BLOCK_COUNT         = "clip.gen.audio.block_count"
@@ -581,6 +583,7 @@ class MODEL_ARCH(IntEnum):
    MELLUM           = auto()
    NANBEIGE         = auto()
    QWEN3TTS         = auto()
+    POCKETTTS        = auto()


 class VISION_PROJECTOR_TYPE(IntEnum):
@@ -1040,6 +1043,38 @@ class MODEL_TENSOR(IntEnum):
    A_GEN_WAV_DAC_RES_CONV2        = auto() # DAC residual unit, pointwise causal conv
    A_GEN_WAV_DAC_POST_SNAKE       = auto() # DAC final SnakeBeta
    A_GEN_WAV_DAC_POST_CONV        = auto() # DAC conv_post -> 1-channel PCM
+    # pocket-tts: SEANet encoder (speaker path) and decoder (a.gen.wav path)
+    A_ENC_SEANET_CONV_IN     = auto()
+    A_ENC_SEANET_CONV_OUT    = auto()
+    A_ENC_SEANET_RES_CONV1   = auto() # residual unit, dilated conv
+    A_ENC_SEANET_RES_CONV2   = auto() # residual unit, pointwise conv
+    A_ENC_SEANET_SCALE_CONV  = auto() # strided downsample conv
+    A_ENC_ATTN_SCALE         = auto() # layer scale (gamma) on the attn output
+    A_ENC_FFN_SCALE_LS       = auto() # layer scale (gamma) on the FFN output
+    A_ENC_SPEAKER_PROJ       = auto() # voice latent -> backbone embd
+    A_GEN_FLOW_INPUT_PROJ    = auto()
+    A_GEN_FLOW_COND_EMBD     = auto()
+    A_GEN_FLOW_TIME_FREQS    = auto() # timestep embedder, stored cos/sin frequencies
+    A_GEN_FLOW_TIME_UP       = auto()
+    A_GEN_FLOW_TIME_DOWN     = auto()
+    A_GEN_FLOW_TIME_NORM     = auto() # RMSNorm alpha
+    A_GEN_FLOW_BLK_NORM      = auto() # AdaLN res block, in_ln
+    A_GEN_FLOW_BLK_UP        = auto()
+    A_GEN_FLOW_BLK_DOWN      = auto()
+    A_GEN_FLOW_BLK_ADA       = auto() # AdaLN modulation, -> shift/scale/gate
+    A_GEN_FLOW_FINAL_ADA     = auto() # final layer AdaLN modulation, -> shift/scale
+    A_GEN_FLOW_FINAL_PROJ    = auto()
+    A_GEN_OUT_EOS            = auto() # end-of-speech head on the backbone hidden state
+    A_GEN_INPUT_LINEAR       = auto() # generated latent -> backbone embd
+    A_GEN_EMB_MEAN           = auto() # latent denormalization stats
+    A_GEN_EMB_STD            = auto()
+    A_GEN_WAV_QUANT_OUT      = auto() # DummyQuantizer output_proj, latent -> decoder dim
+    A_GEN_WAV_UPSAMPLE       = auto() # frame rate -> encoder frame rate, depthwise convtr
+    A_GEN_WAV_SEANET_CONV_IN   = auto()
+    A_GEN_WAV_SEANET_CONV_OUT  = auto() # -> 1-channel PCM
+    A_GEN_WAV_SEANET_RES_CONV1 = auto()
+    A_GEN_WAV_SEANET_RES_CONV2 = auto()
+    A_GEN_WAV_SEANET_SCALE_CONV = auto() # strided upsample convtr
    A_MMPROJ              = auto()
    A_MMPROJ_FC           = auto()
    A_MM_NORM_PRE         = auto()
@@ -1255,6 +1290,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
    MODEL_ARCH.MELLUM:           "mellum",
    MODEL_ARCH.NANBEIGE:         "nanbeige",
    MODEL_ARCH.QWEN3TTS:         "qwen3tts",
+    MODEL_ARCH.POCKETTTS:        "pockettts",
 }

 VISION_PROJECTOR_TYPE_NAMES: dict[VISION_PROJECTOR_TYPE, str] = {
@@ -1709,6 +1745,37 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
    MODEL_TENSOR.A_GEN_WAV_DAC_RES_CONV2:   "a.gen.wav.dac.blk.{bid}.res.{xid}.conv2",
    MODEL_TENSOR.A_GEN_WAV_DAC_POST_SNAKE:  "a.gen.wav.dac.post_snake",
    MODEL_TENSOR.A_GEN_WAV_DAC_POST_CONV:   "a.gen.wav.dac.post_conv",
+    MODEL_TENSOR.A_ENC_SEANET_CONV_IN:      "a.seanet.conv_in",
+    MODEL_TENSOR.A_ENC_SEANET_CONV_OUT:     "a.seanet.conv_out",
+    MODEL_TENSOR.A_ENC_SEANET_RES_CONV1:    "a.seanet.blk.{bid}.res_conv1",
+    MODEL_TENSOR.A_ENC_SEANET_RES_CONV2:    "a.seanet.blk.{bid}.res_conv2",
+    MODEL_TENSOR.A_ENC_SEANET_SCALE_CONV:   "a.seanet.blk.{bid}.scale_conv",
+    MODEL_TENSOR.A_ENC_ATTN_SCALE:          "a.blk.{bid}.ls1",
+    MODEL_TENSOR.A_ENC_FFN_SCALE_LS:        "a.blk.{bid}.ls2",
+    MODEL_TENSOR.A_ENC_SPEAKER_PROJ:        "a.speaker_proj",
+    MODEL_TENSOR.A_GEN_FLOW_INPUT_PROJ:     "a.gen.flow.input_proj",
+    MODEL_TENSOR.A_GEN_FLOW_COND_EMBD:      "a.gen.flow.cond_embd",
+    MODEL_TENSOR.A_GEN_FLOW_TIME_FREQS:     "a.gen.flow.time.{bid}.freqs",
+    MODEL_TENSOR.A_GEN_FLOW_TIME_UP:        "a.gen.flow.time.{bid}.up",
+    MODEL_TENSOR.A_GEN_FLOW_TIME_DOWN:      "a.gen.flow.time.{bid}.down",
+    MODEL_TENSOR.A_GEN_FLOW_TIME_NORM:      "a.gen.flow.time.{bid}.norm",
+    MODEL_TENSOR.A_GEN_FLOW_BLK_NORM:       "a.gen.flow.blk.{bid}.norm",
+    MODEL_TENSOR.A_GEN_FLOW_BLK_UP:         "a.gen.flow.blk.{bid}.up",
+    MODEL_TENSOR.A_GEN_FLOW_BLK_DOWN:       "a.gen.flow.blk.{bid}.down",
+    MODEL_TENSOR.A_GEN_FLOW_BLK_ADA:        "a.gen.flow.blk.{bid}.ada",
+    MODEL_TENSOR.A_GEN_FLOW_FINAL_ADA:      "a.gen.flow.final.ada",
+    MODEL_TENSOR.A_GEN_FLOW_FINAL_PROJ:     "a.gen.flow.final.proj",
+    MODEL_TENSOR.A_GEN_OUT_EOS:             "a.gen.out_eos",
+    MODEL_TENSOR.A_GEN_INPUT_LINEAR:        "a.gen.input_linear",
+    MODEL_TENSOR.A_GEN_EMB_MEAN:            "a.gen.emb_mean",
+    MODEL_TENSOR.A_GEN_EMB_STD:             "a.gen.emb_std",
+    MODEL_TENSOR.A_GEN_WAV_QUANT_OUT:       "a.gen.wav.quant_out",
+    MODEL_TENSOR.A_GEN_WAV_UPSAMPLE:        "a.gen.wav.upsample",
+    MODEL_TENSOR.A_GEN_WAV_SEANET_CONV_IN:  "a.gen.wav.seanet.conv_in",
+    MODEL_TENSOR.A_GEN_WAV_SEANET_CONV_OUT: "a.gen.wav.seanet.conv_out",
+    MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1: "a.gen.wav.seanet.blk.{bid}.res_conv1",
+    MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2: "a.gen.wav.seanet.blk.{bid}.res_conv2",
+    MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV: "a.gen.wav.seanet.blk.{bid}.scale_conv",
    MODEL_TENSOR.A_MMPROJ:                  "mm.a.mlp.{bid}",
    MODEL_TENSOR.A_MMPROJ_FC:               "mm.a.fc",
    MODEL_TENSOR.A_MM_NORM_PRE:             "mm.a.norm_pre",
@@ -2020,6 +2087,37 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.A_GEN_WAV_DAC_RES_CONV2,
        MODEL_TENSOR.A_GEN_WAV_DAC_POST_SNAKE,
        MODEL_TENSOR.A_GEN_WAV_DAC_POST_CONV,
+        MODEL_TENSOR.A_ENC_SEANET_CONV_IN,
+        MODEL_TENSOR.A_ENC_SEANET_CONV_OUT,
+        MODEL_TENSOR.A_ENC_SEANET_RES_CONV1,
+        MODEL_TENSOR.A_ENC_SEANET_RES_CONV2,
+        MODEL_TENSOR.A_ENC_SEANET_SCALE_CONV,
+        MODEL_TENSOR.A_ENC_ATTN_SCALE,
+        MODEL_TENSOR.A_ENC_FFN_SCALE_LS,
+        MODEL_TENSOR.A_ENC_SPEAKER_PROJ,
+        MODEL_TENSOR.A_GEN_FLOW_INPUT_PROJ,
+        MODEL_TENSOR.A_GEN_FLOW_COND_EMBD,
+        MODEL_TENSOR.A_GEN_FLOW_TIME_FREQS,
+        MODEL_TENSOR.A_GEN_FLOW_TIME_UP,
+        MODEL_TENSOR.A_GEN_FLOW_TIME_DOWN,
+        MODEL_TENSOR.A_GEN_FLOW_TIME_NORM,
+        MODEL_TENSOR.A_GEN_FLOW_BLK_NORM,
+        MODEL_TENSOR.A_GEN_FLOW_BLK_UP,
+        MODEL_TENSOR.A_GEN_FLOW_BLK_DOWN,
+        MODEL_TENSOR.A_GEN_FLOW_BLK_ADA,
+        MODEL_TENSOR.A_GEN_FLOW_FINAL_ADA,
+        MODEL_TENSOR.A_GEN_FLOW_FINAL_PROJ,
+        MODEL_TENSOR.A_GEN_OUT_EOS,
+        MODEL_TENSOR.A_GEN_INPUT_LINEAR,
+        MODEL_TENSOR.A_GEN_EMB_MEAN,
+        MODEL_TENSOR.A_GEN_EMB_STD,
+        MODEL_TENSOR.A_GEN_WAV_QUANT_OUT,
+        MODEL_TENSOR.A_GEN_WAV_UPSAMPLE,
+        MODEL_TENSOR.A_GEN_WAV_SEANET_CONV_IN,
+        MODEL_TENSOR.A_GEN_WAV_SEANET_CONV_OUT,
+        MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1,
+        MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2,
+        MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV,
        MODEL_TENSOR.A_ENC_CONV_NORM_MEAN,
        MODEL_TENSOR.A_ENC_CONV_NORM_VAR,
        MODEL_TENSOR.A_ENC_MEL_FILTERS,
@@ -4628,6 +4726,7 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.D2T,
    ],
    MODEL_ARCH.DFLASH: [
+        MODEL_TENSOR.TOKEN_EMBD,
        MODEL_TENSOR.OUTPUT_NORM,
        MODEL_TENSOR.ATTN_NORM,
        MODEL_TENSOR.ATTN_Q,
@@ -4903,6 +5002,18 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.FFN_DOWN,
        MODEL_TENSOR.FFN_UP,
    ],
+    MODEL_ARCH.POCKETTTS: [
+        MODEL_TENSOR.TOKEN_EMBD,
+        MODEL_TENSOR.OUTPUT_NORM,
+        MODEL_TENSOR.ATTN_NORM,
+        MODEL_TENSOR.ATTN_Q,
+        MODEL_TENSOR.ATTN_K,
+        MODEL_TENSOR.ATTN_V,
+        MODEL_TENSOR.ATTN_OUT,
+        MODEL_TENSOR.FFN_NORM,
+        MODEL_TENSOR.FFN_DOWN,
+        MODEL_TENSOR.FFN_UP,
+    ],
 }

 # tensors that will not be serialized
@@ -5179,6 +5290,8 @@ class VisionProjectorType:
    NEMOTRON_V2_VL = "nemotron_v2_vl"
    QWEN3TTS_SPKENC = "qwen3tts_spkenc" # audio: ECAPA-TDNN speaker encoder
    QWEN3TTS_GEN = "qwen3tts_gen" # audio generation: code_predictor
+    POCKETTTS_SPKENC = "pockettts_spkenc" # audio: mimi encoder as voice-prompt encoder
+    POCKETTTS_GEN = "pockettts_gen" # audio generation: flow-matching decoder + mimi decoder
    HUNYUANVL      = "hunyuanvl"
    PARAKEET       = "parakeet"  # audio
    MINIMAXM3      = "minimax_m3"
@@ -1453,6 +1453,9 @@ class GGUFWriter:
    def add_gen_audio_attention_layernorm_eps(self, value: float) -> None:
        self.add_float32(Keys.ClipGenAudio.Attention.LAYERNORM_EPS, value)

+    def add_gen_audio_model_variant(self, value: str) -> None:
+        self.add_string(Keys.ClipGenAudio.MODEL_VARIANT, value)
+
    def add_xielu_alpha_p(self, values: Sequence[float]):
        self.add_array(Keys.xIELU.ALPHA_P, values)

@@ -203,11 +203,12 @@ extern "C" {
    };

    enum llama_load_mode {
-        LLAMA_LOAD_MODE_NONE       = 0, // no special loading mode
-        LLAMA_LOAD_MODE_MMAP       = 1, // memory map the model
-        LLAMA_LOAD_MODE_MLOCK      = 2, // force system to keep model in RAM rather than swapping or compressing
-        LLAMA_LOAD_MODE_MMAP_MLOCK = 3, // mmap + force system to keep model in RAM rather than swapping or compressing
-        LLAMA_LOAD_MODE_DIRECT_IO  = 4, // use direct I/O if available
+        LLAMA_LOAD_MODE_AUTO       = -1, // auto-detect based on device capabilities
+        LLAMA_LOAD_MODE_NONE       =  0, // no special loading mode
+        LLAMA_LOAD_MODE_MMAP       =  1, // memory map the model
+        LLAMA_LOAD_MODE_MLOCK      =  2, // force system to keep model in RAM rather than swapping or compressing
+        LLAMA_LOAD_MODE_MMAP_MLOCK =  3, // mmap + force system to keep model in RAM rather than swapping or compressing
+        LLAMA_LOAD_MODE_DIRECT_IO  =  4, // use direct I/O if available
    };

    LLAMA_API const char * llama_load_mode_name(enum llama_load_mode load_mode);
@@ -348,14 +349,15 @@ extern "C" {
    // NOTE: changing the default values of parameters marked as [EXPERIMENTAL] may cause crashes or incorrect results in certain configurations
    //       https://github.com/ggml-org/llama.cpp/pull/7544
    struct llama_context_params {
-        uint32_t n_ctx;             // text context, 0 = from model
-        uint32_t n_batch;           // logical maximum batch size that can be submitted to llama_decode
-        uint32_t n_ubatch;          // physical maximum batch size
-        uint32_t n_seq_max;         // max number of sequences (i.e. distinct states for recurrent models)
-        uint32_t n_rs_seq;          // number of recurrent-state snapshots per seq for rollback (0 = no rollback) [EXPERIMENTAL]
-        uint32_t n_outputs_max;     // max outputs in a ubatch (0 = n_batch)
-        int32_t  n_threads;         // number of threads to use for generation
-        int32_t  n_threads_batch;   // number of threads to use for batch processing
+        uint32_t n_ctx;                 // text context, 0 = from model
+        uint32_t n_batch;               // logical maximum batch size that can be submitted to llama_decode
+        uint32_t n_ubatch;              // physical maximum batch size
+        uint32_t n_seq_max;             // max number of sequences (i.e. distinct states for recurrent models)
+        uint32_t n_rs_seq;              // number of recurrent-state snapshots per seq for rollback (0 = no rollback) [EXPERIMENTAL]
+        uint32_t n_outputs_max;         // max outputs in a ubatch (0 = n_batch)
+        uint32_t n_outputs_max_per_seq; // max outputs per sequence (0 = n_outputs_max)
+        int32_t  n_threads;             // number of threads to use for generation
+        int32_t  n_threads_batch;       // number of threads to use for batch processing

        enum llama_context_type      ctx_type;          // set the context type (e.g. MTP)
        enum llama_rope_scaling_type rope_scaling_type; // RoPE scaling type, from `enum llama_rope_scaling_type`
@@ -455,6 +457,8 @@ extern "C" {
    // lora adapter
    struct llama_adapter_lora;

+    LLAMA_API const char * llama_version(void);
+
    // Helpers for getting default parameters
    // TODO: update API to start accepting pointers to params structs (https://github.com/ggml-org/llama.cpp/discussions/9172)
    LLAMA_API struct llama_model_params          llama_model_default_params(void);
@@ -881,6 +885,7 @@ extern "C" {
               const llama_token * tokens,
                          size_t   n_token_count);

+    // If tokens_out is NULL, only the token count is reported through n_token_count_out and no state is loaded
    LLAMA_API size_t llama_state_seq_load_file(
            struct llama_context * ctx,
                      const char * filepath,
@@ -1054,6 +1059,9 @@ extern "C" {
    //

    // Get the backend sampled token for the ith token.
+    // With multiple outputs, sampler state advances when the token is accepted,
+    // not when it is read through this function.
+    // When accepting multiple outputs, accept a contiguous prefix in output order.
    // Returns LLAMA_TOKEN_NULL if no token was sampled.
    LLAMA_API llama_token llama_get_sampled_token_ith(struct llama_context * ctx, int32_t i);

@@ -1270,9 +1278,12 @@ extern "C" {
        // [EXPERIMENTAL]
        // backend sampling interface:

-        // return true if the backend supports all ops needed by the sampler
+        // return true if the backend supports all ops needed by the sampler and can handle up to n_outputs_max_per_seq outputs per sequence
        // note: call once per sampler
-        bool (*backend_init)(struct llama_sampler * smpl, ggml_backend_buffer_type_t buft);
+        bool (*backend_init)(
+                struct llama_sampler       * smpl,
+                ggml_backend_buffer_type_t   buft,
+                uint32_t                     n_outputs_max_per_seq);

        // call after .backend_apply()
        void (*backend_accept)(
@@ -1290,6 +1301,13 @@ extern "C" {

        // called before graph execution to set inputs for the current ubatch
        void (*backend_set_input)(struct llama_sampler * smpl);
+
+        // called before rebuilding a sampling graph to clear any internal sampler state
+        void (*backend_reset)(struct llama_sampler * smpl);
+
+        // copy mutable state from src into dst while keeping dst's references to the current sampling graph
+        // src and dst must have the same type and configuration
+        void (*copy_state)(const struct llama_sampler * src, struct llama_sampler * dst);
    };

    struct llama_sampler {
@@ -1310,6 +1328,7 @@ extern "C" {
    LLAMA_API void                   llama_sampler_apply (      struct llama_sampler * smpl, llama_token_data_array * cur_p);
    LLAMA_API void                   llama_sampler_reset (      struct llama_sampler * smpl);
    LLAMA_API struct llama_sampler * llama_sampler_clone (const struct llama_sampler * smpl);
+    LLAMA_API void                   llama_sampler_copy  (const struct llama_sampler * src, struct llama_sampler * dst);
    // important: do not free if the sampler has been added to a llama_sampler_chain (via llama_sampler_chain_add)
    LLAMA_API void                   llama_sampler_free  (      struct llama_sampler * smpl);

@@ -1499,6 +1518,7 @@ extern "C" {
    LLAMA_API uint32_t llama_sampler_get_seed(const struct llama_sampler * smpl);

    /// @details Sample and accept a token from the idx-th output of the last evaluation
+    // For multiple outputs from one sampler, call this function in output order without gaps.
    //
    // Shorthand for:
    //    const auto * logits = llama_get_logits_ith(ctx, idx);
--- a/Show More
+++ b/Show More