cmake : fix build when GGML_CPU=OFF and GGML_CUDA=ON (#29026 )

* fix: build fails when GGML_CPU=OFF and GGML_CUDA=ON * fix: eol in examples/convert-llama2c-to-ggml/CMakeLists.txt file
gguf-py: fix Q8_1 block size in GGML_QUANT_SIZES (2+2+32) (#29036 )
2026-10-30 08:27:34 -05:00 · 2026-09-18 11:44:03 +03:00 · 2026-09-18 11:43:06 +03:00 · 2026-09-18 10:20:53 +02:00 · 2026-09-18 09:04:14 +02:00 · 2026-09-18 09:00:15 +02:00
82 changed files with 7390 additions and 6640 deletions
@@ -1,5 +1,5 @@
-ARG OPENVINO_VERSION_MAJOR=2026.3.1
-ARG OPENVINO_VERSION_FULL=2026.3.1.22476.56d9685302d
+ARG OPENVINO_VERSION_MAJOR=2026.4
+ARG OPENVINO_VERSION_FULL=2026.4.0.22959.99c81491cc3
 ARG UBUNTU_VERSION=24.04

 # Intel GPU driver versions. https://github.com/intel/compute-runtime/releases
@@ -10,9 +10,9 @@ ARG COMPUTE_RUNTIME_VERSION_FULL=26.31.39395.13-0
 ARG IGDGMM_VERSION=22.10.0

 # Intel NPU driver versions. https://github.com/intel/linux-npu-driver/releases
-ARG NPU_DRIVER_VERSION=v1.35.0
-ARG NPU_DRIVER_FULL=v1.35.0.20260722-29947505341
-ARG LIBZE1_VERSION=1.28.2-1~24.04~ppa1
+ARG NPU_DRIVER_VERSION=v1.38.0
+ARG NPU_DRIVER_FULL=v1.38.0.20260910-34487311128
+ARG LIBZE1_VERSION=1.32.0-1~24.04~ppa1

 # Optional proxy build arguments
 ARG http_proxy=
@@ -173,7 +173,7 @@ RUN --mount=type=cache,target=/var/cache/intel-npu,sharing=locked \
    fi; \
    DEB=/var/cache/intel-npu/libze1_${LIBZE1_VERSION}_amd64.deb; \
    if [ ! -f "$DEB" ]; then \
-        wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260606T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
+        wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260830T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
    fi; \
    mkdir /tmp/npu/ && cd /tmp/npu/ && tar -xf "$TGZ" && cp "$DEB" .; \
    apt-get update; \
@@ -49,7 +49,7 @@ jobs:
          distribution: zulu

      - name: Setup Android SDK
-        uses: android-actions/setup-android@40fd30fb8d7440372e1316f5d1809ec01dcd3699 # v4.0.1
+        uses: android-actions/setup-android@be39fa834029ff78f1a44aa3bb0819b8fc2bd8fd # v4.0.4
        with:
          log-accepted-android-sdk-licenses: false

@@ -123,7 +123,7 @@ jobs:
          distribution: temurin

      - name: Setup Android SDK
-        uses: android-actions/setup-android@40fd30fb8d7440372e1316f5d1809ec01dcd3699 # v4.0.1
+        uses: android-actions/setup-android@be39fa834029ff78f1a44aa3bb0819b8fc2bd8fd # v4.0.4
        with:
          log-accepted-android-sdk-licenses: false

@@ -41,8 +41,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

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

    env:
      # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

    steps:
      - name: Clone
@@ -41,8 +41,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

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

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

    steps:
      - name: Clone
@@ -412,8 +412,8 @@ jobs:

    env:
      # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

    steps:
      - name: Clone
@@ -19,7 +19,7 @@ on:

 jobs:
  check-vendor:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim

    steps:
      - name: Checkout
@@ -11,10 +11,12 @@ on:
    paths:
      - .github/workflows/copilot-setup-steps.yml

+cache-mode: none
+
 jobs:
  # The job MUST be called `copilot-setup-steps` or it will not be picked up by Copilot.
  copilot-setup-steps:
-    runs-on: ubuntu-latest
+    runs-on: ubuntu-24.04

    # Set the permissions to the lowest permissions possible needed for your steps.
    # Copilot will be given its own token for its operations.
@@ -31,8 +33,8 @@ jobs:
      - name: ccache
        uses: ggml-org/ccache-action@v1.2.24
        with:
-          key: copilot-setup-steps
-          evict-old-files: 1d
+          restore: false
+          save: false

      - name: Dependencies
        id: depends
@@ -13,6 +13,16 @@ on:
        required: true
        type: boolean
        default: true
+      skip_apiabi_check:
+        description: 'Skip API/ABI compatibility check'
+        required: false
+        type: boolean
+        default: false
+      apiabi_compare_tag:
+        description: 'Tag to compare against for API/ABI check (default: latest release)'
+        required: false
+        type: string
+        default: ''

 env:
  GH_TOKEN: ${{ github.token }}
@@ -33,12 +43,18 @@ jobs:
          ref: ${{ inputs.commit != '' && inputs.commit || github.ref_name }}
          fetch-depth: 0

+      - name: Install API/ABI check tools
+        if: ${{ github.event.inputs.skip_apiabi_check != 'true' }}
+        run: sudo apt-get install -y abi-compliance-checker abigail-tools
+
      - 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 }}
          RELEASE_BRANCH: ${{ github.ref_name }}
+          SKIP_APIABI_CHECK: ${{ github.event.inputs.skip_apiabi_check }}
+          APIABI_COMPARE_TAG: ${{ github.event.inputs.apiabi_compare_tag }}

      - name: Create release tag
        if: ${{ github.event.inputs.dry_run == 'false' }}
@@ -12,7 +12,7 @@ on:

 jobs:
    pre-tokenizer-hashes:
-        runs-on: [self-hosted, fast]
+        runs-on: ubuntu-slim

        steps:
        - name: Checkout repository
@@ -20,7 +20,7 @@ concurrency:

 jobs:
  python-check-requirements:
-    runs-on: [self-hosted, CPU, fast]
+    runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }}
    name: check-requirements
    steps:
      - name: Check out source repository
@@ -21,7 +21,7 @@ concurrency:

 jobs:
  flake8-lint:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
    name: Lint
    steps:
      - name: Check out source repository
@@ -22,7 +22,7 @@ concurrency:

 jobs:
  python-type-check:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
    name: python type-check
    steps:
      - name: Check out source repository
@@ -106,6 +106,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-${{ matrix.os }}-${{ matrix.arch }}
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -190,6 +191,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-${{ matrix.os }}-cpu
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -275,6 +277,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-${{ matrix.os }}-vulkan
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -481,10 +484,9 @@ jobs:
          distribution: temurin

      - name: Setup Android SDK
-        uses: android-actions/setup-android@40fd30fb8d7440372e1316f5d1809ec01dcd3699 # v4.0.1
+        uses: android-actions/setup-android@be39fa834029ff78f1a44aa3bb0819b8fc2bd8fd # v4.0.4
        with:
          log-accepted-android-sdk-licenses: false
-          packages: 'platform-tools'

      - name: Install NDK
        run: |
@@ -501,6 +503,7 @@ jobs:
      #  uses: ggml-org/ccache-action@v1.2.24
      #  with:
      #    key: release-android-arm64
+      #    evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -555,8 +558,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

    steps:
      - name: Set OpenVINO version output
@@ -579,6 +582,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-ubuntu-24.04-openvino-release-no-preset-v1
+          evict-old-files: 1d

      - name: Dependencies
        run: |
@@ -669,8 +673,8 @@ jobs:

    env:
      # Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
-      OPENVINO_VERSION_MAJOR: "2026.3.1"
-      OPENVINO_VERSION_FULL: "2026.3.1.22476.56d9685302d"
+      OPENVINO_VERSION_MAJOR: "2026.4"
+      OPENVINO_VERSION_FULL: "2026.4.0.22959.99c81491cc3"

    steps:
      - name: Set OpenVINO version output
@@ -822,6 +826,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-windows-2025-vs2026-${{ matrix.arch }}-cpu
+          evict-old-files: 1d

      - name: Build
        shell: cmd
@@ -1066,6 +1071,7 @@ jobs:
      #  uses: ggml-org/ccache-action@v1.2.24
      #  with:
      #    key: release-windows-2025-${{ matrix.arch }}-${{ matrix.backend }}
+      #    evict-old-files: 1d

      - name: Install OpenCL Headers and Libs
        id: install_opencl
@@ -1154,6 +1160,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -1250,6 +1257,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-windows-2022-x64-sycl
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -1368,6 +1376,7 @@ jobs:
        uses: ggml-org/ccache-action@v1.2.24
        with:
          key: release-ubuntu-24.04-sycl-${{ matrix.build }}
+          evict-old-files: 1d

      - name: Build
        id: cmake_build
@@ -16,7 +16,7 @@ on:

 jobs:
    update-ops-docs:
-        runs-on: [self-hosted, fast, ARM64]
+        runs-on: ubuntu-slim

        steps:
        - name: Checkout repository
@@ -97,7 +97,6 @@
 /src/models/                            @CISC
 /tests/                                 @ggerganov
 /tests/test-chat.*                      @pwilkin
-/tests/test-llama-archs.cpp             @JohannesGaessler
 /tools/batched-bench/                   @ggerganov
 /tools/cli/                             @ngxson
 /tools/completion/                      @ggerganov
@@ -104,6 +104,12 @@ common_chat_params common_chat_params_init_deepseek_v3_2(const common_chat_templ
    const std::string GEN_PROMPT   = "<｜Assistant｜>";
    const std::string TC_SEPARATOR = "\n\n";

+    // lets the server find user turns in the prompt and place context checkpoints there
+    data.message_delimiters = {
+        { COMMON_CHAT_ROLE_ASSISTANT, GEN_PROMPT },
+        { COMMON_CHAT_ROLE_USER, "<｜User｜>" },
+    };
+
    data.prompt = common_chat_template_direct_apply_impl(
        tmpl, inputs, adjusted_messages, std::nullopt, additional_context);
    data.generation_prompt = common_chat_template_generation_prompt_impl(
@@ -12,6 +12,8 @@ The OpenVINO backend is implemented in `ggml/src/ggml-openvino` and provides a t
 - Compiles and caches the model for the target device.
 - Binds GGML tensor memory to OpenVINO inference tensors and runs inference.

+For guidance on contributing to the OpenVINO backend, see the [OpenVINO Backend Contributing Guide](https://github.com/ravi9/llamacpp-ov-dev-guide/blob/main/contributing-llamacpp-ov.md).
+
 ## Contents

 - [Supported Devices](#supported-devices)
@@ -96,7 +98,7 @@ Although, the validated models below were tested with `llama-cli` using the `Q4_
  - **SL** = Stateless (`GGML_OPENVINO_STATEFUL_EXECUTION=0`)
  - **SF** = Stateful (`GGML_OPENVINO_STATEFUL_EXECUTION=1`)
  - Note: The NPU operates in stateless mode only.
- **Validation system:** Intel® Core™ Ultra 5 238V (Lunar Lake) | 32 GB RAM | Ubuntu 24.04 | Intel OpenCL GPU Driver 26.31.39395.13-0 | Intel NPU Driver 1.35.0.
+- **Validation system:** Intel® Core™ Ultra 5 238V (Lunar Lake) | 32 GB RAM | Ubuntu 24.04 | Intel Graphics Compiler 2.41.5 | Intel OpenCL GPU Driver 26.31.39395.13-0 | Intel NPU Driver 1.38.0.
 - See [Known Limitations](#known-limitations) for context on observed failures.

 | Model | CPU (SL / SF) | GPU (SL / SF) | NPU (SL) |
@@ -117,9 +119,9 @@ Although, the validated models below were tested with `llama-cli` using the `Q4_
 | [lmstudio-community/Qwen3.5-9B-Q4_K_M](https://huggingface.co/lmstudio-community/Qwen3.5-9B-GGUF) | ✓ / ✗ | ✓ / ✗ | ✗ |
 |  |  |  |  |
 | [unsloth/gemma-3-4b-it-Q4_K_M](https://huggingface.co/unsloth/gemma-3-4b-it-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
-| [bartowski/google_gemma-4-E2B-it-Q4_K_M](https://huggingface.co/bartowski/google_gemma-4-E2B-it-GGUF) | ✓ / ✗ | ✓ / ✗ | ✗ |
-| [bartowski/google_gemma-4-E4B-it-Q4_K_M](https://huggingface.co/bartowski/google_gemma-4-E4B-it-GGUF) | ✓ / ✗ | ✓ / ✗ | ✓ |
-| [bartowski/gemma-4-12B-it-Q4_K_M](https://huggingface.co/bartowski/gemma-4-12B-it-GGUF) | ✓ / ✗ | ✓ / ✗ | ✓ |
+| [bartowski/google_gemma-4-E2B-it-Q4_K_M](https://huggingface.co/bartowski/google_gemma-4-E2B-it-GGUF) | ✓ / ✓ | ✓ / ✓ | ✗ |
+| [bartowski/google_gemma-4-E4B-it-Q4_K_M](https://huggingface.co/bartowski/google_gemma-4-E4B-it-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
+| [bartowski/gemma-4-12B-it-Q4_K_M](https://huggingface.co/bartowski/gemma-4-12B-it-GGUF) | ✓ / ✓ | ✓ / ✓ | ✗ |
 |  |  |  |  |
 | [bartowski/Phi-3-mini-4k-instruct-Q4_K_M](https://huggingface.co/bartowski/Phi-3-mini-4k-instruct-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
 | [bartowski/Phi-3.5-mini-instruct-Q4_K_M](https://huggingface.co/bartowski/Phi-3.5-mini-instruct-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
@@ -132,9 +134,9 @@ Although, the validated models below were tested with `llama-cli` using the `Q4_
 | [bartowski/DeepSeek-R1-Distill-Llama-8B-Q4_K_M](https://huggingface.co/bartowski/DeepSeek-R1-Distill-Llama-8B-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
 | [bartowski/DeepSeek-R1-Distill-Qwen-7B-Q4_K_M](https://huggingface.co/bartowski/DeepSeek-R1-Distill-Qwen-7B-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
 |  |  |  |  |
-| [ibm-granite/granite-4.0-350m-Q4_K_M](https://huggingface.co/ibm-granite/granite-4.0-350m-GGUF) | ✓ / ✓ | ✗ / ✗ | ✓ |
+| [ibm-granite/granite-4.0-350m-Q4_K_M](https://huggingface.co/ibm-granite/granite-4.0-350m-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
 | [ibm-granite/granite-4.0-micro-Q4_K_M](https://huggingface.co/ibm-granite/granite-4.0-micro-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
-| [ibm-granite/granite-4.0-1b-Q4_K_M](https://huggingface.co/ibm-granite/granite-4.0-1b-GGUF) | ✓ / ✓ | ✗ / ✗ | ✗ |
+| [ibm-granite/granite-4.0-1b-Q4_K_M](https://huggingface.co/ibm-granite/granite-4.0-1b-GGUF) | ✓ / ✓ | ✓ / ✓ | ✗ |
 | [ibm-research/granite-3.2-8b-instruct-Q4_K_M](https://huggingface.co/ibm-research/granite-3.2-8b-instruct-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
 |  |  |  |  |
 | [HuggingFaceTB/smollm2-1.7b-instruct-q4_k_m](https://huggingface.co/HuggingFaceTB/SmolLM2-1.7B-Instruct-GGUF) | ✓ / ✓ | ✓ / ✓ | ✓ |
@@ -242,8 +244,8 @@ chmod +x build-llamacpp-ov.sh
 # ============================================
 set -euo pipefail

-OPENVINO_VERSION_MAJOR="2026.3.1"
-OPENVINO_VERSION_FULL="2026.3.1.22476.56d9685302d"
+OPENVINO_VERSION_MAJOR="2026.4"
+OPENVINO_VERSION_FULL="2026.4.0.22959.99c81491cc3"

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

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

 </details>

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

-set "OPENVINO_VERSION_MAJOR=2026.3.1"
-set "OPENVINO_VERSION_FULL=2026.3.1.22476.56d9685302d"
+set "OPENVINO_VERSION_MAJOR=2026.4"
+set "OPENVINO_VERSION_FULL=2026.4.0.22959.99c81491cc3"

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

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

 </details>

@@ -1,5 +1,7 @@
-set(TARGET llama-convert-llama2c-to-ggml)
-add_executable(${TARGET} convert-llama2c-to-ggml.cpp)
-install(TARGETS ${TARGET} RUNTIME)
-target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
-target_compile_features(${TARGET} PRIVATE cxx_std_17)
+if (GGML_CPU)
+    set(TARGET llama-convert-llama2c-to-ggml)
+    add_executable(${TARGET} convert-llama2c-to-ggml.cpp)
+    install(TARGETS ${TARGET} RUNTIME)
+    target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
+    target_compile_features(${TARGET} PRIVATE cxx_std_17)
+endif()
@@ -25,7 +25,7 @@ GGML_BACKEND_API bool ggml_backend_is_sycl(ggml_backend_t backend);
 GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_sycl_buffer_type(int device);

 // split tensor buffer that splits matrices by rows across multiple devices
-GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split);
+GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type([[maybe_unused]] int main_device, const float * tensor_split);

 // Tensor parallelism (--split-mode tensor): comm_init/free/allreduce_tensor
 // trio queried by the meta-backend via ggml_backend_reg_get_proc_address.
@@ -1440,6 +1440,7 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) {

 static bool use_adreno_bin_kernels(ggml_backend_opencl_context * backend_ctx) {
 #ifndef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+    GGML_UNUSED(backend_ctx);
    return false;
 #else
    if (backend_ctx->gpu_family != GPU_FAMILY::ADRENO) {
@@ -6309,6 +6310,8 @@ static void ggml_opencl_print_backend_info(ggml_backend_opencl_device_context *

    auto * backend_ctx = dev_ctx->backend_ctx;

+    GGML_LOG_INFO("ggml_opencl: OpenCL device: %s\n",
+        backend_ctx->device_name.c_str());
    GGML_LOG_INFO("ggml_opencl: OpenCL driver: %s\n",
        backend_ctx->driver_version.c_str());
    GGML_LOG_INFO("ggml_opencl: vector subgroup broadcast support: %s\n",
@@ -6325,11 +6328,11 @@ static void ggml_opencl_print_backend_info(ggml_backend_opencl_device_context *
        backend_ctx->global_mem_size/1024/1024);
    GGML_LOG_INFO("ggml_opencl: max mem alloc size: %zu MB\n",
        backend_ctx->max_alloc_size/1024/1024);
-    GGML_LOG_INFO("ggml_opencl: device max image buffer size (pixels): %lu\n",
+    GGML_LOG_INFO("ggml_opencl: device max image buffer size (pixels): %zu\n",
        backend_ctx->image_max_buffer_size);
-    GGML_LOG_INFO("ggml_opencl: device max image2d size: %lu x %lu\n",
+    GGML_LOG_INFO("ggml_opencl: device max image2d size: %zu x %zu\n",
        backend_ctx->image2d_max_width, backend_ctx->image2d_max_height);
-    GGML_LOG_INFO("ggml_opencl: device max workgroup size: %lu\n",
+    GGML_LOG_INFO("ggml_opencl: device max workgroup size: %zu\n",
        backend_ctx->max_workgroup_size);
    GGML_LOG_INFO("ggml_opencl: SVM coarse grain buffer support: %s\n",
        backend_ctx->svm_caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER ? "true" : "false");
@@ -7627,7 +7630,7 @@ static void ggml_cl_moe_bias_glu_fused(ggml_backend_t backend, ggml_tensor * gat
    size_t global_work_size[] = { (size_t)glu->ne[1]*nth, (size_t)glu->ne[2], 1 };
    size_t local_work_size[]  = { (size_t)nth, 1, 1 };

-    backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, (ggml_tensor *)glu);
+    backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, glu);
 }

 // Fusion B: the MoE down-projection bias add feeding the combine.
@@ -7771,7 +7774,7 @@ static void ggml_cl_moe_bias_combine_fused(ggml_backend_t backend, const ggml_te

    size_t lws[2] = { 64, 1 };
    size_t gws[2] = { (size_t)(((n_embd4 + 63) / 64) * 64), (size_t)nt };
-    backend_ctx->enqueue_ndrange_kernel(kernel, 2, gws, lws, (ggml_tensor *)dst);
+    backend_ctx->enqueue_ndrange_kernel(kernel, 2, gws, lws, dst);
 }


@@ -8114,7 +8117,6 @@ static void ggml_cl_mul_mat_q4_k_glu_fused(ggml_backend_t backend, ggml_tensor *
    GGML_UNUSED(gate_tensor);
    GGML_UNUSED(up_tensor);
    GGML_UNUSED(glu_tensor);
-    GGML_ABORT("q4_K GLU fusion requires GGML_OPENCL_USE_ADRENO_KERNELS");
 #endif
 }

@@ -11336,7 +11338,6 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,

 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
        if (use_adreno_moe_kernels(backend_ctx, tensor)) {
-            cl_int err;
            cl_kernel kernel = backend_ctx->kernel_restore_block_q4_0_trans4_ns;

            cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
@@ -11535,7 +11536,6 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,

 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
        if (use_adreno_moe_kernels(backend_ctx, tensor)) {
-            cl_int err;
            // TODO: use ggml_cl_buffer to manage this temporary buffer
            cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
            GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
@@ -11638,7 +11638,6 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,

 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
        if (use_adreno_moe_kernels(backend_ctx, tensor)) {
-            cl_int err;
            // TODO: use ggml_cl_buffer to manage this temporary buffer
            cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
            GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
@@ -245,7 +245,7 @@ void GgmlOvDecoder::set_input_output() {
            if (src->op == GGML_OP_VIEW) {
                // Traverse upward through nested VIEW operations
                std::remove_reference_t<decltype(current_node_info.node_inputs_views[src_name])> view_chain;
-                auto current = src;
+                auto * current = src;

                while (current != nullptr) {
                    auto current_name = get_tensor_ov_name(m_cgraph, current);
@@ -612,9 +612,8 @@ std::pair<ModelParams, ComputeParams> GgmlOvDecoder::compute_llm_params(ggml_cgr
                if (node->src[1]->view_src != nullptr) {
                    if (node->src[3] != nullptr) {
                        return 4;  // decoder self-attention
-                    } else {
-                        return 5;  // cross-attention or encoder self-attention
-                    };
+                    }
+                    return 5;      // cross-attention or encoder self-attention
                }
                break;
            default:
@@ -736,8 +735,7 @@ std::pair<ModelParams, ComputeParams> GgmlOvDecoder::compute_llm_params(ggml_cgr

    bool rope_seen = false;
    for (int i = 0; i < cgraph->n_nodes; i++) {
-        auto * node = cgraph->nodes[i];
-        std::string name = std::string(node->name);
+        ggml_tensor * node = cgraph->nodes[i];
        const int attention_pattern_case = get_attention_pattern_case(node);
        if (attention_pattern_case != -1) {
            ggml_tensor * cache_k_permute = nullptr;
@@ -948,7 +946,6 @@ ov::PartialShape GgmlOvDecoder::get_graph_input_shape(const ggml_tensor * op,
    if (m_naive) {
        return input != nullptr ? ov::PartialShape{get_shape(input)} : ov::PartialShape{get_shape(op)};
    }
-    auto name = std::string(input->name);
    ov::PartialShape input_shape;

    if (is_inp_tok(input, op) || is_inp_pos(input, op)) {
@@ -1474,7 +1471,7 @@ std::shared_ptr<ov::Node> GgmlOvDecoder::create_weight_node(ggml_tensor * tensor
 void GgmlOvDecoder::dump_cgraph(const ggml_cgraph * cgraph, std::string & filename) {
    std::ofstream file(filename);
    if (!file.is_open()) {
-        std::cerr << "Failed to open file" << std::endl;
+        std::cerr << "Failed to open file" << '\n';
        return;
    }

@@ -1580,11 +1577,11 @@ void print_tensor_address_map(const ggml_cgraph * cgraph) {
        }
    }
    for (const auto & pair : address_map) {
-        std::cout << "Address: " << pair.first << std::endl;
+        std::cout << "Address: " << pair.first << '\n';
        for (const auto & name : pair.second) {
            std::cout << name << " ; ";
        }
-        std::cout << std::endl << std::endl;
+        std::cout << "\n\n";
    }
 }

@@ -2226,7 +2223,7 @@ void GgmlOvDecoder::compute_node_dynamic_dims() {
                    std::cout << ", ";
                }
            }
-            std::cout << "]" << std::endl;
+            std::cout << "]" << '\n';
            // print the src name & shape with the dynamic dim for debugging
            for (int j = 0; j < GGML_MAX_SRC; j++) {
                ggml_tensor * src = node->src[j];
@@ -2245,9 +2242,9 @@ void GgmlOvDecoder::compute_node_dynamic_dims() {
                        std::cout << ", ";
                    }
                }
-                std::cout << "]" << std::endl;
+                std::cout << "]" << '\n';
            }
-            std::cout << std::endl;
+            std::cout << '\n';
        }
    }
 }
@@ -354,41 +354,41 @@ public:

    void update_io(ggml_cgraph * cgraph);

-    inline static bool is_inp_tok(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_tok(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_GET_ROWS && tensor == op->src[1] && op->src[0]->op == GGML_OP_NONE;
    }

-    inline static bool is_inp_pos(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_pos(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_ROPE && tensor == op->src[1];
    }

    // IMROPE packs 4 stacked position planes (t/h/w/e) into inp_pos, each of length
    // n_tokens; other modes carry a single position per token.
-    inline static int get_inp_pos_n_planes(const ggml_tensor * op) {
+    static int get_inp_pos_n_planes(const ggml_tensor * op) {
        return op->op_params[2] == GGML_ROPE_TYPE_IMROPE ? 4 : 1;
    }

-    inline static bool is_inp_emb(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_emb(const ggml_tensor * tensor, const ggml_tensor * op) {
        return tensor->op == GGML_OP_GET_ROWS && op->op == GGML_OP_RMS_NORM;
    }

-    inline static bool is_inp_mask(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_mask(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_CPY || (op->op == GGML_OP_FLASH_ATTN_EXT && tensor == op->src[3]) ||
               (op->op == GGML_OP_SOFT_MAX && tensor == op->src[1]);
    }

-    inline static bool is_inp_mean(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_mean(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_MUL_MAT && tensor == op->src[1] && tensor->op == GGML_OP_NONE &&
               (tensor->flags & GGML_TENSOR_FLAG_INPUT) && tensor->type == GGML_TYPE_F32 &&
               op->src[0] != nullptr && op->src[0]->op != GGML_OP_NONE;
    }

-    inline static bool is_rope_freqs_weight(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_rope_freqs_weight(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_ROPE && tensor == op->src[2];
    }

    // also returns true for cache_s and cache_r in SSM/DeltaNet models
-    inline static bool is_kvcache(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_kvcache(const ggml_tensor * tensor, const ggml_tensor * op) {
        if (tensor == nullptr) {
            return false;
        }
@@ -396,14 +396,14 @@ public:
               (op != nullptr && op->op == GGML_OP_SET_ROWS && op->src[2] == tensor);
    }

-    inline static bool is_conv_state_writeback(const ggml_tensor * node) {
+    static bool is_conv_state_writeback(const ggml_tensor * node) {
        return node->op == GGML_OP_CPY && node->view_src != nullptr && is_kvcache(node->view_src, nullptr) &&
               node->src[0] != nullptr && node->src[0]->op == GGML_OP_VIEW && node->src[0]->src[0] != nullptr &&
               node->src[0]->src[0]->op == GGML_OP_CONCAT && node->src[1] != nullptr &&
               node->src[1]->op == GGML_OP_VIEW && node->src[1]->view_src == node->view_src;
    }

-    inline static bool is_kv_idx(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_kv_idx(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_SET_ROWS && op->src[1] == tensor;
    }

@@ -411,13 +411,13 @@ public:
        return m_model_params.swa_mask != nullptr && tensor == m_model_params.swa_mask;
    }

-    inline static bool is_output_idx(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_output_idx(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_GET_ROWS && tensor == op->src[1] && op->src[0]->op != GGML_OP_NONE &&
               op->src[1]->op == GGML_OP_NONE;
    }

    // the state permutation index input used in SSM/DeltaNet models (inp->s_copy in llama-graph.cpp)
-    inline static bool is_inp_s_copy(const ggml_tensor * tensor, const ggml_tensor * op) {
+    static bool is_inp_s_copy(const ggml_tensor * tensor, const ggml_tensor * op) {
        return op->op == GGML_OP_GET_ROWS && tensor == op->src[1] &&
               op->src[0]->buffer->usage == GGML_BACKEND_BUFFER_USAGE_ANY;
    }
@@ -481,5 +481,3 @@ private:
 };

 void print_tensor_address_map(const ggml_cgraph * cgraph);
-
-std::optional<int> extract_layer_from_name(const std::string & name);
@@ -472,10 +472,6 @@ ggml_openvino_extracted_layout ggml_openvino_get_extracted_layout(const ggml_ten

    switch (tensor->type) {
    case GGML_TYPE_MXFP4:
-        layout.is_u4 = true;
-        layout.is_symmetric = true;
-        break;
-
    case GGML_TYPE_Q4_0:
        layout.is_u4 = true;
        layout.is_symmetric = true;
@@ -28,12 +28,7 @@
 #include <string>
 #include <vector>

-#ifndef _WIN32
-#    include <sys/mman.h>
-#    include <unistd.h>
-#endif
-
-#if defined(_WIN32)
+#ifdef _WIN32
 #    define WIN32_LEAN_AND_MEAN
 #    ifndef NOMINMAX
 #        define NOMINMAX
@@ -61,6 +56,7 @@
 // - CPU repack buffer: tensor->extra stores tensor_traits with repacked data
 // =====================================================

+namespace {
 // Buffer context that manages per-tensor allocations (no contiguous buffer for weights)
 struct ggml_backend_openvino_buffer_context {
    int device;
@@ -199,6 +195,7 @@ struct ggml_backend_openvino_buffer_type_context {
    int device;
    std::string name;
 };
+}  // namespace

 // =====================================================
 // Host weight-buffer release (GGML_OPENVINO_RELEASE_WEIGHTS)
@@ -258,14 +255,16 @@ void ggml_openvino_release_weight_buffers() {
    for (const auto & b : reg.buffers) {
        // Align down/up to page boundaries so madvise only drops whole pages
        // fully owned by this buffer.
-        const long page = sysconf(_SC_PAGESIZE);
-        uintptr_t start = reinterpret_cast<uintptr_t>(b.first);
-        uintptr_t end = start + b.second;
-        uintptr_t astart = (start + page - 1) & ~(uintptr_t) (page - 1);
-        uintptr_t aend = end & ~(uintptr_t) (page - 1);
-        if (aend > astart) {
-            if (madvise(reinterpret_cast<void *>(astart), aend - astart, MADV_DONTNEED) == 0) {
-                total += aend - astart;
+        const size_t page = (size_t) sysconf(_SC_PAGESIZE);
+        const uintptr_t ustart = reinterpret_cast<uintptr_t>(b.first);
+        const size_t offset_to_page = (page - (ustart & (page - 1))) & (page - 1);
+        if (b.second > offset_to_page) {
+            const size_t aligned_len = (b.second - offset_to_page) & ~(page - 1);
+            if (aligned_len > 0) {
+                char * astart = static_cast<char *>(b.first) + offset_to_page;
+                if (madvise(astart, aligned_len, MADV_DONTNEED) == 0) {
+                    total += aligned_len;
+                }
            }
        }
    }
@@ -876,11 +875,13 @@ GGML_BACKEND_API bool ggml_backend_is_openvino(ggml_backend_t backend) {
    return backend != NULL && ggml_guid_matches(backend->guid, ggml_backend_openvino_guid());
 }

+namespace {
 struct ggml_backend_openvino_device_context {
    int device;
    std::string name;
    std::string description;
 };
+}

 static const char * ggml_backend_openvino_device_get_name(ggml_backend_dev_t dev) {
    ggml_backend_openvino_device_context * ctx = (ggml_backend_openvino_device_context *) dev->context;
@@ -1588,9 +1589,11 @@ static const struct ggml_backend_device_i ggml_backend_openvino_device_interface
    /* .event_synchronize    = */ NULL,
 };

+namespace {
 struct ggml_backend_openvino_reg_context {
    std::vector<ggml_backend_dev_t> devices;
 };
+}

 static const char * ggml_backend_openvino_reg_get_name(ggml_backend_reg_t reg) {
    return GGML_OPENVINO_NAME;
@@ -34,6 +34,15 @@
 #include <string>
 #include <vector>

+// From <openvino>/src/common/transformations/include/transformations/utils/utils.hpp
+namespace ov::op::util {
+// From <openvino>/src/common/transformations/include/transformations/utils/utils.hpp
+bool get_single_value(const std::shared_ptr<ov::op::v0::Constant> & const_node,
+                      float & value,
+                      bool check_value_range = true);
+}  // namespace ov::op::util
+
+namespace {
 void unpack_32_4(const uint8_t * data, uint8_t * dst) {
    std::fill_n(dst, 16, 0);
    for (int j = 0; j < 16; ++j) {
@@ -48,11 +57,11 @@ void unpack_32_4(const uint8_t * data, uint8_t * dst) {
    }
 }

-static constexpr size_t MXFP4_BLOCK_SIZE = 32;
-static constexpr size_t MXFP4_BLOCK_QS_SIZE = MXFP4_BLOCK_SIZE / 2;
-static constexpr size_t MXFP4_BLOCK_BYTES = sizeof(uint8_t) + MXFP4_BLOCK_QS_SIZE;
+constexpr size_t MXFP4_BLOCK_SIZE = 32;
+constexpr size_t MXFP4_BLOCK_QS_SIZE = MXFP4_BLOCK_SIZE / 2;
+constexpr size_t MXFP4_BLOCK_BYTES = sizeof(uint8_t) + MXFP4_BLOCK_QS_SIZE;

-static void pack_32_mxfp4_for_openvino(const uint8_t * data, uint8_t * dst) {
+void pack_32_mxfp4_for_openvino(const uint8_t * data, uint8_t * dst) {
    for (int j = 0; j < static_cast<int>(MXFP4_BLOCK_QS_SIZE); j += 2) {
        const uint8_t v0 = data[j] & 0x0F;
        const uint8_t v1 = (data[j + 1] & 0x0F) << 4;
@@ -419,7 +428,7 @@ void extract_q6_k_data(const ggml_tensor * tensor,
    }
 }

-static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t * d, uint8_t * m) {
+inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t * d, uint8_t * m) {
    if (j < 4) {
        *d = q[j] & 63;
        *m = q[j + 4] & 63;
@@ -514,9 +523,9 @@ void extract_q5_k_data(const ggml_tensor * tensor,
 ov::Output<ov::Node> make_int8_weights(ov::Tensor & weight,
                                       ov::Tensor & scales,
                                       ov::Tensor & zp,
-                                       size_t group_size,
-                                       bool use_bias,
-                                       bool for_gather_matmul) {
+                                       size_t group_size = GGML_QUANTIZATION_GROUP_SIZE,
+                                       bool use_bias = false,
+                                       bool for_gather_matmul = false) {
    ov::Shape orig_shape = weight.get_shape();
    bool is_signed = (weight.get_element_type() == ov::element::i8);  // Symmetric: signed weights, no ZP

@@ -611,13 +620,24 @@ ov::Output<ov::Node> make_int8_weights(ov::Tensor & weight,
    return std::make_shared<ov::op::v0::Convert>(result, ov::element::f32);
 }

+// If for_gather_matmul is true, the weight tensor may be N-D (e.g. 3D MoE expert weights
+// [n_expert, rows, cols]). The dequantization chain (Convert->[Subtract]->Multiply) is built as
+// usual but left in f16 (no final Convert to f32) -- ov::pass::MarkDequantization (registered in
+// translate_session.cpp) marks the chain so it survives model-build-time ConstantFolding -- see
+// make_int8_weights.cpp/make_int4_weights.cpp. mul_mat_id.cpp constructs ov::op::internal::GatherMatmul
+// directly from the resulting f16 dequant chain.
+//
+// When use_bias is true (explicitly, or implicitly because for_gather_matmul is true), the zp
+// tensor is expected to hold an exact f16 bias value (rather than a rounded integer zero point);
+// it is converted in place into an exact zero_point = -bias/scale and consumed via Subtract, not
+// Add, so the chain still matches OpenVINO's Convert->Subtract->Multiply decompression pattern.
 // See make_int8_weights for the meaning of for_gather_matmul.
 ov::Output<ov::Node> make_int4_weights(ov::Tensor & weight,
                                       ov::Tensor & scales,
                                       ov::Tensor & zp,
-                                       size_t group_size,
-                                       bool use_bias,
-                                       bool for_gather_matmul) {
+                                       size_t group_size = GGML_QUANTIZATION_GROUP_SIZE,
+                                       bool use_bias = false,
+                                       bool for_gather_matmul = false) {
    ov::Shape orig_weight_shape = weight.get_shape();
    bool is_signed = (weight.get_element_type() == ov::element::i4);  // Symmetric: signed weights, no ZP

@@ -746,357 +766,6 @@ ov::Output<ov::Node> make_mxfp4_moe_packed_weights(ov::Tensor & weight) {
    return weights_node;
 }

-// Extract quantized weights from tensor and create weight subgraph
-std::shared_ptr<ov::Node> extract_quantized_weights(const ggml_tensor * tensor,
-                                                    const void * data,
-                                                    ov::Tensor & weights,
-                                                    ov::Tensor & scales,
-                                                    ov::Tensor & zp,
-                                                    bool use_bias) {
-    // Create a temporary tensor for extraction functions that read from tensor->data
-    ggml_tensor temp_tensor = *tensor;
-    temp_tensor.data = const_cast<void *>(data);
-
-    if (tensor->type == GGML_TYPE_MXFP4) {
-        extract_mxfp4_data(&temp_tensor, weights, scales);
-        auto result = make_mxfp4_weights(weights, scales).get_node_shared_ptr();
-        result->set_friendly_name(tensor->name);
-        return result;
-    }
-
-    // Determine block size based on tensor type
-    int64_t weights_per_block;
-    bool is_u4;
-    switch (tensor->type) {
-    case GGML_TYPE_Q4_0:
-    case GGML_TYPE_Q4_1:
-    case GGML_TYPE_Q4_K:
-        is_u4 = true;
-        weights_per_block = 32;
-        break;
-    case GGML_TYPE_Q8_0:
-    case GGML_TYPE_Q5_1:
-    case GGML_TYPE_Q5_K:
-        is_u4 = false;
-        weights_per_block = 32;
-        break;
-    case GGML_TYPE_Q6_K:
-        is_u4 = false;
-        weights_per_block = 16;
-        break;
-    default:
-        throw std::runtime_error("Unsupported quantized type for extraction: " +
-                                 std::string(ggml_type_name(tensor->type)));
-    }
-
-    // 3D MoE expert weights (for_gather_matmul) always use the exact f16 zero-point extraction
-    // (see make_int8_weights/make_int4_weights) rather than the rounded integer zero point --
-    // round(min/scale) error is what corrupts Q4_K/Q5_1 experts, and the f16-zp form still fuses
-    // into GatherMatmulCompressed since it stays a Subtract, not an Add.
-    const bool for_gather_matmul = tensor->ne[2] > 1;
-    use_bias = use_bias || for_gather_matmul;
-
-    // Extract quantized data
-    switch (tensor->type) {
-    case GGML_TYPE_Q4_0:
-        extract_q4_0_data(&temp_tensor, weights, scales, zp);
-        break;
-    case GGML_TYPE_Q4_1:
-        extract_q4_1_data(&temp_tensor, weights, scales, zp, use_bias);
-        break;
-    case GGML_TYPE_Q4_K:
-        extract_q4_k_data(&temp_tensor, weights, scales, zp, use_bias);
-        break;
-    case GGML_TYPE_Q5_1:
-        extract_q5_1_data(&temp_tensor, weights, scales, zp, use_bias);
-        break;
-    case GGML_TYPE_Q8_0:
-        extract_q8_0_data(&temp_tensor, weights, scales, zp);
-        break;
-    case GGML_TYPE_Q6_K:
-        extract_q6_k_data(&temp_tensor, weights, scales, zp);
-        break;
-    case GGML_TYPE_Q5_K:
-        extract_q5_k_data(&temp_tensor, weights, scales, zp, use_bias);
-        break;
-    default:
-        throw std::runtime_error("Unsupported quantized type: " + std::string(ggml_type_name(tensor->type)));
-    }
-
-    // Create the OpenVINO weight subgraph. 3D expert weights (MoE) are routed through the
-    // GatherMatmul-oriented path: dequantized in f16, with constant folding disabled on the chain.
-    ov::Output<ov::Node> weight_node;
-    if (is_u4) {
-        weight_node = make_int4_weights(weights, scales, zp, weights_per_block, use_bias, for_gather_matmul);
-    } else {
-        weight_node = make_int8_weights(weights, scales, zp, weights_per_block, use_bias, for_gather_matmul);
-    }
-
-    auto result = weight_node.get_node_shared_ptr();
-    result->set_friendly_name(tensor->name);
-    return result;
-}
-
-// Requantize weights to target format, writing to provided buffers
-std::shared_ptr<ov::Node> requantize_to_buffers(const ggml_tensor * tensor,
-                                                const void * data,
-                                                ExtraQuantType requant_type,
-                                                int64_t block_size,
-                                                ov::Tensor & weights,
-                                                ov::Tensor & scales,
-                                                ov::Tensor & zp) {
-    int64_t n_elements = ggml_nelements(tensor);
-    const int64_t ne0 = tensor->ne[0];                 // elements per row
-    const int64_t n_rows = n_elements / ne0;
-    const auto * type_traits = ggml_get_type_traits(tensor->type);
-    const size_t src_row_bytes = ggml_row_size(tensor->type, ne0);
-
-    bool is_u4 = (requant_type == ExtraQuantType::Q4_0_C || requant_type == ExtraQuantType::Q4_0_128 ||
-                  requant_type == ExtraQuantType::Q4_0_64 || requant_type == ExtraQuantType::Q4_1_64);
-
-    // Streaming dequant (opt-in via GGML_OPENVINO_REDUCE_COMPILE_MEM or
-    // GGML_OPENVINO_MEMORY_OPTIMIZE): instead of
-    // materializing the full n_elements F32 array (e.g. ~1 GB for token_embd), dequantize
-    // a chunk of complete rows into a small scratch and quantize/convert it straight into
-    // the output buffers, capping the transient F32 footprint at CHUNK_ROWS*ne0 floats.
-    //
-    // Only valid (and only used) for the Q8_0_C / Q8_1_C / F16 targets whose block size
-    // divides a row (channel-wise _C uses block_size == ne0) so no target block straddles
-    // a row boundary, and Q8/F16 have no cross-block packing. The u4 (Q4_0) path packs two
-    // weights per byte with running zp ORs that assume a single whole-array call, so it is
-    // never streamed. When the flag is off, behavior is identical to the original
-    // full-materialization path.
-    const bool stream_requant = ggml_openvino_reduce_compile_mem_enabled() && !is_u4 &&
-                                !(block_size > 0 && ne0 % block_size != 0);
-
-    if (!stream_requant) {
-        // Full materialization (original behavior): dequantize the whole tensor to F32,
-        // then convert/quantize in one call.
-        std::vector<float> weights_f32(n_elements);
-        type_traits->to_float(data, weights_f32.data(), n_elements);
-        if (requant_type == ExtraQuantType::F16) {
-            ggml_get_type_traits(GGML_TYPE_F16)->from_float_ref(weights_f32.data(), weights.data(), n_elements);
-            auto result = std::make_shared<ov::op::v0::Constant>(weights);
-            result->set_friendly_name(tensor->name);
-            return result;
-        }
-        if (requant_type == ExtraQuantType::Q4_1_64) {
-            quantize_q4_1_asym(weights_f32.data(), weights, scales, zp, n_elements, block_size);
-        } else if (is_u4) {
-            quantize_q4_0(weights_f32.data(), weights, scales, zp, n_elements, block_size);
-        } else if (requant_type == ExtraQuantType::Q8_1_C) {
-            quantize_q8_1(weights_f32.data(), weights, scales, zp, n_elements, block_size);
-        } else {
-            quantize_q8_0(weights_f32.data(), weights, scales, zp, n_elements, block_size);
-        }
-    } else {
-        // Streaming path for Q8_0_C / Q8_1_C / F16 (covers token_embd, output.weight,
-        // and per-layer Q6_K/Q5_K requant — the large transient cases).
-        const int64_t CHUNK_ROWS = std::min<int64_t>(n_rows, 256);
-        std::vector<float> scratch(CHUNK_ROWS * ne0);
-        // F16 destination: 2 bytes/element, advanced per chunk by r0*ne0 elements.
-        auto * f16_base = static_cast<uint8_t *>(weights.data());
-        for (int64_t r0 = 0; r0 < n_rows; r0 += CHUNK_ROWS) {
-            const int64_t rows = std::min(CHUNK_ROWS, n_rows - r0);
-            const int64_t elems = rows * ne0;
-            const auto * src = static_cast<const uint8_t *>(data) + r0 * src_row_bytes;
-            type_traits->to_float(src, scratch.data(), elems);
-
-            if (requant_type == ExtraQuantType::F16) {
-                ggml_get_type_traits(GGML_TYPE_F16)
-                    ->from_float_ref(scratch.data(), f16_base + (r0 * ne0) * sizeof(uint16_t), elems);
-            } else {
-                const int64_t block_offset = (r0 * ne0) / block_size;
-                if (requant_type == ExtraQuantType::Q8_1_C) {
-                    quantize_q8_1(scratch.data(), weights, scales, zp, elems, block_size, block_offset);
-                } else {
-                    quantize_q8_0(scratch.data(), weights, scales, zp, elems, block_size, block_offset);
-                }
-            }
-        }
-        if (requant_type == ExtraQuantType::F16) {
-            auto result = std::make_shared<ov::op::v0::Constant>(weights);
-            result->set_friendly_name(tensor->name);
-            return result;
-        }
-    }
-
-    // Create the OpenVINO weight subgraph
-    ov::Output<ov::Node> weight_node;
-    if (is_u4) {
-        weight_node = make_int4_weights(weights, scales, zp, block_size);
-    } else {
-        weight_node = make_int8_weights(weights, scales, zp, block_size);
-    }
-
-    auto result = weight_node.get_node_shared_ptr();
-    result->set_friendly_name(tensor->name);
-    return result;
-}
-
-OvWeight process_weight_tensor(const ggml_tensor * tensor, const void * data, void * output_base_ptr, bool use_bias) {
-    GGML_ASSERT(tensor != nullptr);
-    GGML_ASSERT(data != nullptr);
-
-    OvWeight result;
-
-    // Get shape for weights: [rows, cols], or [n_expert, rows, cols] for 3D MoE expert weights.
-    ov::Shape node_shape = (tensor->ne[2] > 1) ?
-                               ov::Shape{static_cast<size_t>(tensor->ne[2]), static_cast<size_t>(tensor->ne[1]),
-                                         static_cast<size_t>(tensor->ne[0])} :
-                               ov::Shape{static_cast<size_t>(tensor->ne[1]), static_cast<size_t>(tensor->ne[0])};
-
-    // Handle F16/F32/BF16 weights
-    if (tensor->type == GGML_TYPE_F32 || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) {
-        ov::element::Type element_type;
-        switch (tensor->type) {
-        case GGML_TYPE_F32:
-            element_type = ov::element::f32;
-            break;
-        case GGML_TYPE_F16:
-            element_type = ov::element::f16;
-            break;
-        case GGML_TYPE_BF16:
-            element_type = ov::element::bf16;
-            break;
-        default:
-            OPENVINO_THROW("Unexpected tensor type in F16/F32/BF16 path");
-        }
-
-        if (output_base_ptr && output_base_ptr != data) {
-            // Using external buffer - copy data and create shared-memory constant
-            size_t tensor_bytes = ggml_nbytes(tensor);
-            memcpy(output_base_ptr, data, tensor_bytes);
-            result.weights = ov::Tensor(element_type, node_shape, output_base_ptr);
-        } else {
-            result.weights = ov::Tensor(element_type, node_shape, data);
-        }
-        result.weight_node = std::make_shared<ov::op::v0::Constant>(result.weights);
-        return result;
-    }
-
-    // Handle quantized weights
-    if (!ggml_is_quantized(tensor->type)) {
-        OPENVINO_THROW("Unsupported weight tensor type: ", ggml_type_name(tensor->type));
-    }
-
-    result.layout = ggml_openvino_get_extracted_layout(tensor, use_bias);
-    const auto & layout = result.layout;
-    if (layout.total_size == 0) {
-        OPENVINO_THROW("Unsupported quantized type: ", ggml_type_name(tensor->type));
-    }
-
-    // 3D MoE expert weights (for_gather_matmul) always use the exact f16 zero-point path (see
-    // extract_quantized_weights) -- must be kept in sync with the "use_bias || for_gather_matmul"
-    // check in ggml_openvino_get_extracted_layout, which sizes/offsets the zp slot accordingly.
-    // Requantized tensors (layout.is_requant) are handled by requantize_to_buffers instead, whose
-    // zp sizing/type is unaffected by for_gather_matmul, so they are excluded here.
-    const bool for_gather_matmul = tensor->ne[2] > 1;
-    const bool zp_is_f16 = !layout.is_requant && (use_bias || for_gather_matmul);
-
-    const bool is_3d_mxfp4_moe = tensor->type == GGML_TYPE_MXFP4 && (tensor->ne[2] > 1 || tensor->ne[3] > 1);
-    if (is_3d_mxfp4_moe) {
-        ov::Shape packed_shape = {static_cast<size_t>(tensor->ne[3]),
-                                  static_cast<size_t>(tensor->ne[2]),
-                                  static_cast<size_t>(tensor->ne[1]),
-                                  static_cast<size_t>(tensor->ne[0] / MXFP4_BLOCK_SIZE),
-                                  MXFP4_BLOCK_BYTES};
-        const size_t tensor_bytes = ggml_nbytes(tensor);
-        if (output_base_ptr) {
-            auto * buf_base = static_cast<uint8_t *>(output_base_ptr);
-            memcpy(buf_base + layout.weights_offset, data, tensor_bytes);
-            result.weights = ov::Tensor(ov::element::u8, packed_shape, buf_base + layout.weights_offset);
-        } else {
-            result.weights = ov::Tensor(ov::element::u8, packed_shape);
-            memcpy(result.weights.data(), data, tensor_bytes);
-        }
-        result.weight_node = make_mxfp4_moe_packed_weights(result.weights).get_node_shared_ptr();
-        result.weight_node->set_friendly_name(tensor->name);
-        return result;
-    }
-
-    if (use_bias) {
-        OPENVINO_ASSERT(!layout.is_requant,
-                        "use_bias is only used for test-backend-ops, which should not have requantization");
-        // bias node will be created on the fly and not use backend buffer
-        output_base_ptr = nullptr;
-    }
-
-    // F16 requant path - no separate scales/zp needed in result
-    if (layout.is_requant && layout.requant_type.has_value() && layout.requant_type.value() == ExtraQuantType::F16) {
-        if (output_base_ptr) {
-            result.weights = ov::Tensor(ov::element::f16, node_shape,
-                                        static_cast<uint8_t *>(output_base_ptr) + layout.weights_offset);
-        } else {
-            result.weights = ov::Tensor(ov::element::f16, node_shape);
-        }
-        ov::Tensor dummy_scales, dummy_zp;  // Not used for F16
-        result.weight_node =
-            requantize_to_buffers(tensor, data, ExtraQuantType::F16, 0, result.weights, dummy_scales, dummy_zp);
-        return result;
-    }
-
-    // Quantized path (normal extraction or quantized requant)
-    // Create weight/scale/zp tensors - shared between both paths
-    // For symmetric quantization, use signed types (i4/i8) and no ZP tensor
-    ov::element::Type weight_type = tensor->type == GGML_TYPE_MXFP4 ?
-                                        ov::element::f4e2m1 :
-                                        (layout.is_symmetric ? (layout.is_u4 ? ov::element::i4 : ov::element::i8) :
-                                                               (layout.is_u4 ? ov::element::u4 : ov::element::u8));
-    ov::Shape scale_shape = node_shape;
-    scale_shape.back() /= layout.weights_per_block;
-
-    if (tensor->type == GGML_TYPE_MXFP4) {
-        if (tensor->ne[2] == 1 && tensor->ne[3] == 1) {
-            node_shape = {static_cast<size_t>(tensor->ne[1]), static_cast<size_t>(tensor->ne[0])};
-        } else {
-            node_shape.clear();
-            for (int i = GGML_MAX_DIMS - 1; i >= 0; --i) {
-                node_shape.push_back(static_cast<size_t>(tensor->ne[i]));
-            }
-        }
-
-        scale_shape = node_shape;
-        scale_shape.back() /= layout.weights_per_block;
-    }
-
-    if (output_base_ptr) {
-        uint8_t * buf_base = static_cast<uint8_t *>(output_base_ptr);
-        result.weights = ov::Tensor(weight_type, node_shape, buf_base + layout.weights_offset);
-        const ov::element::Type scale_type = tensor->type == GGML_TYPE_MXFP4 ? ov::element::f8e8m0 : ov::element::f16;
-        result.scales = ov::Tensor(scale_type, scale_shape, buf_base + layout.scales_offset);
-        if (!layout.is_symmetric) {
-            ov::element::Type zp_type =
-                zp_is_f16 ? ov::element::f16 : (layout.is_u4 ? ov::element::u4 : ov::element::u8);
-            result.zp = ov::Tensor(zp_type, scale_shape, buf_base + layout.zp_offset);
-        }
-        // else: result.zp remains default-constructed (empty) for symmetric
-    } else {
-        result.weights = ov::Tensor(weight_type, node_shape);
-        const ov::element::Type scale_type = tensor->type == GGML_TYPE_MXFP4 ? ov::element::f8e8m0 : ov::element::f16;
-        result.scales = ov::Tensor(scale_type, scale_shape);
-        if (!layout.is_symmetric) {
-            if (zp_is_f16) {
-                result.zp = ov::Tensor(ov::element::f16, scale_shape);
-            } else {
-                ov::element::Type zp_type = layout.is_u4 ? ov::element::u4 : ov::element::u8;
-                result.zp = ov::Tensor(zp_type, scale_shape);
-            }
-        }
-        // else: result.zp remains default-constructed (empty) for symmetric
-    }
-
-    if (layout.is_requant && layout.requant_type.has_value()) {
-        result.weight_node = requantize_to_buffers(tensor, data, layout.requant_type.value(), layout.weights_per_block,
-                                                   result.weights, result.scales, result.zp);
-    } else {
-        result.weight_node =
-            extract_quantized_weights(tensor, data, result.weights, result.scales, result.zp, use_bias);
-    }
-
-    return result;
-}
-
 void quantize_q4_0(const float * x,
                   ov::Tensor & weights_arr,
                   ov::Tensor & scales_arr,
@@ -1252,7 +921,7 @@ void quantize_q8_0(const float * x,
                   ov::Tensor & zp_arr,
                   int64_t k,
                   int64_t qk,
-                   int64_t block_offset) {
+                   int64_t block_offset = 0) {
    assert(k % qk == 0);
    const int nb = k / qk;

@@ -1308,7 +977,7 @@ void quantize_q8_1(const float * x,
                   ov::Tensor & zp_arr,
                   int64_t k,
                   int64_t qk,
-                   int64_t block_offset) {
+                   int64_t block_offset = 0) {
    assert(k % qk == 0);
    const int nb = k / qk;

@@ -1339,3 +1008,366 @@ void quantize_q8_1(const float * x,
        }
    }
 }
+
+// Extract quantized weights from tensor and create weight subgraph
+// If weights/scales/zp are provided (non-empty), uses them as output buffers
+// Otherwise allocates new ov::Tensors internally
+// Returns the weight node (make_int4_weights or make_int8_weights result)
+std::shared_ptr<ov::Node> extract_quantized_weights(const ggml_tensor * tensor,
+                                                    const void * data,  // Source data pointer (may differ from tensor->data)
+                                                    ov::Tensor & weights,
+                                                    ov::Tensor & scales,
+                                                    ov::Tensor & zp,
+                                                    // Use an exact f16 zero point (vs. a rounded integer one); always
+                                                    // used for for_gather_matmul (3D MoE expert) weights regardless of
+                                                    // this flag, and also settable explicitly for test-backend-ops.
+                                                    bool use_bias = false) {
+    // Create a temporary tensor for extraction functions that read from tensor->data
+    ggml_tensor temp_tensor = *tensor;
+    temp_tensor.data = const_cast<void *>(data);
+
+    if (tensor->type == GGML_TYPE_MXFP4) {
+        extract_mxfp4_data(&temp_tensor, weights, scales);
+        auto result = make_mxfp4_weights(weights, scales).get_node_shared_ptr();
+        result->set_friendly_name(tensor->name);
+        return result;
+    }
+
+    // Determine block size based on tensor type
+    int64_t weights_per_block;
+    bool is_u4;
+    switch (tensor->type) {
+    case GGML_TYPE_Q4_0:
+    case GGML_TYPE_Q4_1:
+    case GGML_TYPE_Q4_K:
+        is_u4 = true;
+        weights_per_block = 32;
+        break;
+    case GGML_TYPE_Q8_0:
+    case GGML_TYPE_Q5_1:
+    case GGML_TYPE_Q5_K:
+        is_u4 = false;
+        weights_per_block = 32;
+        break;
+    case GGML_TYPE_Q6_K:
+        is_u4 = false;
+        weights_per_block = 16;
+        break;
+    default:
+        throw std::runtime_error("Unsupported quantized type for extraction: " +
+                                 std::string(ggml_type_name(tensor->type)));
+    }
+
+    // 3D MoE expert weights (for_gather_matmul) always use the exact f16 zero-point extraction
+    // (see make_int8_weights/make_int4_weights) rather than the rounded integer zero point --
+    // round(min/scale) error is what corrupts Q4_K/Q5_1 experts, and the f16-zp form still fuses
+    // into GatherMatmulCompressed since it stays a Subtract, not an Add.
+    const bool for_gather_matmul = tensor->ne[2] > 1;
+    use_bias = use_bias || for_gather_matmul;
+
+    // Extract quantized data
+    switch (tensor->type) {
+    case GGML_TYPE_Q4_0:
+        extract_q4_0_data(&temp_tensor, weights, scales, zp);
+        break;
+    case GGML_TYPE_Q4_1:
+        extract_q4_1_data(&temp_tensor, weights, scales, zp, use_bias);
+        break;
+    case GGML_TYPE_Q4_K:
+        extract_q4_k_data(&temp_tensor, weights, scales, zp, use_bias);
+        break;
+    case GGML_TYPE_Q5_1:
+        extract_q5_1_data(&temp_tensor, weights, scales, zp, use_bias);
+        break;
+    case GGML_TYPE_Q8_0:
+        extract_q8_0_data(&temp_tensor, weights, scales, zp);
+        break;
+    case GGML_TYPE_Q6_K:
+        extract_q6_k_data(&temp_tensor, weights, scales, zp);
+        break;
+    case GGML_TYPE_Q5_K:
+        extract_q5_k_data(&temp_tensor, weights, scales, zp, use_bias);
+        break;
+    default:
+        throw std::runtime_error("Unsupported quantized type: " + std::string(ggml_type_name(tensor->type)));
+    }
+
+    // Create the OpenVINO weight subgraph. 3D expert weights (MoE) are routed through the
+    // GatherMatmul-oriented path: dequantized in f16, with constant folding disabled on the chain.
+    ov::Output<ov::Node> weight_node;
+    if (is_u4) {
+        weight_node = make_int4_weights(weights, scales, zp, weights_per_block, use_bias, for_gather_matmul);
+    } else {
+        weight_node = make_int8_weights(weights, scales, zp, weights_per_block, use_bias, for_gather_matmul);
+    }
+
+    auto result = weight_node.get_node_shared_ptr();
+    result->set_friendly_name(tensor->name);
+    return result;
+}
+
+// Requantize weights from tensor to target format, writing to provided buffers
+// For F16 target, only weights buffer is used (scales/zp ignored)
+// Returns the weight node
+std::shared_ptr<ov::Node> requantize_to_buffers(const ggml_tensor * tensor,
+                                                const void * data,  // Source data pointer
+                                                ExtraQuantType requant_type,
+                                                int64_t block_size,
+                                                ov::Tensor & weights,
+                                                ov::Tensor & scales,
+                                                ov::Tensor & zp) {
+    int64_t n_elements = ggml_nelements(tensor);
+    const int64_t ne0 = tensor->ne[0];                 // elements per row
+    const int64_t n_rows = n_elements / ne0;
+    const auto * type_traits = ggml_get_type_traits(tensor->type);
+    const size_t src_row_bytes = ggml_row_size(tensor->type, ne0);
+
+    bool is_u4 = (requant_type == ExtraQuantType::Q4_0_C || requant_type == ExtraQuantType::Q4_0_128 ||
+                  requant_type == ExtraQuantType::Q4_0_64 || requant_type == ExtraQuantType::Q4_1_64);
+
+    // Streaming dequant (opt-in via GGML_OPENVINO_REDUCE_COMPILE_MEM or
+    // GGML_OPENVINO_MEMORY_OPTIMIZE): instead of
+    // materializing the full n_elements F32 array (e.g. ~1 GB for token_embd), dequantize
+    // a chunk of complete rows into a small scratch and quantize/convert it straight into
+    // the output buffers, capping the transient F32 footprint at CHUNK_ROWS*ne0 floats.
+    //
+    // Only valid (and only used) for the Q8_0_C / Q8_1_C / F16 targets whose block size
+    // divides a row (channel-wise _C uses block_size == ne0) so no target block straddles
+    // a row boundary, and Q8/F16 have no cross-block packing. The u4 (Q4_0) path packs two
+    // weights per byte with running zp ORs that assume a single whole-array call, so it is
+    // never streamed. When the flag is off, behavior is identical to the original
+    // full-materialization path.
+    const bool stream_requant = ggml_openvino_reduce_compile_mem_enabled() && !is_u4 &&
+                                !(block_size > 0 && ne0 % block_size != 0);
+
+    if (!stream_requant) {
+        // Full materialization (original behavior): dequantize the whole tensor to F32,
+        // then convert/quantize in one call.
+        std::vector<float> weights_f32(n_elements);
+        type_traits->to_float(data, weights_f32.data(), n_elements);
+        if (requant_type == ExtraQuantType::F16) {
+            ggml_get_type_traits(GGML_TYPE_F16)->from_float_ref(weights_f32.data(), weights.data(), n_elements);
+            auto result = std::make_shared<ov::op::v0::Constant>(weights);
+            result->set_friendly_name(tensor->name);
+            return result;
+        }
+        if (requant_type == ExtraQuantType::Q4_1_64) {
+            quantize_q4_1_asym(weights_f32.data(), weights, scales, zp, n_elements, block_size);
+        } else if (is_u4) {
+            quantize_q4_0(weights_f32.data(), weights, scales, zp, n_elements, block_size);
+        } else if (requant_type == ExtraQuantType::Q8_1_C) {
+            quantize_q8_1(weights_f32.data(), weights, scales, zp, n_elements, block_size);
+        } else {
+            quantize_q8_0(weights_f32.data(), weights, scales, zp, n_elements, block_size);
+        }
+    } else {
+        // Streaming path for Q8_0_C / Q8_1_C / F16 (covers token_embd, output.weight,
+        // and per-layer Q6_K/Q5_K requant — the large transient cases).
+        const int64_t CHUNK_ROWS = std::min<int64_t>(n_rows, 256);
+        std::vector<float> scratch(CHUNK_ROWS * ne0);
+        // F16 destination: 2 bytes/element, advanced per chunk by r0*ne0 elements.
+        auto * f16_base = static_cast<uint8_t *>(weights.data());
+        for (int64_t r0 = 0; r0 < n_rows; r0 += CHUNK_ROWS) {
+            const int64_t rows = std::min(CHUNK_ROWS, n_rows - r0);
+            const int64_t elems = rows * ne0;
+            const auto * src = static_cast<const uint8_t *>(data) + r0 * src_row_bytes;
+            type_traits->to_float(src, scratch.data(), elems);
+
+            if (requant_type == ExtraQuantType::F16) {
+                ggml_get_type_traits(GGML_TYPE_F16)
+                    ->from_float_ref(scratch.data(), f16_base + (r0 * ne0) * sizeof(uint16_t), elems);
+            } else {
+                const int64_t block_offset = (r0 * ne0) / block_size;
+                if (requant_type == ExtraQuantType::Q8_1_C) {
+                    quantize_q8_1(scratch.data(), weights, scales, zp, elems, block_size, block_offset);
+                } else {
+                    quantize_q8_0(scratch.data(), weights, scales, zp, elems, block_size, block_offset);
+                }
+            }
+        }
+        if (requant_type == ExtraQuantType::F16) {
+            auto result = std::make_shared<ov::op::v0::Constant>(weights);
+            result->set_friendly_name(tensor->name);
+            return result;
+        }
+    }
+
+    // Create the OpenVINO weight subgraph
+    ov::Output<ov::Node> weight_node;
+    if (is_u4) {
+        weight_node = make_int4_weights(weights, scales, zp, block_size);
+    } else {
+        weight_node = make_int8_weights(weights, scales, zp, block_size);
+    }
+
+    auto result = weight_node.get_node_shared_ptr();
+    result->set_friendly_name(tensor->name);
+    return result;
+}
+}  // namespace
+
+OvWeight process_weight_tensor(const ggml_tensor * tensor, const void * data, void * output_base_ptr, bool use_bias) {
+    GGML_ASSERT(tensor != nullptr);
+    GGML_ASSERT(data != nullptr);
+
+    OvWeight result;
+
+    // Get shape for weights: [rows, cols], or [n_expert, rows, cols] for 3D MoE expert weights.
+    ov::Shape node_shape = (tensor->ne[2] > 1) ?
+                               ov::Shape{static_cast<size_t>(tensor->ne[2]), static_cast<size_t>(tensor->ne[1]),
+                                         static_cast<size_t>(tensor->ne[0])} :
+                               ov::Shape{static_cast<size_t>(tensor->ne[1]), static_cast<size_t>(tensor->ne[0])};
+
+    // Handle F16/F32/BF16 weights
+    if (tensor->type == GGML_TYPE_F32 || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) {
+        ov::element::Type element_type;
+        switch (tensor->type) {
+        case GGML_TYPE_F32:
+            element_type = ov::element::f32;
+            break;
+        case GGML_TYPE_F16:
+            element_type = ov::element::f16;
+            break;
+        case GGML_TYPE_BF16:
+            element_type = ov::element::bf16;
+            break;
+        default:
+            OPENVINO_THROW("Unexpected tensor type in F16/F32/BF16 path");
+        }
+
+        if (output_base_ptr && output_base_ptr != data) {
+            // Using external buffer - copy data and create shared-memory constant
+            size_t tensor_bytes = ggml_nbytes(tensor);
+            memcpy(output_base_ptr, data, tensor_bytes);
+            result.weights = ov::Tensor(element_type, node_shape, output_base_ptr);
+        } else {
+            result.weights = ov::Tensor(element_type, node_shape, data);
+        }
+        result.weight_node = std::make_shared<ov::op::v0::Constant>(result.weights);
+        return result;
+    }
+
+    // Handle quantized weights
+    if (!ggml_is_quantized(tensor->type)) {
+        OPENVINO_THROW("Unsupported weight tensor type: ", ggml_type_name(tensor->type));
+    }
+
+    result.layout = ggml_openvino_get_extracted_layout(tensor, use_bias);
+    const auto & layout = result.layout;
+    if (layout.total_size == 0) {
+        OPENVINO_THROW("Unsupported quantized type: ", ggml_type_name(tensor->type));
+    }
+
+    // 3D MoE expert weights (for_gather_matmul) always use the exact f16 zero-point path (see
+    // extract_quantized_weights) -- must be kept in sync with the "use_bias || for_gather_matmul"
+    // check in ggml_openvino_get_extracted_layout, which sizes/offsets the zp slot accordingly.
+    // Requantized tensors (layout.is_requant) are handled by requantize_to_buffers instead, whose
+    // zp sizing/type is unaffected by for_gather_matmul, so they are excluded here.
+    const bool for_gather_matmul = tensor->ne[2] > 1;
+    const bool zp_is_f16 = !layout.is_requant && (use_bias || for_gather_matmul);
+
+    const bool is_3d_mxfp4_moe = tensor->type == GGML_TYPE_MXFP4 && (tensor->ne[2] > 1 || tensor->ne[3] > 1);
+    if (is_3d_mxfp4_moe) {
+        ov::Shape packed_shape = {static_cast<size_t>(tensor->ne[3]),
+                                  static_cast<size_t>(tensor->ne[2]),
+                                  static_cast<size_t>(tensor->ne[1]),
+                                  static_cast<size_t>(tensor->ne[0] / MXFP4_BLOCK_SIZE),
+                                  MXFP4_BLOCK_BYTES};
+        const size_t tensor_bytes = ggml_nbytes(tensor);
+        if (output_base_ptr) {
+            auto * buf_base = static_cast<uint8_t *>(output_base_ptr);
+            memcpy(buf_base + layout.weights_offset, data, tensor_bytes);
+            result.weights = ov::Tensor(ov::element::u8, packed_shape, buf_base + layout.weights_offset);
+        } else {
+            result.weights = ov::Tensor(ov::element::u8, packed_shape);
+            memcpy(result.weights.data(), data, tensor_bytes);
+        }
+        result.weight_node = make_mxfp4_moe_packed_weights(result.weights).get_node_shared_ptr();
+        result.weight_node->set_friendly_name(tensor->name);
+        return result;
+    }
+
+    if (use_bias) {
+        OPENVINO_ASSERT(!layout.is_requant,
+                        "use_bias is only used for test-backend-ops, which should not have requantization");
+        // bias node will be created on the fly and not use backend buffer
+        output_base_ptr = nullptr;
+    }
+
+    // F16 requant path - no separate scales/zp needed in result
+    if (layout.is_requant && layout.requant_type.has_value() && layout.requant_type.value() == ExtraQuantType::F16) {
+        if (output_base_ptr) {
+            result.weights = ov::Tensor(ov::element::f16, node_shape,
+                                        static_cast<uint8_t *>(output_base_ptr) + layout.weights_offset);
+        } else {
+            result.weights = ov::Tensor(ov::element::f16, node_shape);
+        }
+        // Not used for F16:
+        ov::Tensor dummy_scales;
+        ov::Tensor dummy_zp;
+        result.weight_node =
+            requantize_to_buffers(tensor, data, ExtraQuantType::F16, 0, result.weights, dummy_scales, dummy_zp);
+        return result;
+    }
+
+    // Quantized path (normal extraction or quantized requant)
+    // Create weight/scale/zp tensors - shared between both paths
+    // For symmetric quantization, use signed types (i4/i8) and no ZP tensor
+    ov::element::Type weight_type;
+    if (tensor->type == GGML_TYPE_MXFP4) {
+        weight_type = ov::element::f4e2m1;
+    } else if (layout.is_symmetric) {
+        weight_type = layout.is_u4 ? ov::element::i4 : ov::element::i8;
+    } else {
+        weight_type = layout.is_u4 ? ov::element::u4 : ov::element::u8;
+    }
+    ov::Shape scale_shape = node_shape;
+    scale_shape.back() /= layout.weights_per_block;
+
+    if (tensor->type == GGML_TYPE_MXFP4) {
+        if (tensor->ne[2] == 1 && tensor->ne[3] == 1) {
+            node_shape = {static_cast<size_t>(tensor->ne[1]), static_cast<size_t>(tensor->ne[0])};
+        } else {
+            node_shape.clear();
+            for (int i = GGML_MAX_DIMS - 1; i >= 0; --i) {
+                node_shape.push_back(static_cast<size_t>(tensor->ne[i]));
+            }
+        }
+
+        scale_shape = node_shape;
+        scale_shape.back() /= layout.weights_per_block;
+    }
+
+    const ov::element::Type scale_type = tensor->type == GGML_TYPE_MXFP4 ? ov::element::f8e8m0 : ov::element::f16;
+    ov::element::Type zp_type = layout.is_u4 ? ov::element::u4 : ov::element::u8;
+    if (zp_is_f16) {
+        zp_type = ov::element::f16;
+    }
+
+    if (output_base_ptr) {
+        uint8_t * buf_base = static_cast<uint8_t *>(output_base_ptr);
+        result.weights = ov::Tensor(weight_type, node_shape, buf_base + layout.weights_offset);
+        result.scales = ov::Tensor(scale_type, scale_shape, buf_base + layout.scales_offset);
+        if (!layout.is_symmetric) {
+            result.zp = ov::Tensor(zp_type, scale_shape, buf_base + layout.zp_offset);
+        }
+        // else: result.zp remains default-constructed (empty) for symmetric
+    } else {
+        result.weights = ov::Tensor(weight_type, node_shape);
+        result.scales = ov::Tensor(scale_type, scale_shape);
+        if (!layout.is_symmetric) {
+            result.zp = ov::Tensor(zp_type, scale_shape);
+        }
+        // else: result.zp remains default-constructed (empty) for symmetric
+    }
+
+    if (layout.is_requant && layout.requant_type.has_value()) {
+        result.weight_node = requantize_to_buffers(tensor, data, layout.requant_type.value(), layout.weights_per_block,
+                                                   result.weights, result.scales, result.zp);
+    } else {
+        result.weight_node =
+            extract_quantized_weights(tensor, data, result.weights, result.scales, result.zp, use_bias);
+    }
+
+    return result;
+}
@@ -2,112 +2,12 @@
 #include "ggml-openvino-extra.h"  // For ExtraQuantType
 #include "ggml.h"

-#include <cstdint>
-#include <openvino/op/constant.hpp>
 #include <openvino/core/node_output.hpp>
+#include <openvino/op/constant.hpp>
 #include <openvino/runtime/tensor.hpp>

-void unpack_32_4(const uint8_t * data, uint8_t * dst);
-
-void extract_q4_0_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr);
-
-void extract_q4_1_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr,
-                       bool use_bias = false);
-
-void extract_q5_1_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr,
-                       bool use_bias = false);
-
-void extract_q8_0_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr);
-
-void unpack_256_4(const uint8_t * data, uint8_t * dst);
-
-void extract_q4_k_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr,
-                       bool use_bias = false);
-
-void extract_q5_k_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr,
-                       bool use_bias = false);
-
-void extract_q6_k_data(const ggml_tensor * tensor,
-                       ov::Tensor & weights_arr,
-                       ov::Tensor & scales_arr,
-                       ov::Tensor & zp_arr);
-
-void extract_mxfp4_data(const ggml_tensor * tensor, ov::Tensor & weights_arr, ov::Tensor & scales_arr);
-
 static constexpr size_t GGML_QUANTIZATION_GROUP_SIZE = 32;

-// If for_gather_matmul is true, the weight tensor may be N-D (e.g. 3D MoE expert weights
-// [n_expert, rows, cols]). The dequantization chain (Convert->[Subtract]->Multiply) is built as
-// usual but left in f16 (no final Convert to f32) -- ov::pass::MarkDequantization (registered in
-// translate_session.cpp) marks the chain so it survives model-build-time ConstantFolding -- see
-// make_int8_weights.cpp/make_int4_weights.cpp. mul_mat_id.cpp constructs ov::op::internal::GatherMatmul
-// directly from the resulting f16 dequant chain.
-//
-// When use_bias is true (explicitly, or implicitly because for_gather_matmul is true), the zp
-// tensor is expected to hold an exact f16 bias value (rather than a rounded integer zero point);
-// it is converted in place into an exact zero_point = -bias/scale and consumed via Subtract, not
-// Add, so the chain still matches OpenVINO's Convert->Subtract->Multiply decompression pattern.
-ov::Output<ov::Node> make_int8_weights(ov::Tensor & weight,
-                                       ov::Tensor & scales,
-                                       ov::Tensor & zp,
-                                       size_t group_size = GGML_QUANTIZATION_GROUP_SIZE,
-                                       bool use_bias = false,
-                                       bool for_gather_matmul = false);
-
-ov::Output<ov::Node> make_int4_weights(ov::Tensor & weight,
-                                       ov::Tensor & scales,
-                                       ov::Tensor & zp,
-                                       size_t group_size = GGML_QUANTIZATION_GROUP_SIZE,
-                                       bool use_bias = false,
-                                       bool for_gather_matmul = false);
-
-ov::Output<ov::Node> make_mxfp4_weights(ov::Tensor & weight, ov::Tensor & scales);
-
-ov::Output<ov::Node> make_mxfp4_moe_packed_weights(ov::Tensor & weight);
-
-// Extract quantized weights from tensor and create weight subgraph
-// If weights/scales/zp are provided (non-empty), uses them as output buffers
-// Otherwise allocates new ov::Tensors internally
-// Returns the weight node (make_int4_weights or make_int8_weights result)
-std::shared_ptr<ov::Node> extract_quantized_weights(
-    const ggml_tensor * tensor,
-    const void * data,  // Source data pointer (may differ from tensor->data)
-    ov::Tensor & weights,
-    ov::Tensor & scales,
-    ov::Tensor & zp,
-    bool use_bias = false);  // Use an exact f16 zero point (vs. a rounded integer one); always
-                             // used for for_gather_matmul (3D MoE expert) weights regardless of
-                             // this flag, and also settable explicitly for test-backend-ops.
-
-// Requantize weights from tensor to target format, writing to provided buffers
-// For F16 target, only weights buffer is used (scales/zp ignored)
-// Returns the weight node
-std::shared_ptr<ov::Node> requantize_to_buffers(const ggml_tensor * tensor,
-                                                const void * data,  // Source data pointer
-                                                ExtraQuantType requant_type,
-                                                int64_t block_size,
-                                                ov::Tensor & weights,
-                                                ov::Tensor & scales,
-                                                ov::Tensor & zp);
-
 inline const char * extra_quant_type_name(ExtraQuantType t) {
    switch (t) {
    case ExtraQuantType::F16:
@@ -156,41 +56,3 @@ OvWeight process_weight_tensor(
                                       // always used for for_gather_matmul (3D MoE expert) weights
                                       // regardless of this flag, and also settable explicitly for
                                       // test-backend-ops.
-
-void quantize_q4_0(const float * x,
-                   ov::Tensor & weights_arr,
-                   ov::Tensor & scales_arr,
-                   ov::Tensor & zp_arr,
-                   int64_t k,
-                   int64_t qk);
-void quantize_q8_1(const float * x,
-                   ov::Tensor & weights_arr,
-                   ov::Tensor & scales_arr,
-                   ov::Tensor & zp_arr,
-                   int64_t k,
-                   int64_t qk,
-                   int64_t block_offset = 0);
-void quantize_q4_1_asym(const float * x,
-                        ov::Tensor & weights_arr,
-                        ov::Tensor & scales_arr,
-                        ov::Tensor & zp_arr,
-                        int64_t k,
-                        int64_t qk);
-void quantize_q8_0(const float * x,
-                   ov::Tensor & weights_arr,
-                   ov::Tensor & scales_arr,
-                   ov::Tensor & zp_arr,
-                   int64_t k,
-                   int64_t qk,
-                   int64_t block_offset = 0);
-
-namespace ov {
-namespace op {
-namespace util {
-// From <openvino>/src/common/transformations/include/transformations/utils/utils.hpp
-bool get_single_value(const std::shared_ptr<ov::op::v0::Constant> & const_node,
-                      float & value,
-                      bool check_value_range = true);
-}  // namespace util
-}  // namespace op
-}  // namespace ov
@@ -237,7 +237,8 @@ bool ggml_openvino_model_cache_verify_manifest(const std::string & path,
    if (!f.is_open()) {
        return false;
    }
-    std::string tag, val;
+    std::string tag;
+    std::string val;
    // header: fingerprint
    if (!(f >> tag >> val) || tag != "fingerprint" || val != hex64(fingerprint)) {
        return false;
@@ -12,7 +12,6 @@ namespace ggml {

 class FrontEnd {
 public:
-    using Ptr = std::shared_ptr<FrontEnd>;
    FrontEnd();

    static std::shared_ptr<Model> convert(const InputModel::Ptr & model, bool naive = false);
@@ -20,7 +20,7 @@ namespace op {
 static ov::Output<ov::Node> reshape_add_id_input_to_2d(const ov::Output<ov::Node> & input,
                                                       const ov::PartialShape & input_shape,
                                                       const std::vector<int> & dims) {
-    const auto actual_shape = input.get_partial_shape();
+    const auto & actual_shape = input.get_partial_shape();
    if (actual_shape.rank().is_static() && actual_shape.rank().get_length() == 2) {
        return input;
    }
@@ -3,12 +3,9 @@
 #include "../op_table.h"
 #include "../utils.h"

-#include <climits>
-#include <cstdint>
 #include <memory>
 #include <openvino/op/reshape.hpp>
 #include <openvino/op/slice.hpp>
-#include <vector>

 namespace ov {
 namespace frontend {
@@ -195,7 +195,9 @@ OutputVector translate_flash_attn_ext(const NodeContext & context) {
    auto tile_kv = [&](int64_t n_heads, int64_t n_heads_kv, int64_t hs, ov::Output<Node> kv) {
        int64_t f = n_heads / n_heads_kv;
        if (f > 1 && n_heads_kv > 1) {
-            ov::Output<ov::Node> kv_broadcast_shape, kv_unsqueezed, new_kv_shape;
+            ov::Output<ov::Node> kv_broadcast_shape;
+            ov::Output<ov::Node> kv_unsqueezed;
+            ov::Output<ov::Node> new_kv_shape;
            auto unsqueeze_axes = ov::op::v0::Constant::create(ov::element::i64, Shape{}, {2});
            kv_unsqueezed = std::make_shared<ov::op::v0::Unsqueeze>(kv, unsqueeze_axes);

@@ -196,7 +196,7 @@ static OutputVector translate_gated_delta_net_ref(const NodeContext & context) {
    }

    // Merge batch and head dims: [B*H_v, T, S_v]
-    auto merge_bh = [&](ov::Output<ov::Node> x, int64_t last_dim) {
+    auto merge_bh = [&](const ov::Output<ov::Node> & x, int64_t last_dim) {
        auto shape = ov::op::v0::Constant::create(ov::element::i64, {3}, std::vector<int64_t>{B * H_v, T, last_dim});
        return std::make_shared<ov::op::v1::Reshape>(x, shape, false);
    };
@@ -1,7 +1,6 @@
 #include "../node_context.h"
 #include "../op_table.h"
 #include "../utils.h"
-#include "ggml-impl.h"

 #include <cstddef>
 #include <memory>
@@ -42,7 +42,7 @@ ov::Output<ov::Node> slice_axis(const ov::Output<ov::Node> & input, int64_t axis

 ov::Output<ov::Node> static_shape_dims_or_shapeof(const ov::Output<ov::Node> & input,
                                                  const std::vector<int> & dims) {
-    const auto partial_shape = input.get_partial_shape();
+    const auto & partial_shape = input.get_partial_shape();
    if (partial_shape.is_static()) {
        std::vector<int64_t> values;
        values.reserve(dims.size());
@@ -8,6 +8,7 @@
 #include <openvino/op/pad.hpp>
 #include <openvino/op/reshape.hpp>
 #include <openvino/op/shape_of.hpp>
+#include <utility>
 #include <vector>

 namespace ov {
@@ -20,7 +21,7 @@ namespace {
 ov::Output<ov::Node> translate_circular_pad(ov::Output<ov::Node> input,
                                            const std::array<int32_t, 8> & pads,
                                            const ov::Shape & input_shape) {
-    ov::Output<ov::Node> result = input;
+    ov::Output<ov::Node> result = std::move(input);

    const std::array<int32_t, 4> pads_begin = {pads[6], pads[4], pads[2], pads[0]};
    const std::array<int32_t, 4> pads_end = {pads[7], pads[5], pads[3], pads[1]};
@@ -1,7 +1,6 @@
 #include "../node_context.h"
 #include "../op_table.h"
 #include "../utils.h"
-#include "ggml.h"

 #include <memory>
 #include <openvino/op/broadcast.hpp>
@@ -25,9 +25,7 @@ OutputVector translate_rms_norm(const NodeContext & context) {
    auto op_case = context.get_op_case();

    ov::Output<ov::Node> input_node;
-    if (op_case == 1) {
-        input_node = process_view_input_new(context, 0);
-    } else if (op_case == 2) {
+    if (op_case == 2) {
        auto ssm_state_size = context.get_ssm_state_size();
        // The GDN op packs [attn | new_state] along the row axis; the state occupies the last
        // ssm_state_size * n_seqs rows. Slice it off (scaling by the active sequence count) to keep
@@ -7,7 +7,6 @@
 #include <openvino/op/reshape.hpp>
 #include <openvino/op/shape_of.hpp>
 #include <openvino/op/slice.hpp>
-#include <set>

 namespace ov {
 namespace frontend {
@@ -153,7 +152,8 @@ OutputVector translate_view(const NodeContext & context) {
        return {input};
    }

-    int64_t src_elems = 1, dst_elems = 1;
+    int64_t src_elems = 1;
+    int64_t dst_elems = 1;
    for (int64_t i = 0; i < src_shape.rank().get_length(); ++i) {
        if (src_shape[i].is_dynamic()) {
            return {input};
@@ -84,7 +84,7 @@ bool KVStateSeqAxis::run_on_model(const std::shared_ptr<ov::Model> & model) {
        // Readers still expect seq at dim 1. A reader that is itself the inverse
        // Transpose wanted seq at dim 2 all along, so drop it; give anything else the
        // inverse Transpose so its input is unchanged.
-        for (auto & reader : readers) {
+        for (const auto & reader : readers) {
            auto * node = reader.get_node();
            if (ov::is_type<ov::op::v6::Assign>(node)) {
                continue;
@@ -344,7 +344,7 @@ std::shared_ptr<Model> TranslateSession::translate_graph(const frontend::InputMo
        }
    };

-    auto node_visitor = [&](std::shared_ptr<GgmlDecoder> decoder, int node_idx) {
+    auto node_visitor = [&](const std::shared_ptr<GgmlDecoder> & decoder, int node_idx) {
        auto converted_outputs = translate_node(decoder, node_idx);
        if (converted_outputs.empty()) {
            return;
@@ -1,7 +1,5 @@
 #include "utils.h"

-#include "ggml-impl.h"
-
 #include <cmath>
 #include <cstddef>
 #include <ctime>
@@ -28,13 +26,6 @@ namespace ov {
 namespace frontend {
 namespace ggml {

-std::string getCurrentTime() {
-    std::time_t now = std::time(nullptr);
-    char buf[100];
-    std::strftime(buf, sizeof(buf), "%Y-%m-%d %H:%M:%S", std::localtime(&now));
-    return buf;
-}
-
 void num_inputs_check(const NodeContext & context, size_t min_inputs, size_t max_inputs) {
    auto input_size = context.get_input_size();
    FRONT_END_OP_CONVERSION_CHECK(input_size >= min_inputs, "Got less inputs than expected");
@@ -82,7 +73,7 @@ namespace {
 ov::Output<ov::Node> rope_yarn_ramp_mix(int n_dims, const float corr_dims[2], float ext_factor) {
    int half_n_dims = n_dims / 2;
    std::vector<float> dim_ids_vec(half_n_dims);
-    std::iota(dim_ids_vec.begin(), dim_ids_vec.end(), 0);
+    std::iota(dim_ids_vec.begin(), dim_ids_vec.end(), 0.0f);
    auto dim_ids = ov::op::v0::Constant::create(ov::element::f32, Shape{1, 1, 1, (size_t) half_n_dims}, dim_ids_vec);
    auto corr_low = ov::op::v0::Constant::create(ov::element::f32, Shape{1, 1, 1, 1}, {corr_dims[0]});
    auto corr_high = ov::op::v0::Constant::create(ov::element::f32, Shape{1, 1, 1, 1}, {corr_dims[1]});
@@ -551,6 +542,7 @@ ov::Output<ov::Node> process_view_input_new(const NodeContext & context, int inp

                if (tail_begin >= 0 && tail_end <= tail_src_elems) {
                    std::vector<int64_t> flat_shape;
+                    flat_shape.reserve(slice_dim);
                    for (int i = 0; i < slice_dim; ++i) {
                        flat_shape.push_back(static_cast<int64_t>(view_src_ggml_shape[i]));
                    }
@@ -14,8 +14,6 @@ namespace ggml {

 std::string getCurrentTime();

-void dump_ov_model(std::shared_ptr<ov::Model> model);
-
 void num_inputs_check(const NodeContext & context, size_t min_inputs, size_t max_inputs);

 int non_cont_dim(std::vector<size_t> ne, std::vector<size_t> nb);
@@ -142,9 +142,6 @@ struct ov_runtime_context {

 enum ggml_status ov_graph_compute(struct ggml_cgraph * cgraph, ggml_backend_t backend);

-enum ggml_status ov_graph_compute_dynamic(struct ggml_cgraph * cgraph, std::shared_ptr<ov_runtime_context> r_ctx);
-enum ggml_status ov_graph_compute_static(struct ggml_cgraph * cgraph, std::shared_ptr<ov_runtime_context> r_ctx);
-
 size_t checksum(const void * data, size_t size);

 bool save_ggml_tensor_data_to_txt(const ggml_tensor * tensor, const std::string & file_path);
@@ -185,18 +182,6 @@ int64_t get_inp_pos_n_tokens(struct ggml_cgraph * cgraph, const ggml_tensor * in

 bool get_is_prefill(struct ggml_cgraph * cgraph, const ggml_tensor * inp_pos);

-ov::Tensor get_ov_input_tensor(std::shared_ptr<GgmlOvDecoder> ggml_decoder, const std::string & param_name);
-ov::Tensor get_ov_input_tensor_static_decode(std::shared_ptr<GgmlOvDecoder> ggml_decoder,
-                                             const std::string & param_name);
-ov::Tensor get_ov_input_tensor_static_prefill(std::shared_ptr<GgmlOvDecoder> ggml_decoder,
-                                              const std::string & param_name,
-                                              int chunk_index);
-
-ov::Tensor create_ov_output_tensor(std::shared_ptr<GgmlOvDecoder> ggml_decoder,
-                                   std::shared_ptr<ov::InferRequest> infer_request,
-                                   int output_index,
-                                   const ggml_tensor * ggml_tensor);
-
 bool is_naive(struct ggml_cgraph * cgraph);

 /**
@@ -205,9 +190,3 @@ bool is_naive(struct ggml_cgraph * cgraph);
 * @return true if the graph is identified as split; otherwise false.
 */
 bool is_model_splitted(struct ggml_cgraph * cgraph);
-
-enum ggml_status naive_compute(struct ggml_cgraph * cgraph,
-                               ov::Core & core,
-                               const std::string & device,
-                               const ov::AnyMap & config,
-                               ov_compiled_model_cache & cache);
@@ -208,5 +208,53 @@ bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializ
        return true;
    }

+    if (ops.size() == 2 && ops.begin()[0] == GGML_OP_SSM_CONV && ops.begin()[1] == GGML_OP_UNARY &&
+        unary_ops.size() == 1 && unary_ops.begin()[0] == GGML_UNARY_OP_SILU) {
+        const ggml_tensor * ssm_conv = cgraph->nodes[node_idx];
+        const ggml_tensor * silu     = cgraph->nodes[node_idx + 1];
+
+        if (ggml_get_unary_op(silu) != unary_ops.begin()[0]) {
+            return false;
+        }
+        if (ssm_conv->type != GGML_TYPE_F32 || silu->type != GGML_TYPE_F32) {
+            return false;
+        }
+        // the fused kernel writes the SiLU output with dense strides, so it must be contiguous
+        if (!ggml_is_contiguous(silu)) {
+            return false;
+        }
+
+        return true;
+    }
+
+    if (ops.size() == 3 && ops.begin()[0] == GGML_OP_SSM_CONV && ops.begin()[1] == GGML_OP_ADD &&
+        ops.begin()[2] == GGML_OP_UNARY && unary_ops.size() == 1 && unary_ops.begin()[0] == GGML_UNARY_OP_SILU) {
+        const ggml_tensor * ssm_conv = cgraph->nodes[node_idx];
+        const ggml_tensor * add      = cgraph->nodes[node_idx + 1];
+        const ggml_tensor * silu     = cgraph->nodes[node_idx + 2];
+
+        if (ggml_get_unary_op(silu) != unary_ops.begin()[0]) {
+            return false;
+        }
+        if (ssm_conv->type != GGML_TYPE_F32 || add->type != GGML_TYPE_F32 || silu->type != GGML_TYPE_F32) {
+            return false;
+        }
+        // the fused kernel writes the SiLU output with dense strides, so it must be contiguous
+        if (!ggml_is_contiguous(silu)) {
+            return false;
+        }
+
+        // ADD must consume ssm_conv's output and broadcast a 1-D channel-wise bias
+        const ggml_tensor * bias = (add->src[0] == ssm_conv) ? add->src[1] : add->src[0];
+        if (bias->type != GGML_TYPE_F32 || !ggml_is_contiguous(bias)) {
+            return false;
+        }
+        if (ggml_nelements(bias) != ssm_conv->ne[0] || bias->ne[0] != ssm_conv->ne[0]) {
+            return false;
+        }
+
+        return true;
+    }
+
    return false;
 }
@@ -1011,12 +1011,25 @@ static size_t ggml_backend_sycl_buffer_type_get_alignment(ggml_backend_buffer_ty
    GGML_UNUSED(buft);
 }

+bool is_bmg_g31_arch(int device) {
+    return ggml_sycl_info().devices[device].hw_info.arch == gpu_arch::intel_gpu_bmg_g31;
+}
+
 static size_t ggml_backend_sycl_buffer_type_get_max_size(ggml_backend_buffer_type_t buft) {
    size_t max_alloc_size = dpct::get_current_device().get_max_mem_alloc_size();
    if (g_ggml_sycl_host_pinned_mem_2g) {
        return std::min(max_alloc_size, (size_t) 2LL*1024*1024*1024);
    } else {
+        ggml_backend_sycl_buffer_type_context * ctx = (ggml_backend_sycl_buffer_type_context *)buft->context;
+        int device = ctx->device;
+        if(is_bmg_g31_arch(device)) {
+            //Todo, it's workaround for BMG-G31, which has a known issue with large allocations.
+            //The max alloc size is reduced to 60% of the reported max alloc size.
+            //remove it after https://github.com/intel/compute-runtime/issues/998 is fixed.
+            max_alloc_size = max_alloc_size*0.6;
+        }
        return max_alloc_size;
+
    }
    GGML_UNUSED(buft);
 }
@@ -1484,7 +1497,7 @@ static ggml_backend_buffer_type_i ggml_backend_sycl_split_buffer_type_interface
    /* .is_host          = */ ggml_backend_sycl_split_buffer_type_is_host,
 };

-ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split) {
+ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type([[maybe_unused]] int main_device, const float * tensor_split) {
    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n");

    static std::mutex mutex;
@@ -3243,7 +3256,7 @@ inline void ggml_sycl_op_scale(ggml_backend_sycl_context & ctx, ggml_tensor * ds
    SYCL_CHECK(0);
 }

-static void ggml_sycl_set_peer_access(const int n_tokens, int main_device) {
+static void ggml_sycl_set_peer_access(const int n_tokens, [[maybe_unused]] int main_device) {
    static bool peer_access_enabled = false;

    const bool enable_peer_access = n_tokens <= GGML_SYCL_PEER_MAX_BATCH_SIZE;
@@ -6034,6 +6047,20 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
            }
        }

+        if (node->op == GGML_OP_SSM_CONV &&
+            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_SSM_CONV, GGML_OP_ADD, GGML_OP_UNARY }, { GGML_UNARY_OP_SILU })) {
+            ggml_sycl_ssm_conv_fused(*sycl_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2]);
+            i += 2;
+            continue;
+        }
+
+        if (node->op == GGML_OP_SSM_CONV &&
+            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_SSM_CONV, GGML_OP_UNARY }, { GGML_UNARY_OP_SILU })) {
+            ggml_sycl_ssm_conv_fused(*sycl_ctx, node, nullptr, cgraph->nodes[i + 1]);
+            i++;
+            continue;
+        }
+
        if (node->op == GGML_OP_MUL_MAT && ggml_sycl_mul_mat_glu_mmvq_fused(*sycl_ctx, cgraph, i)) {
            i += 2;
            continue;
@@ -1,11 +1,71 @@
 #include "ssm_conv.hpp"
 #include "common.hpp"
+#include "element_wise.hpp"

 #include <cstdio>

 using namespace sycl;

-static void kernel_ssm_conv(
+// One output element of the conv. DC is d_conv as a compile-time constant (0 keeps the
+// runtime loop); unfused callers pass literal false/nullptr so the epilogue folds away.
+template <int DC>
+static __dpct_inline__ void ssm_conv_element(
+    size_t idx,
+    const float *src_data,
+    const float *weights,
+    float *dst_data,
+    int d_conv,
+    int d_inner,
+    int n_t,
+    int src_stride_inner,
+    int src_stride_seq,
+    int dst_stride_token,
+    int dst_stride_seq,
+    bool apply_silu,
+    const float *bias
+) {
+    // src is token-contiguous per channel, dst is channel-contiguous per token,
+    // so indexing token-fastest coalesces the d_conv loads.
+    const int token   = static_cast<int>(idx % n_t);
+    const int channel = static_cast<int>((idx / n_t) % d_inner);
+    const int seq     = static_cast<int>(idx / (static_cast<size_t>(n_t) * static_cast<size_t>(d_inner)));
+
+    const float *s = src_data
+        + static_cast<size_t>(seq) * static_cast<size_t>(src_stride_seq)
+        + static_cast<size_t>(channel) * static_cast<size_t>(src_stride_inner)
+        + static_cast<size_t>(token);
+
+    const float *c = weights + static_cast<size_t>(channel) * static_cast<size_t>(d_conv);
+
+    float sumf = 0.0f;
+    if constexpr (DC > 0) {
+#pragma unroll
+        for (int i0 = 0; i0 < DC; ++i0) {
+            sumf += s[i0] * c[i0];
+        }
+    } else {
+        for (int i0 = 0; i0 < d_conv; ++i0) {
+            sumf += s[i0] * c[i0];
+        }
+    }
+
+    // fused bias add: the ADD node broadcasts a 1-D channel bias over tokens
+    if (bias != nullptr) {
+        sumf += bias[channel];
+    }
+
+    const size_t dst_idx =
+        static_cast<size_t>(seq) * static_cast<size_t>(dst_stride_seq) +
+        static_cast<size_t>(token) * static_cast<size_t>(dst_stride_token) +
+        static_cast<size_t>(channel);
+
+    dst_data[dst_idx] = apply_silu ? op_silu(sumf) : sumf;
+}
+
+// FUSED=false keeps apply_silu/bias out of the kernel capture list, so the unfused launch
+// takes the pre-fusion argument list; matters at n_t == 1, where the op is launch-bound.
+template <int DC, bool FUSED>
+static void kernel_ssm_conv_impl(
    queue &q,
    const float *src_data,
    const float *weights,
@@ -18,7 +78,9 @@ static void kernel_ssm_conv(
    int src_stride_inner,
    int src_stride_seq,
    int dst_stride_token,
-    int dst_stride_seq
+    int dst_stride_seq,
+    bool apply_silu,
+    const float *bias
 ) {
    const size_t total_work = static_cast<size_t>(d_inner) * static_cast<size_t>(n_t) * static_cast<size_t>(n_s);
    const size_t work_group_size = 256;
@@ -27,53 +89,199 @@ static void kernel_ssm_conv(
    const range<1> global_range(num_work_groups * work_group_size);
    const range<1> local_range(work_group_size);

-    q.submit([&](handler &h) {
-        h.parallel_for(
-            nd_range<1>(global_range, local_range),
-            [=](nd_item<1> item) {
-                const size_t idx = item.get_global_id(0);
-                if (idx >= total_work) {
-                    return;
+    if constexpr (FUSED) {
+        q.submit([&](handler &h) {
+            h.parallel_for(
+                nd_range<1>(global_range, local_range),
+                [=](nd_item<1> item) {
+                    const size_t idx = item.get_global_id(0);
+                    if (idx >= total_work) {
+                        return;
+                    }
+
+                    ssm_conv_element<DC>(idx, src_data, weights, dst_data, d_conv, d_inner, n_t,
+                                         src_stride_inner, src_stride_seq, dst_stride_token,
+                                         dst_stride_seq, apply_silu, bias);
                }
+            );
+        });
+    } else {
+        GGML_UNUSED(apply_silu);
+        GGML_UNUSED(bias);

-                // src has the tokens of one channel contiguous, dst has the channels of one
-                // token contiguous, so either the loads or the store must be strided. Indexing
-                // token-fastest coalesces the d_conv loads, which measured faster except for
-                // short, cache-resident rows.
-                const int token   = static_cast<int>(idx % n_t);
-                const int channel = static_cast<int>((idx / n_t) % d_inner);
-                const int seq     = static_cast<int>(idx / (static_cast<size_t>(n_t) * static_cast<size_t>(d_inner)));
+        q.submit([&](handler &h) {
+            h.parallel_for(
+                nd_range<1>(global_range, local_range),
+                [=](nd_item<1> item) {
+                    const size_t idx = item.get_global_id(0);
+                    if (idx >= total_work) {
+                        return;
+                    }

-                const float *s = src_data
-                    + static_cast<size_t>(seq) * static_cast<size_t>(src_stride_seq)
-                    + static_cast<size_t>(channel) * static_cast<size_t>(src_stride_inner)
-                    + static_cast<size_t>(token);
-
-                const float *c = weights + static_cast<size_t>(channel) * static_cast<size_t>(d_conv);
-
-                float sumf = 0.0f;
-                for (int i0 = 0; i0 < d_conv; ++i0) {
-                    sumf += s[i0] * c[i0];
+                    ssm_conv_element<DC>(idx, src_data, weights, dst_data, d_conv, d_inner, n_t,
+                                         src_stride_inner, src_stride_seq, dst_stride_token,
+                                         dst_stride_seq, false, nullptr);
                }
-
-                const size_t dst_idx =
-                    static_cast<size_t>(seq) * static_cast<size_t>(dst_stride_seq) +
-                    static_cast<size_t>(token) * static_cast<size_t>(dst_stride_token) +
-                    static_cast<size_t>(channel);
-
-                dst_data[dst_idx] = sumf;
-            }
-        );
-    });
+            );
+        });
+    }
 }

-inline void ggml_sycl_op_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+// SLM transpose tile: coalesces both the loads and the stores. The +1 pad makes the row
+// stride 33, coprime with 32 banks, so both phases are bank-conflict-free.
+template <int DC, int TT, int TC, int WG>
+static __dpct_inline__ void ssm_conv_tile(
+    nd_item<1> it, local_accessor<float, 1> tile, const float *src_data, const float *weights,
+    float *dst_data, int n_t, int nt_tiles, int nc_tiles, int src_stride_inner,
+    int src_stride_seq, int dst_stride_token, int dst_stride_seq, bool apply_silu,
+    const float *bias
+) {
+    const int    lid = static_cast<int>(it.get_local_id(0));
+    const size_t g   = it.get_group(0);
+    const int    tt  = static_cast<int>(g % nt_tiles);
+    const int    ct  = static_cast<int>((g / nt_tiles) % nc_tiles);
+    const int    seq = static_cast<int>(g / (static_cast<size_t>(nt_tiles) * nc_tiles));
+    const int    t0 = tt * TT, c0 = ct * TC;
+
+    const int ti = lid % TT;
+    const int cj = lid / TT;
+#pragma unroll
+    for (int r = 0; r < TC / (WG / TT); ++r) {
+        const int c   = cj + r * (WG / TT);
+        const int tok = t0 + ti;
+        float sumf = 0.0f;
+        if (tok < n_t) {
+            const float *s = src_data + static_cast<size_t>(seq) * src_stride_seq
+                           + static_cast<size_t>(c0 + c) * src_stride_inner + tok;
+            const float *cw = weights + static_cast<size_t>(c0 + c) * DC;
+#pragma unroll
+            for (int i = 0; i < DC; ++i) sumf += s[i] * cw[i];
+            if (bias != nullptr) sumf += bias[c0 + c];
+            if (apply_silu) sumf = op_silu(sumf);
+        }
+        tile[c * (TT + 1) + ti] = sumf;
+    }
+    it.barrier(access::fence_space::local_space);
+
+    const int cc = lid % TC;
+    const int tj = lid / TC;
+#pragma unroll
+    for (int r = 0; r < TT / (WG / TC); ++r) {
+        const int t   = tj + r * (WG / TC);
+        const int tok = t0 + t;
+        if (tok < n_t) {
+            dst_data[static_cast<size_t>(seq) * dst_stride_seq
+                     + static_cast<size_t>(tok) * dst_stride_token + c0 + cc]
+                = tile[cc * (TT + 1) + t];
+        }
+    }
+}
+
+// Same FUSED split as kernel_ssm_conv_impl. The fused instantiation keeps the runtime
+// apply_silu/bias branches: at n_t >= 32 they are amortized over the whole tile.
+template <int DC, bool FUSED>
+static void kernel_ssm_conv_tiled(
+    queue &q, const float *src_data, const float *weights, float *dst_data,
+    int d_inner, int n_t, int n_s, int src_stride_inner, int src_stride_seq,
+    int dst_stride_token, int dst_stride_seq, bool apply_silu, const float *bias
+) {
+    constexpr int TT = 32, TC = 32, WG = 256;
+    const int nt_tiles = (n_t + TT - 1) / TT;
+    const int nc_tiles = d_inner / TC;
+    const size_t groups = static_cast<size_t>(nt_tiles) * nc_tiles * n_s;
+
+    if constexpr (FUSED) {
+        q.submit([&](handler &h) {
+            local_accessor<float, 1> tile(range<1>(TC * (TT + 1)), h);
+            h.parallel_for(nd_range<1>(range<1>(groups * WG), range<1>(WG)), [=](nd_item<1> it) {
+                ssm_conv_tile<DC, TT, TC, WG>(it, tile, src_data, weights, dst_data, n_t, nt_tiles,
+                                              nc_tiles, src_stride_inner, src_stride_seq,
+                                              dst_stride_token, dst_stride_seq, apply_silu, bias);
+            });
+        });
+    } else {
+        GGML_UNUSED(apply_silu);
+        GGML_UNUSED(bias);
+
+        q.submit([&](handler &h) {
+            local_accessor<float, 1> tile(range<1>(TC * (TT + 1)), h);
+            h.parallel_for(nd_range<1>(range<1>(groups * WG), range<1>(WG)), [=](nd_item<1> it) {
+                ssm_conv_tile<DC, TT, TC, WG>(it, tile, src_data, weights, dst_data, n_t, nt_tiles,
+                                              nc_tiles, src_stride_inner, src_stride_seq,
+                                              dst_stride_token, dst_stride_seq, false, nullptr);
+            });
+        });
+    }
+}
+
+static void kernel_ssm_conv(
+    queue &q,
+    const float *src_data,
+    const float *weights,
+    float *dst_data,
+    int d_conv,
+    int d_inner,
+    int n_t,
+    int n_s,
+    int ncs,
+    int src_stride_inner,
+    int src_stride_seq,
+    int dst_stride_token,
+    int dst_stride_seq,
+    bool apply_silu,
+    const float *bias
+) {
+    // Only the fused instantiations carry apply_silu/bias as kernel arguments; the plain
+    // ssm_conv launch keeps the argument list it had before the fusion landed.
+    const bool fused = apply_silu || bias != nullptr;
+
+    // d_inner must be a multiple of 32 so the channel tiles are exact; the transpose is only
+    // worth it for n_t >= 32. d_conv == 4 is the only window with a DC-specialized kernel.
+    if (d_conv == 4 && n_t >= 32 && (d_inner % 32) == 0) {
+        if (fused) {
+            kernel_ssm_conv_tiled<4, true>(q, src_data, weights, dst_data, d_inner, n_t, n_s,
+                                           src_stride_inner, src_stride_seq, dst_stride_token,
+                                           dst_stride_seq, apply_silu, bias);
+        } else {
+            kernel_ssm_conv_tiled<4, false>(q, src_data, weights, dst_data, d_inner, n_t, n_s,
+                                            src_stride_inner, src_stride_seq, dst_stride_token,
+                                            dst_stride_seq, apply_silu, bias);
+        }
+        return;
+    }
+
+    if (d_conv == 4) {
+        if (fused) {
+            kernel_ssm_conv_impl<4, true>(q, src_data, weights, dst_data, d_conv, d_inner, n_t, n_s,
+                                          ncs, src_stride_inner, src_stride_seq, dst_stride_token,
+                                          dst_stride_seq, apply_silu, bias);
+        } else {
+            kernel_ssm_conv_impl<4, false>(q, src_data, weights, dst_data, d_conv, d_inner, n_t, n_s,
+                                           ncs, src_stride_inner, src_stride_seq, dst_stride_token,
+                                           dst_stride_seq, apply_silu, bias);
+        }
+        return;
+    }
+
+    if (fused) {
+        kernel_ssm_conv_impl<0, true>(q, src_data, weights, dst_data, d_conv, d_inner, n_t, n_s,
+                                      ncs, src_stride_inner, src_stride_seq, dst_stride_token,
+                                      dst_stride_seq, apply_silu, bias);
+    } else {
+        kernel_ssm_conv_impl<0, false>(q, src_data, weights, dst_data, d_conv, d_inner, n_t, n_s,
+                                       ncs, src_stride_inner, src_stride_seq, dst_stride_token,
+                                       dst_stride_seq, apply_silu, bias);
+    }
+}
+
+inline void ggml_sycl_op_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor * dst, ggml_tensor * silu_dst = nullptr, const float * bias = nullptr) {
    ggml_tensor * src0 = dst->src[0];
    ggml_tensor * src1 = dst->src[1];

    GGML_ASSERT(src0->type == GGML_TYPE_F32);
    GGML_ASSERT(src1->type == GGML_TYPE_F32);
    GGML_ASSERT(dst->type  == GGML_TYPE_F32);
+    GGML_ASSERT(bias == nullptr || silu_dst != nullptr);

    const int d_conv   = src1->ne[0];
    const int ncs      = src0->ne[0];
@@ -104,7 +312,8 @@ inline void ggml_sycl_op_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor *

        const float *src_data = static_cast<const float *>(src0->data);
        const float *weights  = static_cast<const float *>(src1->data);
-        float *dst_data       = static_cast<float *>(dst->data);
+        const bool apply_silu = silu_dst != nullptr;
+        float *dst_data       = static_cast<float *>((silu_dst ? silu_dst : dst)->data);

        GGML_ASSERT(src_data && weights && dst_data);

@@ -121,7 +330,9 @@ inline void ggml_sycl_op_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor *
            src_stride_inner,
            src_stride_seq,
            dst_stride_token,
-            dst_stride_seq
+            dst_stride_seq,
+            apply_silu,
+            bias
        );

    } catch (const std::exception &e) {
@@ -134,3 +345,17 @@ void ggml_sycl_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
    ggml_sycl_op_ssm_conv(ctx, dst);
 }
+
+// Fused ssm_conv + ADD + SiLU: write silu(conv(x) + b) straight into silu_dst, eliding the
+// standalone SiLU launch and its HBM round-trip of the conv output.
+void ggml_sycl_ssm_conv_fused(ggml_backend_sycl_context & ctx, ggml_tensor * dst, ggml_tensor * add, ggml_tensor * silu_dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    GGML_ASSERT(silu_dst && ggml_are_same_shape(dst, silu_dst) && silu_dst->type == GGML_TYPE_F32);
+    // the fused kernel reads only the ADD's bias operand; the ADD result is never written
+    const float * bias = nullptr;
+    if (add != nullptr) {
+        const ggml_tensor * bias_t = (add->src[0] == dst) ? add->src[1] : add->src[0];
+        bias = static_cast<const float *>(bias_t->data);
+    }
+    ggml_sycl_op_ssm_conv(ctx, dst, silu_dst, bias);
+}
@@ -3,3 +3,4 @@
 #include "common.hpp"

 void ggml_sycl_ssm_conv(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+void ggml_sycl_ssm_conv_fused(ggml_backend_sycl_context & ctx, ggml_tensor * dst, ggml_tensor * add, ggml_tensor * silu_dst);
@@ -62,6 +62,11 @@ if (Vulkan_FOUND)
    ggml_add_backend_library(ggml-vulkan
                             ggml-vulkan.cpp
                             ../../include/ggml-vulkan.h
+                             ggml-vulkan-types.h
+                             ggml-vulkan-push-constants.h
+                             ggml-vulkan-common.h
+                             ggml-vulkan-buffers.cpp
+                             ggml-vulkan-debug.cpp
                            )

    set(VULKAN_SHADER_GEN_CMAKE_ARGS "")
@@ -0,0 +1,783 @@
+#include "ggml-vulkan-common.h"
+
+ggml_backend_buffer_type_i ggml_backend_vk_buffer_type_interface = {
+    /* .get_name         = */ ggml_backend_vk_buffer_type_name,
+    /* .alloc_buffer     = */ ggml_backend_vk_buffer_type_alloc_buffer,
+    /* .get_alignment    = */ ggml_backend_vk_buffer_type_get_alignment,
+    /* .get_max_size     = */ ggml_backend_vk_buffer_type_get_max_size,
+    /* .get_alloc_size   = */ ggml_backend_vk_buffer_type_get_alloc_size,
+    /* .is_host          = */ NULL,
+};
+
+static std::vector<uint32_t> ggml_vk_find_memory_properties(const vk::PhysicalDeviceMemoryProperties* mem_props, vk::MemoryRequirements* mem_req, vk::MemoryPropertyFlags flags) {
+    std::vector<uint32_t> indices;
+
+    for (uint32_t i = 0; i < mem_props->memoryTypeCount; ++i) {
+        vk::MemoryType memory_type = mem_props->memoryTypes[i];
+        if ((mem_req->memoryTypeBits & ((uint64_t)1 << i)) &&
+            (flags & memory_type.propertyFlags) == flags &&
+            mem_props->memoryHeaps[memory_type.heapIndex].size >= mem_req->size) {
+            indices.push_back(i);
+        }
+    }
+    return indices;
+}
+
+static vk_buffer ggml_vk_create_buffer(vk_device& device, size_t size, const std::initializer_list<vk::MemoryPropertyFlags> & req_flags_list,
+                                       void *import_ptr = nullptr) {
+    VK_LOG_DEBUG("ggml_vk_create_buffer(" << device->name << ", " << size << ", " << to_string(req_flags_list.begin()[0]) << ", " << to_string(req_flags_list.begin()[req_flags_list.size()-1]) << ")");
+    if (size > device->max_buffer_size) {
+        throw vk::OutOfDeviceMemoryError("Requested buffer size exceeds device buffer size limit");
+    }
+
+    vk_buffer buf = std::make_shared<vk_buffer_struct>();
+
+    if (size == 0) {
+        buf->size = 0;
+        return buf;
+    }
+
+    vk::BufferUsageFlags usage_flags = vk::BufferUsageFlagBits::eStorageBuffer | vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eTransferDst;
+    vk::MemoryAllocateFlags mem_flags {};
+    if (device->buffer_device_address) {
+        usage_flags |= vk::BufferUsageFlagBits::eShaderDeviceAddress;
+        mem_flags |= vk::MemoryAllocateFlagBits::eDeviceAddress;
+    }
+
+    vk::BufferCreateInfo buffer_create_info{
+        vk::BufferCreateFlags(),
+        size,
+        usage_flags,
+        vk::SharingMode::eExclusive,
+        0,
+        nullptr,
+    };
+
+    vk::ExternalMemoryBufferCreateInfo external_memory_bci;
+    if (import_ptr) {
+        external_memory_bci.handleTypes = vk::ExternalMemoryHandleTypeFlagBits::eHostAllocationEXT;
+        buffer_create_info.setPNext(&external_memory_bci);
+    }
+
+    buf->buffer = device->device.createBuffer(buffer_create_info);
+
+    vk::MemoryRequirements mem_req = device->device.getBufferMemoryRequirements(buf->buffer);
+
+    vk::PhysicalDeviceMemoryProperties mem_props = device->physical_device.getMemoryProperties();
+
+    const vk::MemoryPriorityAllocateInfoEXT mem_priority_info { 1.0f };
+
+    vk::MemoryAllocateFlagsInfo mem_flags_info { mem_flags };
+
+    if (device->memory_priority) {
+        mem_flags_info.setPNext(&mem_priority_info);
+    }
+
+    if (import_ptr) {
+        vk::MemoryHostPointerPropertiesEXT host_pointer_props;
+        try {
+            host_pointer_props = device->device.getMemoryHostPointerPropertiesEXT(vk::ExternalMemoryHandleTypeFlagBits::eHostAllocationEXT, import_ptr);
+        } catch (vk::SystemError& e) {
+            GGML_LOG_WARN("ggml_vulkan: Failed getMemoryHostPointerPropertiesEXT (%s)\n", e.what());
+            device->device.destroyBuffer(buf->buffer);
+            return {};
+        }
+        vk::PhysicalDeviceMemoryProperties mem_props = device->physical_device.getMemoryProperties();
+
+        uint32_t memory_type_idx;
+        vk::MemoryPropertyFlags property_flags = *req_flags_list.begin();
+        for (memory_type_idx = 0; memory_type_idx < 32; ++memory_type_idx) {
+            if (!(host_pointer_props.memoryTypeBits & (1u << memory_type_idx))) {
+                continue;
+            }
+            if (!(mem_req.memoryTypeBits & (1u << memory_type_idx))) {
+                continue;
+            }
+
+            vk::MemoryType memory_type = mem_props.memoryTypes[memory_type_idx];
+            // check for visible+coherent+cached. Other flags (e.g. devicelocal) are allowed
+            if ((memory_type.propertyFlags & property_flags) == property_flags) {
+                property_flags = memory_type.propertyFlags;
+                break;
+            }
+        }
+        if (memory_type_idx == 32) {
+            GGML_LOG_WARN("ggml_vulkan: Memory type for host allocation not found\n");
+            device->device.destroyBuffer(buf->buffer);
+            return {};
+        }
+
+        buf->memory_property_flags = mem_props.memoryTypes[memory_type_idx].propertyFlags;
+        try {
+            vk::ImportMemoryHostPointerInfoEXT import_info;
+            import_info.handleType = vk::ExternalMemoryHandleTypeFlagBits::eHostAllocationEXT;
+            import_info.pHostPointer = import_ptr;
+            import_info.setPNext(&mem_flags_info);
+            buf->device_memory = device->device.allocateMemory({ size, memory_type_idx, &import_info });
+        } catch (const vk::SystemError& e) {
+        }
+    } else {
+        for (auto it = req_flags_list.begin(); it != req_flags_list.end(); it++) {
+            const auto & req_flags = *it;
+
+            const std::vector<uint32_t> memory_type_indices = ggml_vk_find_memory_properties(&mem_props, &mem_req, req_flags);
+
+            if (memory_type_indices.empty()) {
+                continue;
+            }
+
+            bool done = false;
+
+            for (auto mtype_it = memory_type_indices.begin(); mtype_it != memory_type_indices.end(); mtype_it++) {
+                try {
+                    buf->device_memory = device->device.allocateMemory({ mem_req.size, *mtype_it, &mem_flags_info });
+                    buf->memory_property_flags = mem_props.memoryTypes[*mtype_it].propertyFlags;
+                    done = true;
+                    break;
+                } catch (const vk::SystemError& e) {
+                    // loop and retry
+                    // during last attempt throw the exception
+                    if (it + 1 == req_flags_list.end() && mtype_it + 1 == memory_type_indices.end()) {
+                        device->device.destroyBuffer(buf->buffer);
+                        throw e;
+                    }
+                }
+            }
+
+            if (done) {
+                break;
+            }
+        }
+    }
+
+    if (!buf->device_memory) {
+        device->device.destroyBuffer(buf->buffer);
+        throw vk::OutOfDeviceMemoryError("No suitable memory type found");
+    }
+
+    buf->ptr = nullptr;
+
+    if (import_ptr) {
+        buf->ptr = import_ptr;
+    } else {
+        if (buf->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible) {
+            buf->ptr = device->device.mapMemory(buf->device_memory, 0, VK_WHOLE_SIZE);
+        }
+    }
+
+    device->device.bindBufferMemory(buf->buffer, buf->device_memory, 0);
+
+    buf->device = device;
+    buf->size = size;
+
+    if (device->buffer_device_address) {
+        const vk::BufferDeviceAddressInfo addressInfo(buf->buffer);
+        buf->bda_addr = device->device.getBufferAddress(addressInfo);
+    }
+
+    device->memory_logger->log_allocation(buf, size);
+
+    return buf;
+}
+
+vk_buffer ggml_vk_create_buffer_check(vk_device& device, size_t size, vk::MemoryPropertyFlags req_flags, vk::MemoryPropertyFlags fallback_flags) {
+    try {
+        return ggml_vk_create_buffer(device, size, {req_flags, fallback_flags});
+    } catch (const vk::SystemError& e) {
+        std::cerr << "ggml_vulkan: Memory allocation of size " << size << " failed." << std::endl;
+        std::cerr << "ggml_vulkan: " << e.what() << std::endl;
+        throw e;
+    }
+}
+
+vk_buffer ggml_vk_create_buffer_device(vk_device& device, size_t size) {
+    vk_buffer buf;
+    try {
+        if (device->prefer_host_memory) {
+            buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent,
+                                                       vk::MemoryPropertyFlagBits::eDeviceLocal});
+        } else if (device->uma) {
+            // On UMA, prefer host-visible memory so direct tensor borrowing works.
+            // If unavailable, fall back to device-local memory.
+            buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eDeviceLocal | vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent,
+                                                       vk::MemoryPropertyFlagBits::eDeviceLocal,
+                                                       vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent});
+        } else if (device->disable_host_visible_vidmem) {
+            if (device->allow_sysmem_fallback) {
+                buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eDeviceLocal,
+                                                           vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent});
+            } else {
+                buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eDeviceLocal});
+            }
+        } else {
+            // use rebar if available, otherwise fallback to device only visible memory
+            if (device->allow_sysmem_fallback) {
+                buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eDeviceLocal | vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent,
+                                                           vk::MemoryPropertyFlagBits::eDeviceLocal,
+                                                           vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent});
+            } else {
+                buf = ggml_vk_create_buffer(device, size, {vk::MemoryPropertyFlagBits::eDeviceLocal | vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent,
+                                                           vk::MemoryPropertyFlagBits::eDeviceLocal});
+            }
+        }
+    } catch (const vk::SystemError& e) {
+        std::cerr << "ggml_vulkan: Device memory allocation of size " << size << " failed." << std::endl;
+        std::cerr << "ggml_vulkan: " << e.what() << std::endl;
+        throw e;
+    }
+
+    return buf;
+}
+
+void ggml_vk_destroy_buffer(vk_buffer& buf) {
+    if (buf == nullptr) {
+        return;
+    }
+
+    if (buf->device != nullptr) {
+        buf->device->memory_logger->log_deallocation(buf);
+    }
+
+    buf.reset();
+}
+
+void * ggml_vk_host_malloc(vk_device& device, size_t size) {
+    VK_LOG_MEMORY("ggml_vk_host_malloc(" << size << ")");
+    vk_buffer buf = ggml_vk_create_buffer(device, size,
+        {vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent | vk::MemoryPropertyFlagBits::eHostCached,
+         vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent});
+
+    if(!(buf->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible)) {
+        fprintf(stderr, "WARNING: failed to allocate %.2f MB of pinned memory\n",
+            size/1024.0/1024.0);
+        device->device.freeMemory(buf->device_memory);
+        device->device.destroyBuffer(buf->buffer);
+        return nullptr;
+    }
+
+    std::lock_guard<std::shared_mutex> guard(device->pinned_memory_mutex);
+    device->pinned_memory.push_back(std::make_tuple(buf->ptr, size, buf));
+
+    return buf->ptr;
+}
+
+void ggml_vk_host_free(vk_device& device, void* ptr) {
+    if (ptr == nullptr) {
+        return;
+    }
+    VK_LOG_MEMORY("ggml_vk_host_free(" << ptr << ")");
+    std::lock_guard<std::shared_mutex> guard(device->pinned_memory_mutex);
+
+    vk_buffer buf;
+    size_t index;
+    for (size_t i = 0; i < device->pinned_memory.size(); i++) {
+        const uint8_t* addr = (const uint8_t*) std::get<0>(device->pinned_memory[i]);
+        const uint8_t* endr = addr + std::get<1>(device->pinned_memory[i]);
+        if (ptr >= addr && ptr < endr) {
+            buf = std::get<2>(device->pinned_memory[i]);
+            index = i;
+            break;
+        }
+    }
+    if (buf == nullptr) {
+        fprintf(stderr, "WARNING: failed to free pinned memory: memory not in map\n");
+        return;
+    }
+
+    ggml_vk_destroy_buffer(buf);
+
+    device->pinned_memory.erase(device->pinned_memory.begin() + index);
+}
+
+void ggml_vk_host_get(const vk_device& device, const void * ptr, vk_buffer& buf, size_t& buf_offset) {
+    std::shared_lock<std::shared_mutex> guard(device->pinned_memory_mutex);
+    buf = nullptr;
+    buf_offset = 0;
+    for (size_t i = 0; i < device->pinned_memory.size(); i++) {
+        const uint8_t* addr = (const uint8_t*) std::get<0>(device->pinned_memory[i]);
+        const uint8_t* endr = addr + std::get<1>(device->pinned_memory[i]);
+        if (ptr >= addr && ptr < endr) {
+            buf = std::get<2>(device->pinned_memory[i]);
+            buf_offset = ((const uint8_t *)ptr) - addr;
+            break;
+        }
+    }
+}
+
+void ggml_vk_ensure_sync_staging_buffer(vk_device& device, size_t size) {
+    if (device->sync_staging == nullptr || device->sync_staging->size < size) {
+        VK_LOG_MEMORY("ggml_vk_ensure_sync_staging_buffer(" << size << ")");
+        ggml_vk_destroy_buffer(device->sync_staging);
+        device->sync_staging = ggml_vk_create_buffer_check(device, size,
+            vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent | vk::MemoryPropertyFlagBits::eHostCached,
+            vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent);
+    }
+}
+
+void ggml_vk_ensure_sync_staging_buffer(ggml_backend_vk_context * ctx, size_t size) {
+    if (ctx->sync_staging == nullptr || ctx->sync_staging->size < size) {
+        VK_LOG_MEMORY("ggml_vk_ensure_sync_staging_buffer(" << size << ")");
+        ggml_vk_destroy_buffer(ctx->sync_staging);
+        ctx->sync_staging = ggml_vk_create_buffer_check(ctx->device, size,
+            vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent | vk::MemoryPropertyFlagBits::eHostCached,
+            vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent);
+    }
+}
+
+static void ggml_vk_buffer_write_nc_async(ggml_backend_vk_context * ctx, vk_context& subctx, vk_buffer& dst, size_t offset, const ggml_tensor * tensor, bool sync_staging = false) {
+    VK_LOG_DEBUG("ggml_vk_buffer_write_nc_async(" << tensor << ")");
+    GGML_ASSERT(!ggml_is_contiguous(tensor));
+    // Buffer is already mapped
+    if(dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible) {
+        std::cerr << "ggml_vulkan: buffer_write_nc_async dst buffer is host_visible. Use synchronous write." << std::endl;
+        GGML_ABORT("fatal error");
+    }
+    // Check if src is pinned memory
+    vk_buffer buf = nullptr;
+    size_t buf_offset = 0;
+    ggml_vk_host_get(ctx->device, tensor->data, buf, buf_offset);
+
+    const uint64_t ne0 = tensor->ne[0];
+    const uint64_t ne1 = tensor->ne[1];
+    const uint64_t ne2 = tensor->ne[2];
+    const uint64_t ne3 = tensor->ne[3];
+    const uint64_t nb0 = tensor->nb[0];
+    const uint64_t nb1 = tensor->nb[1];
+    const uint64_t nb2 = tensor->nb[2];
+    const uint64_t nb3 = tensor->nb[3];
+    const ggml_type type = tensor->type;
+    const uint64_t ts = ggml_type_size(type);
+    const uint64_t bs = ggml_blck_size(type);
+
+    const uint64_t dstnb0 = ts;
+    const uint64_t dstnb1 = dstnb0*(ne0/bs);
+    const uint64_t dstnb2 = dstnb1*ne1;
+    const uint64_t dstnb3 = dstnb2*ne2;
+
+    const uint64_t ne = ggml_nelements(tensor);
+
+    if (buf != nullptr) {
+        // Memory is pinned, use as staging buffer
+        std::vector<vk::BufferCopy> slices;
+
+        for (uint64_t i3 = 0; i3 < ne3; i3++) {
+            for (uint64_t i2 = 0; i2 < ne2; i2++) {
+                // Find longest contiguous slice
+                if (ne1*nb1 == dstnb2) {
+                    slices.push_back({ buf_offset + i3*nb3 + i2*nb2, offset + i3*dstnb3 + i2*dstnb2, dstnb2 });
+                } else {
+                    for (uint64_t i1 = 0; i1 < ne1; i1++) {
+                        if (ne0*nb0/bs == dstnb1) {
+                            slices.push_back({ buf_offset + i3*nb3 + i2*nb2 + i1*nb1, offset + i3*dstnb3 + i2*dstnb2 + i1*dstnb1, dstnb1 });
+                        } else {
+                            const uint64_t s_off = buf_offset + i3*nb3 + i2*nb2 + i1*nb1;
+                            const uint64_t d_off = offset + i3*dstnb3 + i2*dstnb2 + i1*dstnb1;
+                            for (uint64_t i0 = 0; i0 < ne0; i0++) {
+                                slices.push_back({ s_off + i0*nb0, d_off + i0*dstnb0, dstnb0 });
+                            }
+                        }
+                    }
+                }
+            }
+        }
+
+        ggml_vk_sync_buffers(ctx, subctx);
+        subctx->s->buffer->buf.copyBuffer(buf->buffer, dst->buffer, slices);
+        return;
+    }
+
+    if (!sync_staging) {
+        GGML_ABORT("Asynchronous write to non-pinned memory not supported");
+    }
+
+    // Staging buffer required
+    vk_buffer& staging = ctx->device->sync_staging;
+    const uint64_t copy_size = ts*ne/bs;
+    ggml_vk_ensure_sync_staging_buffer(ctx->device, copy_size);
+    VkBufferCopy buf_copy{ 0, offset, copy_size };
+
+    ggml_vk_sync_buffers(ctx, subctx);
+    vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)staging->buffer, (VkBuffer)dst->buffer, 1, &buf_copy);
+
+    for (uint64_t i3 = 0; i3 < ne3; i3++) {
+        for (uint64_t i2 = 0; i2 < ne2; i2++) {
+            // Find longest contiguous slice
+            if (ne1*nb1 == dstnb2) {
+                deferred_memcpy((uint8_t *)staging->ptr + i3*dstnb3 + i2*dstnb2, (const uint8_t *) tensor->data + buf_offset + i3*nb3 + i2*nb2, dstnb2, &subctx->in_memcpys);
+            } else {
+                for (uint64_t i1 = 0; i1 < ne1; i1++) {
+                    if (ne0*nb0/bs == dstnb1) {
+                        deferred_memcpy((uint8_t *)staging->ptr + i3*dstnb3 + i2*dstnb2 + i1*dstnb1, (const uint8_t *) tensor->data + buf_offset + i3*nb3 + i2*nb2 + i1*nb1, dstnb1, &subctx->in_memcpys);
+                    } else {
+                        const uint64_t s_off = buf_offset + i3*nb3 + i2*nb2 + i1*nb1;
+                        const uint64_t d_off = i3*dstnb3 + i2*dstnb2 + i1*dstnb1;
+                        for (uint64_t i0 = 0; i0 < ne0; i0++) {
+                            deferred_memcpy((uint8_t *)staging->ptr + d_off + i0*dstnb0, (const uint8_t *) tensor->data + s_off + i0*nb0, dstnb0, &subctx->in_memcpys);
+                        }
+                    }
+                }
+            }
+        }
+    }
+}
+
+bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging) {
+    VK_LOG_DEBUG("ggml_vk_buffer_write_2d_async(" << width << ", " << height << ")");
+    // Check if src is pinned memory
+    vk_buffer buf = nullptr;
+    size_t buf_offset = 0;
+    ggml_vk_host_get(dst->device, src, buf, buf_offset);
+
+    if (buf != nullptr) {
+        // Memory is pinned, use as staging buffer
+        std::vector<vk::BufferCopy> slices(1);
+        if (width == spitch && width == dpitch) {
+            // Only do single write if stride is equal
+            slices[0].srcOffset = buf_offset;
+            slices[0].dstOffset = offset;
+            slices[0].size = width * height;
+        } else {
+            slices.resize(height);
+            for (size_t i = 0; i < height; i++) {
+                slices[i].srcOffset = buf_offset + i * spitch;
+                slices[i].dstOffset = offset + i * dpitch;
+                slices[i].size = width;
+            }
+        }
+
+        ggml_vk_sync_buffers(nullptr, subctx);
+        subctx->s->buffer->buf.copyBuffer(buf->buffer, dst->buffer, slices);
+        return true;
+    }
+    VK_LOG_DEBUG("STAGING");
+
+    if (!sync_staging) {
+        // copy was not handled caller needs to fall back
+        return false;
+    }
+
+    // Staging buffer required
+    const size_t staging_size = width * height;
+    ggml_vk_ensure_sync_staging_buffer(dst->device, staging_size);
+
+    vk_buffer& staging_buffer = dst->device->sync_staging;
+
+    std::vector<vk::BufferCopy> slices(1);
+    if (width == dpitch) {
+        slices[0].srcOffset = 0;
+        slices[0].dstOffset = offset;
+        slices[0].size = staging_size;
+    } else {
+        slices.resize(height);
+        for (size_t i = 0; i < height; i++) {
+            slices[i].srcOffset = i * width;
+            slices[i].dstOffset = offset + i * dpitch;
+            slices[i].size = width;
+        }
+    }
+
+    ggml_vk_sync_buffers(nullptr, subctx);
+    subctx->s->buffer->buf.copyBuffer(staging_buffer->buffer, dst->buffer, slices);
+
+    if (width == spitch) {
+        deferred_memcpy((uint8_t *)staging_buffer->ptr, src, staging_size, &subctx->in_memcpys);
+    } else {
+        for (size_t i = 0; i < height; i++) {
+            deferred_memcpy((uint8_t *)staging_buffer->ptr + i * width, (const uint8_t *) src + i * spitch, width, &subctx->in_memcpys);
+        }
+    }
+    return true;
+}
+
+bool ggml_vk_buffer_write_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t size, bool sync_staging) {
+    VK_LOG_DEBUG("ggml_vk_buffer_write_async(" << size << ")");
+    return ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, size, size, size, 1, sync_staging);
+}
+
+void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height) {
+    VK_LOG_DEBUG("ggml_vk_buffer_write_2d(" << width << ", " << height << ")");
+    // Buffer is already mapped
+    if(dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible) {
+        GGML_ASSERT(dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCoherent);
+
+        if (width == spitch && width == dpitch) {
+            memcpy((uint8_t *)dst->ptr + offset, src, width * height);
+        } else {
+            for (size_t i = 0; i < height; i++) {
+                memcpy((uint8_t *)dst->ptr + offset + i * dpitch, (const uint8_t *) src + i * spitch, width);
+            }
+        }
+    } else {
+        std::lock_guard<std::recursive_mutex> guard(dst->device->mutex);
+
+        vk_context subctx = ggml_vk_create_temporary_context(dst->device->transfer_queue->cmd_pool);
+        ggml_vk_ctx_begin(dst->device, subctx);
+        bool ret = ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, spitch, dpitch, width, height, true);
+        GGML_ASSERT(ret);
+        ggml_vk_ctx_end(subctx);
+
+        for (auto& cpy : subctx->in_memcpys) {
+            memcpy(cpy.dst, cpy.src, cpy.n);
+        }
+
+        for (auto& mset : subctx->memsets) {
+            memset(mset.dst, mset.val, mset.n);
+        }
+
+        ggml_vk_submit(subctx, dst->device->fence);
+        VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_buffer_write_2d waitForFences", dst->device);
+        dst->device->device.resetFences({ dst->device->fence });
+        ggml_vk_queue_command_pools_cleanup(dst->device);
+    }
+}
+
+void ggml_vk_buffer_write(vk_buffer& dst, size_t offset, const void * src, size_t size) {
+    VK_LOG_DEBUG("ggml_vk_buffer_write(" << size << ")");
+    ggml_vk_buffer_write_2d(dst, offset, src, size, size, size, 1);
+}
+
+bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging) {
+    VK_LOG_DEBUG("ggml_vk_buffer_read_2d_async(offset=" << offset << ", width=" << width << ", height=" << height << ")");
+    GGML_ASSERT(width > 0);
+    GGML_ASSERT(height > 0);
+    GGML_ASSERT(src != nullptr);
+
+    // TODO: staging_offset is not used
+
+    // Check if dst is pinned memory
+    vk_buffer buf = nullptr;
+    size_t buf_offset = 0;
+    ggml_vk_host_get(src->device, dst, buf, buf_offset);
+
+    std::vector<vk::BufferCopy> slices(1);
+    if (width == spitch && width == dpitch) {
+        // Only do single write if stride is equal
+        slices[0].srcOffset = offset;
+        slices[0].dstOffset = buf_offset;
+        slices[0].size = width * height;
+    } else {
+        slices.resize(height);
+        for (size_t i = 0; i < height; i++) {
+            slices[i].srcOffset = offset + i * spitch;
+            slices[i].dstOffset = buf_offset + i * dpitch;
+            slices[i].size = width;
+        }
+    }
+
+    if (buf != nullptr) {
+        // Memory is pinned, use as staging buffer
+        ggml_vk_sync_buffers(nullptr, subctx);
+        subctx->s->buffer->buf.copyBuffer(src->buffer, buf->buffer, slices);
+
+        return true;
+    }
+    VK_LOG_DEBUG("STAGING");
+
+    if (!sync_staging) {
+        // copy was not handled caller needs to fall back
+        return false;
+    }
+
+    // Fall back to staging buffer
+    const size_t staging_size = width * height;
+    ggml_vk_ensure_sync_staging_buffer(src->device, staging_size);
+
+    vk_buffer& staging_buffer = src->device->sync_staging;
+
+    std::vector<vk::BufferCopy> staging_slices(1);
+    if (width == spitch) {
+        staging_slices[0].srcOffset = offset;
+        staging_slices[0].dstOffset = 0;
+        staging_slices[0].size = staging_size;
+    } else {
+        staging_slices.resize(height);
+        for (size_t i = 0; i < height; i++) {
+            staging_slices[i].srcOffset = offset + i * spitch;
+            staging_slices[i].dstOffset = i * width;
+            staging_slices[i].size = width;
+        }
+    }
+
+    ggml_vk_sync_buffers(nullptr, subctx);
+    subctx->s->buffer->buf.copyBuffer(src->buffer, staging_buffer->buffer, staging_slices);
+
+    if (width == dpitch) {
+        deferred_memcpy(dst, staging_buffer->ptr, staging_size, &subctx->out_memcpys);
+    } else {
+        for (size_t i = 0; i < height; i++) {
+            deferred_memcpy((uint8_t *) dst + i * dpitch, (const uint8_t *) staging_buffer->ptr + i * width, width, &subctx->out_memcpys);
+        }
+    }
+    return true;
+}
+
+static bool ggml_vk_buffer_read_async(vk_context subctx, vk_buffer& src, size_t offset, void * dst, size_t size, bool sync_staging = false) {
+    return ggml_vk_buffer_read_2d_async(subctx, src, offset, dst, size, size, size, 1, sync_staging);
+}
+
+void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height) {
+    VK_LOG_DEBUG("ggml_vk_buffer_read_2d(" << src->buffer << ", " << offset << ", " << width << ", " << height << ")");
+
+    // If the device is not an UMA device the memory is host-accessible through rebar. While writing
+    // through PCIe is sufficient fast reading back data from PCIe is slower than going through
+    // the HW device to host copy path.
+    if(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible && src->device->uma) {
+        GGML_ASSERT(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCoherent);
+
+        std::lock_guard<std::recursive_mutex> guard(src->device->mutex);
+        vk_context subctx = ggml_vk_create_temporary_context(src->device->compute_queue->cmd_pool);
+        ggml_vk_ctx_begin(src->device, subctx);
+        subctx->s->buffer->buf.pipelineBarrier(
+            vk::PipelineStageFlagBits::eComputeShader | vk::PipelineStageFlagBits::eTransfer,
+            vk::PipelineStageFlagBits::eHost,
+            {},
+            { { vk::AccessFlagBits::eShaderWrite | vk::AccessFlagBits::eTransferWrite,
+                vk::AccessFlagBits::eHostRead } },
+            {}, {});
+        ggml_vk_ctx_end(subctx);
+        ggml_vk_submit(subctx, src->device->fence);
+        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX),
+                 "vk_buffer_read_2d uma waitForFences", src->device);
+        src->device->device.resetFences({ src->device->fence });
+        ggml_vk_queue_command_pools_cleanup(src->device);
+
+        if (width == spitch && width == dpitch) {
+            memcpy(dst, (const uint8_t *) src->ptr + offset, width * height);
+        } else {
+            for (size_t i = 0; i < height; i++) {
+                memcpy((uint8_t *) dst + i * dpitch, (const uint8_t *) src->ptr + offset + i * spitch, width);
+            }
+        }
+    } else {
+        std::lock_guard<std::recursive_mutex> guard(src->device->mutex);
+
+        vk_context subctx = ggml_vk_create_temporary_context(src->device->transfer_queue->cmd_pool);
+        ggml_vk_ctx_begin(src->device, subctx);
+        bool ret = ggml_vk_buffer_read_2d_async(subctx, src, offset, dst, spitch, dpitch, width, height, true);
+        GGML_ASSERT(ret);
+        ggml_vk_ctx_end(subctx);
+
+        ggml_vk_submit(subctx, src->device->fence);
+        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_read_2d waitForFences", src->device);
+        src->device->device.resetFences({ src->device->fence });
+        ggml_vk_queue_command_pools_cleanup(src->device);
+
+        for (auto& cpy : subctx->out_memcpys) {
+            memcpy(cpy.dst, cpy.src, cpy.n);
+        }
+    }
+}
+
+void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_t size) {
+    VK_LOG_DEBUG("ggml_vk_buffer_read(" << src->buffer << ", " << offset << ", " << size << ")");
+    ggml_vk_buffer_read_2d(src, offset, dst, size, size, size, 1);
+}
+
+void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) {
+    VK_LOG_DEBUG("ggml_vk_buffer_copy_async(" << size << ")");
+    // Make sure both buffers are on same device
+    GGML_ASSERT(src->device == dst->device);
+
+    VkBufferCopy bc{ src_offset, dst_offset, size };
+
+    vkCmdCopyBuffer(ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)dst->buffer, 1, &bc);
+}
+
+void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) {
+    if (src->device == dst->device) {
+        std::lock_guard<std::recursive_mutex> guard(src->device->mutex);
+        VK_LOG_DEBUG("ggml_vk_buffer_copy(SINGLE_DEVICE, " << size << ")");
+        // Copy within the device
+        vk_context subctx = ggml_vk_create_temporary_context(src->device->transfer_queue->cmd_pool);
+        ggml_vk_ctx_begin(src->device, subctx);
+        ggml_vk_buffer_copy_async(subctx, dst, dst_offset, src, src_offset, size);
+        ggml_vk_ctx_end(subctx);
+        ggml_vk_submit(subctx, src->device->fence);
+        VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_copy waitForFences", src->device);
+        src->device->device.resetFences({ src->device->fence });
+        ggml_vk_queue_command_pools_cleanup(src->device);
+    } else {
+        VK_LOG_DEBUG("ggml_vk_buffer_copy(MULTI_DEVICE, " << size << ")");
+        // Copy device to device
+        ggml_vk_ensure_sync_staging_buffer(src->device, size);
+
+        // Copy to src staging buffer
+        ggml_vk_buffer_copy(src->device->sync_staging, 0, src, src_offset, size);
+        // Copy to dst buffer
+        ggml_vk_buffer_write(dst, dst_offset, src->device->sync_staging->ptr, size);
+    }
+}
+
+void ggml_vk_buffer_memset_async(vk_context& ctx, vk_buffer& dst, size_t offset, uint32_t c, size_t size) {
+    VK_LOG_DEBUG("ggml_vk_buffer_memset_async(" << offset << ", " << c << ", " << size << ")");
+
+    if (dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible &&
+        dst->device->uma) {
+        deferred_memset((uint8_t*)dst->ptr + offset, c, size, &ctx->memsets);
+        return;
+    }
+
+    // Fall back to GPU fillBuffer for non-UMA or non-host-visible buffers
+    ctx->s->buffer->buf.fillBuffer(dst->buffer, offset, size, c);
+}
+
+void ggml_vk_buffer_memset(vk_buffer& dst, size_t offset, uint32_t c, size_t size) {
+    VK_LOG_DEBUG("ggml_vk_buffer_memset(" << offset << ", " << c << ", " << size << ")");
+
+    if (dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible &&
+        dst->device->uma) {
+        memset((uint8_t*)dst->ptr + offset, c, size);
+        return;
+    }
+
+    std::lock_guard<std::recursive_mutex> guard(dst->device->mutex);
+    vk_context subctx = ggml_vk_create_temporary_context(dst->device->transfer_queue->cmd_pool);
+    ggml_vk_ctx_begin(dst->device, subctx);
+    subctx->s->buffer->buf.fillBuffer(dst->buffer, offset, size, c);
+    ggml_vk_ctx_end(subctx);
+
+    ggml_vk_submit(subctx, dst->device->fence);
+    VK_CHECK(dst->device->device.waitForFences({ dst->device->fence }, true, UINT64_MAX), "vk_memset waitForFences", dst->device);
+    dst->device->device.resetFences({ dst->device->fence });
+    ggml_vk_queue_command_pools_cleanup(dst->device);
+}
+
+ggml_backend_buffer_i ggml_backend_vk_buffer_interface = {
+    /* .free_buffer     = */ ggml_backend_vk_buffer_free_buffer,
+    /* .get_base        = */ ggml_backend_vk_buffer_get_base,
+    /* .init_tensor     = */ ggml_backend_vk_buffer_init_tensor,
+    /* .memset_tensor   = */ ggml_backend_vk_buffer_memset_tensor,
+    /* .set_tensor      = */ ggml_backend_vk_buffer_set_tensor,
+    /* .get_tensor      = */ ggml_backend_vk_buffer_get_tensor,
+    /* .set_tensor_2d   = */ ggml_backend_vk_buffer_set_tensor_2d,
+    /* .get_tensor_2d   = */ ggml_backend_vk_buffer_get_tensor_2d,
+    /* .cpy_tensor      = */ ggml_backend_vk_buffer_cpy_tensor,
+    /* .clear           = */ ggml_backend_vk_buffer_clear,
+    /* .reset           = */ NULL,
+};
+
+vk_buffer ggml_vk_buffer_from_host_ptr(vk_device & device, void * ptr, size_t size) {
+    if (!device->external_memory_host) {
+        return {};
+    }
+
+    uintptr_t uptr = reinterpret_cast<uintptr_t>(ptr);
+    if (uptr & (device->min_imported_host_pointer_alignment - 1)) {
+        return {};
+    }
+    if (size & (device->min_imported_host_pointer_alignment - 1)) {
+        return {};
+    }
+
+    const vk::MemoryPropertyFlags property_flags = vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent | vk::MemoryPropertyFlagBits::eHostCached;
+
+    vk_buffer buf {};
+    try {
+        buf = ggml_vk_create_buffer(device, size, { property_flags }, ptr);
+    } catch (vk::SystemError& e) {
+        GGML_LOG_WARN("ggml_vulkan: Failed ggml_vk_create_buffer (%s)\n", e.what());
+    }
+
+    return buf;
+}
+
@@ -0,0 +1,282 @@
+#pragma once
+#include "ggml-vulkan-push-constants.h"
+
+// shared globals
+extern ggml_backend_buffer_type_i ggml_backend_vk_buffer_type_interface;
+extern bool vk_memory_logger_enabled;
+extern bool vk_perf_logger_enabled;
+extern bool vk_perf_logger_concurrent;
+extern bool vk_enable_sync_logger;
+extern uint32_t vk_perf_logger_frequency;
+extern std::string vk_pipeline_stats_filter;
+extern void * const vk_ptr_base;
+extern vk_instance_t vk_instance;
+extern ggml_backend_buffer_i ggml_backend_vk_buffer_interface;
+
+// instance
+vk_device ggml_vk_get_device(size_t idx);
+DispatchLoaderDynamic & ggml_vk_default_dispatcher();
+void ggml_vk_instance_init();
+void ggml_vk_init(ggml_backend_vk_context * ctx, size_t idx);
+int ggml_vk_get_device_count();
+void ggml_vk_get_device_description(int device, char * description, size_t description_size);
+bool ggml_vk_instance_layer_settings_available();
+bool ggml_vk_instance_portability_enumeration_ext_available(const std::vector<vk::ExtensionProperties>& instance_extensions);
+bool ggml_vk_instance_debug_utils_ext_available(const std::vector<vk::ExtensionProperties> & instance_extensions);
+bool ggml_vk_device_is_supported(const vk::PhysicalDevice & vkdev);
+bool ggml_vk_khr_cooperative_matrix_support(const vk::PhysicalDeviceProperties& props, const vk::PhysicalDeviceDriverProperties& driver_props, vk_device_architecture arch);
+uint32_t ggml_vk_intel_shader_core_count(const vk::PhysicalDevice& vkdev);
+bool ggml_vk_intel_windows_driver_in_range(uint32_t driver_version, uint32_t lower_major, uint32_t lower_minor, uint32_t upper_major, uint32_t upper_minor);
+
+// shaders
+void ggml_vk_destroy_pipeline(vk::Device& device, vk_pipeline& pipeline);
+vk_fa_tuning_params get_fa_tuning_params(const vk_device& device, uint32_t hsk, uint32_t hsv, uint32_t n_rows, uint32_t n_kv, ggml_type k_type, ggml_type v_type, bool f32acc);
+vk_fa_pipeline_state get_fa_pipeline_state(const vk_device& device, const vk_fa_tuning_params& params, uint32_t hsk, uint32_t hsv, bool aligned, bool f32acc, bool use_mask, bool use_mask_opt, bool use_logit_softcap, bool use_sparse, ggml_type k_type, ggml_type v_type);
+uint32_t get_subgroup_size(const std::string &pipeline_name, const vk_device_architecture &arch);
+void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested = nullptr);
+bool ggml_vk_flash_attn_scalar_shmem_support(const vk_device& device, const vk_fa_tuning_params& params, uint32_t hsk, uint32_t hsv, bool f32acc, ggml_type k_type, ggml_type v_type);
+bool ggml_vk_flash_attn_coopmat_shmem_support(const vk_device& device, const vk_fa_tuning_params& params, uint32_t hsk, uint32_t hsv, bool f32acc, ggml_type k_type = GGML_TYPE_F16, ggml_type v_type = GGML_TYPE_F16);
+
+// buffers
+vk_buffer ggml_vk_create_buffer_check(vk_device& device, size_t size, vk::MemoryPropertyFlags req_flags, vk::MemoryPropertyFlags fallback_flags = vk::MemoryPropertyFlags(0));
+vk_buffer ggml_vk_create_buffer_device(vk_device& device, size_t size);
+void ggml_vk_destroy_buffer(vk_buffer& buf);
+void * ggml_vk_host_malloc(vk_device& device, size_t size);
+void ggml_vk_host_free(vk_device& device, void* ptr);
+void ggml_vk_host_get(const vk_device& device, const void * ptr, vk_buffer& buf, size_t& buf_offset);
+void ggml_vk_ensure_sync_staging_buffer(vk_device& device, size_t size);
+void ggml_vk_ensure_sync_staging_buffer(ggml_backend_vk_context * ctx, size_t size);
+bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging = false);
+bool ggml_vk_buffer_write_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t size, bool sync_staging = false);
+void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height);
+void ggml_vk_buffer_write(vk_buffer& dst, size_t offset, const void * src, size_t size);
+bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging = false);
+void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height);
+void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_t size);
+void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size);
+void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size);
+void ggml_vk_buffer_memset_async(vk_context& ctx, vk_buffer& dst, size_t offset, uint32_t c, size_t size);
+void ggml_vk_buffer_memset(vk_buffer& dst, size_t offset, uint32_t c, size_t size);
+vk_buffer ggml_vk_buffer_from_host_ptr(vk_device & device, void * ptr, size_t size);
+
+// pipelines
+uint64_t vk_tensor_offset(const ggml_tensor * tensor);
+uint32_t get_misalign_bytes(const ggml_backend_vk_context * ctx, const ggml_tensor * t);
+void ggml_vk_wait_for_fence(ggml_backend_vk_context * ctx);
+void ggml_pipeline_request_descriptor_sets(ggml_backend_vk_context *ctx, vk_pipeline& pipeline, uint32_t n);
+void ggml_pipeline_allocate_descriptor_sets(ggml_backend_vk_context * ctx);
+void ggml_vk_submit(vk_context& ctx, vk::Fence fence);
+uint32_t ggml_vk_find_queue_family_index(std::vector<vk::QueueFamilyProperties>& queue_family_props, const vk::QueueFlags& required, const vk::QueueFlags& avoid, int32_t compute_index, uint32_t min_num_queues);
+std::unique_ptr<vk_queue> ggml_vk_create_queue(vk_device& device, uint32_t queue_family_index, uint32_t queue_index, vk::PipelineStageFlags&& stage_flags, bool transfer_only);
+std::unique_ptr<vk_queue> ggml_vk_create_aliased_queue(vk_device& device, const std::unique_ptr<vk_queue>& source);
+vk_context ggml_vk_create_context(ggml_backend_vk_context * ctx, vk_command_pool& p);
+vk_context ggml_vk_create_temporary_context(vk_command_pool& p);
+void ggml_vk_command_pool_cleanup(vk_device& device, vk_command_pool& p);
+void ggml_vk_queue_command_pools_cleanup(vk_device& device);
+vk_subbuffer ggml_vk_subbuffer(const ggml_backend_vk_context* ctx, const vk_buffer& buf, size_t offset = 0);
+void ggml_vk_sync_buffers(ggml_backend_vk_context* ctx, vk_context& subctx);
+void ggml_vk_set_event(vk_context& ctx, vk::Event& event);
+void ggml_vk_wait_events(vk_context& ctx, std::vector<vk::Event>&& events);
+vk_subbuffer ggml_vk_tensor_subbuffer(const ggml_backend_vk_context * ctx, const ggml_tensor * tensor, bool allow_misalign = false);
+void ggml_vk_cmd_label_begin(vk::CommandBuffer buf, const char * name);
+void ggml_vk_ctx_end(vk_context& ctx);
+void ggml_vk_ctx_begin(vk_device& device, vk_context& subctx);
+vk_context ggml_vk_get_compute_ctx(ggml_backend_vk_context * ctx);
+vk_context ggml_vk_get_transfer_ctx(ggml_backend_vk_context * ctx);
+bool ggml_vk_submit_transfer_ctx(ggml_backend_vk_context * ctx);
+size_t ggml_vk_align_size(size_t width, size_t align);
+void deferred_memcpy(void * dst, const void * src, size_t size, std::vector<vk_staging_memcpy>* memcpys = nullptr);
+void deferred_memset(void * dst, uint32_t val, size_t size, std::vector<vk_staging_memset>* memsets = nullptr);
+
+// matmul
+vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type type);
+void ggml_vk_matmul(ggml_backend_vk_context * ctx, vk_context& subctx, vk_pipeline& pipeline, vk_subbuffer&& a, vk_subbuffer&& b, vk_subbuffer&& d, vk_subbuffer&& split_k_buffer, uint32_t m, uint32_t n, uint32_t k, uint32_t stride_a, uint32_t stride_b, uint32_t stride_d, uint32_t batch_stride_a, uint32_t batch_stride_b, uint32_t batch_stride_d, uint32_t split_k, uint32_t batch, uint32_t ne02, uint32_t ne12, uint32_t broadcast2, uint32_t broadcast3, uint32_t padded_n);
+bool ggml_vk_dim01_contiguous(const ggml_tensor * tensor);
+vk_pipeline ggml_vk_get_cpy_pipeline(ggml_backend_vk_context * ctx, const ggml_tensor * src, const ggml_tensor * dst, ggml_type to);
+vk_pipeline ggml_vk_get_quantize_pipeline(ggml_backend_vk_context * ctx, ggml_type type);
+void ggml_vk_quantize_q8_1(ggml_backend_vk_context * ctx, vk_context& subctx, const vk_subbuffer & in, const vk_subbuffer & out, uint32_t ne);
+void ggml_vk_dsv4_hc_comb(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * mixes, const ggml_tensor * scale, const ggml_tensor * base, ggml_tensor * dst);
+void ggml_vk_dsv4_hc_pre(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * weights, ggml_tensor * dst);
+void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst);
+void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
+bool ggml_vk_use_mul_mat_vec_id(const struct ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_mul_mat_id(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
+
+// flash-attn
+void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * q, const ggml_tensor * k, const ggml_tensor * v, const ggml_tensor * mask, const ggml_tensor * sinks, ggml_tensor * dst);
+
+// operators
+void ggml_vk_cpy_to_contiguous(ggml_backend_vk_context * ctx, vk_context& subctx, vk_pipeline pipeline, const ggml_tensor * tensor, const vk_subbuffer & in, const vk_subbuffer & out);
+bool ggml_vk_can_use_fwht(const ggml_backend_vk_context * ctx, const ggml_tensor * src1, const ggml_tensor * dst);
+void ggml_vk_fwht(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src, ggml_tensor * dst);
+void ggml_vk_get_rows(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_get_rows_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_acc(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_multi_add(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_add(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_out_prod(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_sub(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_mul(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+int ggml_vk_unary_mul_op_index(ggml_unary_op op);
+void ggml_vk_unary_mul(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_div(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_add_id(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, ggml_tensor * dst);
+void ggml_vk_rwkv_wkv6(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_rwkv_wkv7(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_gated_linear_attn(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_lightning_indexer(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_gated_delta_net(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_ssm_scan(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_ssm_conv(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_opt_step_adamw(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_opt_step_sgd(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, ggml_tensor * dst);
+void ggml_vk_concat(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_upscale(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_scale(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_sqr(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_sqrt(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_add1(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_arange(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_fill(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_sin(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_cos(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_log(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_tri(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_diag(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_clamp(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_pad(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_pad_reflect_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_roll(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_repeat(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_repeat_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_cpy(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_set_rows(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_silu_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_norm(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_group_norm(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+uint32_t ggml_vk_rms_partials_size(ggml_backend_vk_context * ctx, const ggml_tensor *node);
+void ggml_vk_rms_norm(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx, float * op_params);
+void ggml_vk_rms_norm_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_l2_norm(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_unary(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_xielu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_glu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_diag_mask_inf(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_soft_max(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, ggml_tensor * dst);
+void ggml_vk_soft_max_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_topk_moe(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_rope(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_cgraph * cgraph, int node_idx, bool backprop);
+void ggml_vk_argsort(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_topk(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_topk_qsa(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_sum(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_sum_rows(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_mean(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_cumsum(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_cross_entropy_loss(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_cross_entropy_loss_back(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
+void ggml_vk_argmax(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_count_equal(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_solve_tri(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_im2col(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_im2col_3d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_timestep_embedding(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_conv_transpose_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_col2im_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_snake_dispatch_fused(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_pool_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_pool_2d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+void ggml_vk_conv_2d(ggml_backend_vk_context * ctx, vk_context & subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_conv_3d(ggml_backend_vk_context * ctx, vk_context & subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_conv_2d_dw(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
+void ggml_vk_leaky_relu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
+
+// graph
+void ggml_vk_preallocate_buffers(ggml_backend_vk_context * ctx, vk_context subctx);
+bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, int node_idx, ggml_tensor *node_begin, int node_idx_begin, bool last_node, bool almost_ready, bool submit);
+void ggml_vk_compute_forward(ggml_backend_vk_context* ctx, ggml_cgraph * cgraph, ggml_tensor* tensor, int tensor_idx, bool almost_ready);
+void ggml_vk_graph_cleanup(ggml_backend_vk_context * ctx);
+void ggml_vk_cleanup(ggml_backend_vk_context * ctx);
+void ggml_vk_synchronize(ggml_backend_vk_context * ctx);
+bool ggml_vk_is_empty(ggml_tensor * node);
+bool ggml_vk_can_fuse(const ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops);
+bool ggml_vk_can_fuse_ssm_conv(const ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx, int num_extra);
+bool ggml_vk_can_fuse_topk_moe(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx, topk_moe_mode mode);
+bool ggml_vk_can_fuse_topk_qsa(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+bool ggml_vk_can_fuse_rope_set_rows(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+bool ggml_vk_can_fuse_rms_norm_set_rows(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+bool ggml_vk_can_fuse_snake(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+bool ggml_vk_tensors_overlap(const ggml_tensor * a, const ggml_tensor * b, bool elementwise);
+bool ggml_vk_can_fuse_rms_norm_mul_rope(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+uint32_t ggml_vk_fuse_multi_add(ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx);
+void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * graph, struct ggml_backend_graph_optimize_params * params);
+
+// backend
+bool ggml_backend_buffer_is_vk(ggml_backend_buffer_t buffer);
+void ggml_backend_vk_buffer_free_buffer(ggml_backend_buffer_t buffer);
+void * ggml_backend_vk_buffer_get_base(ggml_backend_buffer_t buffer);
+enum ggml_status ggml_backend_vk_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor);
+void ggml_backend_vk_buffer_memset_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, uint8_t value, size_t offset, size_t size);
+void ggml_backend_vk_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size);
+void ggml_backend_vk_buffer_set_tensor_2d(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data);
+void ggml_backend_vk_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size);
+void ggml_backend_vk_buffer_get_tensor_2d(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data);
+bool ggml_backend_vk_buffer_cpy_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * src, ggml_tensor * dst);
+void ggml_backend_vk_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value);
+const char * ggml_backend_vk_buffer_type_name(ggml_backend_buffer_type_t buft);
+ggml_backend_buffer_t ggml_backend_vk_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size);
+size_t ggml_backend_vk_buffer_type_get_alignment(ggml_backend_buffer_type_t buft);
+size_t ggml_backend_vk_buffer_type_get_max_size(ggml_backend_buffer_type_t buft);
+size_t ggml_backend_vk_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor);
+void ggml_backend_vk_free(ggml_backend_t backend);
+ggml_backend_reg_t ggml_backend_vk_reg();
+
+// debug
+int64_t ggml_vk_get_op_batch_size(const ggml_tensor * op);
+
+// ggml-vulkan.cpp (residual)
+bool ggml_vk_lightning_indexer_k_type_supported(ggml_type type);
+void ggml_vk_print_device_fault_info(const vk_device& device);
+uint64_t ggml_vk_get_node_flops(const ggml_tensor * node);
+void ggml_vk_print_node_list(const ggml_cgraph * cgraph, int start, int end);
+void ggml_vk_print_device_lost_info(const vk_device& device);
+size_t ggml_vk_tensor_buffer_offset(const ggml_backend_vk_context * ctx, const ggml_tensor * t);
+size_t ggml_vk_descriptor_offset(size_t tensor_offset, size_t alignment, size_t type_size);
+uint32_t ggml_vk_concat_unit_size(ggml_type type);
+bool ggml_vk_concat_supported(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst);
+
+template <typename T>
+inline void ggml_vk_dispatch_pipeline(ggml_backend_vk_context* ctx, vk_context& subctx, vk_pipeline& pipeline, std::initializer_list<vk::DescriptorBufferInfo> const& descriptor_buffer_infos, const T &push_constants, std::array<uint32_t, 3> elements) {
+    const uint32_t wg0 = CEIL_DIV(elements[0], pipeline->wg_denoms[0]);
+    const uint32_t wg1 = CEIL_DIV(elements[1], pipeline->wg_denoms[1]);
+    const uint32_t wg2 = CEIL_DIV(elements[2], pipeline->wg_denoms[2]);
+    VK_LOG_DEBUG("ggml_vk_dispatch_pipeline(" << pipeline->name << ", {";
+    for (auto& buffer : descriptor_buffer_infos) {
+        std::cerr << "(" << buffer.buffer << ", " << buffer.offset << ", " << buffer.range << "), ";
+    }
+    std::cerr << "}, (" << wg0 << "," << wg1 << "," << wg2 << "))");
+    GGML_ASSERT(wg0 <= ctx->device->properties.limits.maxComputeWorkGroupCount[0] &&
+                wg1 <= ctx->device->properties.limits.maxComputeWorkGroupCount[1] &&
+                wg2 <= ctx->device->properties.limits.maxComputeWorkGroupCount[2]);
+    GGML_ASSERT(ctx->descriptor_set_idx < ctx->descriptor_sets.size());
+    GGML_ASSERT(descriptor_buffer_infos.size() <= MAX_PARAMETER_COUNT);
+    GGML_ASSERT(pipeline->parameter_count == descriptor_buffer_infos.size());
+    GGML_ASSERT(pipeline->push_constant_size == push_constant_size(push_constants));
+
+    vk::DescriptorSet& descriptor_set = ctx->descriptor_sets[ctx->descriptor_set_idx++];
+    vk::WriteDescriptorSet write_descriptor_set{ descriptor_set, 0, 0, pipeline->parameter_count, vk::DescriptorType::eStorageBuffer, nullptr, descriptor_buffer_infos.begin() };
+    ctx->device->device.updateDescriptorSets({ write_descriptor_set }, {});
+
+    subctx->s->buffer->buf.pushConstants(pipeline->layout, vk::ShaderStageFlagBits::eCompute, 0, push_constant_size(push_constants), push_constant_data(push_constants));
+    subctx->s->buffer->buf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline->pipeline);
+    subctx->s->buffer->buf.bindDescriptorSets(vk::PipelineBindPoint::eCompute,
+                                pipeline->layout,
+                                0,
+                                { descriptor_set },
+                                {});
+    {
+        ggml_vk_debug_label dbg(subctx, pipeline->name, wg0, wg1, wg2);
+        subctx->s->buffer->buf.dispatch(wg0, wg1, wg2);
+    }
+}
+
@@ -27,6 +27,8 @@ layout (push_constant) uniform parameter {
    uint inner_end;
 } p;

+shared int s;
+
 void argsort(bool needs_bounds_check, const uint row) {
    // bitonic sort
    int col = int(gl_GlobalInvocationID.x);
@@ -35,6 +37,12 @@ void argsort(bool needs_bounds_check, const uint row) {
    const uint row_offset = row * p.ncols;
    uint idx_offset = row * p.ncols_padded;

+    // workaround for NV driver/compiler bug - dummy use of shared memory
+    if (gl_LocalInvocationIndex == 0) {
+        s = 0;
+        col += s;
+    }
+
    bool need_barrier = false;

    // initialize indices
@@ -30,9 +30,14 @@ layout(local_size_x = BLOCK_SIZE, local_size_y = 1, local_size_z = 1) in;
 layout (binding = 0) readonly buffer A {uint data_a[];};
 layout (binding = 1) writeonly buffer D {uint data_d[];};

-shared uint vals[BLOCK_SIZE];
-shared uint offsets[BLOCK_SIZE];
-shared uint cursors[BLOCK_SIZE];
+// Upper bound on n_experts for the hoisted row-id path. Must match the limit in
+// ggml_vk_mul_mat_id_q_f16 (hoist_row_ids). The non-hoisted reduction below only
+// needs BLOCK_SIZE entries.
+#define MAX_EXPERTS 1024
+
+shared uint vals[MAX_EXPERTS];
+shared uint offsets[MAX_EXPERTS];
+shared uint cursors[MAX_EXPERTS];

 // data_d layout when p.hoist_row_ids is set:
 //   [0,              n_experts)   per-expert row count
@@ -46,8 +51,8 @@ void main() {
    const uint tid = gl_LocalInvocationID.x;

    if (p.hoist_row_ids != 0) {
-        if (tid < p.n_experts) {
-            vals[tid] = 0;
+        for (uint e = tid; e < p.n_experts; e += BLOCK_SIZE) {
+            vals[e] = 0;
        }
        barrier();

@@ -7,8 +7,13 @@
 //
 //   dst[i0, idst, it] = x[i0, it]*post[idst, it]
 //                     + sum_isrc residual[i0, isrc, it]*comb[idst, isrc, it]
+//
+// HAS_COMB == 0: identity mixing, each stream keeps its own residual:
+//
+//   dst[i0, idst, it] = x[i0, it]*post[idst, it] + residual[i0, idst, it]

 layout(constant_id = 0) const uint BLOCK_SIZE = 256;
+layout(constant_id = 1) const uint HAS_COMB = 1;

 layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;

@@ -48,7 +53,7 @@ void main() {
    if (tid < hc) {
        post_s[tid] = data_p[p_offset + tid * nbp0 + it * nbp1];
    }
-    if (tid < hc * hc) {
+    if (HAS_COMB == 1 && tid < hc * hc) {
        const uint idst = tid & 3;
        const uint isrc = tid >> 2;
        comb_s[tid] = data_c[c_offset + idst * nbc0 + isrc * nbc1 + it * nbc2];
@@ -74,9 +79,13 @@ void main() {
    [[unroll]]
    for (uint idst = 0; idst < hc; ++idst) {
        float result = xv * post_s[idst];
-        [[unroll]]
-        for (uint isrc = 0; isrc < hc; ++isrc) {
-            result = fma(r[isrc], comb_s[idst + hc * isrc], result);
+        if (HAS_COMB == 1) {
+            [[unroll]]
+            for (uint isrc = 0; isrc < hc; ++isrc) {
+                result = fma(r[isrc], comb_s[idst + hc * isrc], result);
+            }
+        } else {
+            result += r[idst];
        }
        data_d[d_offset + i0 * nbd0 + idst * nbd1 + it * nbd2] = result;
    }
@@ -4,9 +4,14 @@

 // Collapse the hc residual streams of a token into one, weighted per stream:
 //
-//   dst[i0, it] = sum_ih x[i0, ih, it] * weights[ih, it]
+//   dst[i0, it] = scale * sum_ih x[i0, ih, it] * weights[ih, it]
+//
+// GATED: weights is a per-element gate [n_embd, hc, n_tokens], applied as sigmoid:
+//
+//   dst[i0, it] = scale * sum_ih x[i0, ih, it] * sigmoid(gate[i0, ih, it])

 layout(constant_id = 0) const uint BLOCK_SIZE = 256;
+layout(constant_id = 1) const uint GATED = 0;

 layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;

@@ -16,12 +21,14 @@ layout(push_constant) uniform parameter
    uint n_tokens;

    uint nbx0; uint nbx1; uint nbx2;   // x
-    uint nbw0; uint nbw1;              // weights
+    uint nbw0; uint nbw1; uint nbw2;   // weights / gate
    uint nbd0; uint nbd1;              // dst

    uint x_offset;
    uint w_offset;
    uint d_offset;
+
+    float scale;
 };

 layout(binding = 0, std430) readonly buffer X { float data_x[]; };
@@ -36,10 +43,12 @@ void main() {
    const uint tid = gl_LocalInvocationID.x;
    const uint it  = gl_WorkGroupID.y;

-    if (tid < hc) {
-        w[tid] = data_w[w_offset + tid * nbw0 + it * nbw1];
+    if (GATED == 0) {
+        if (tid < hc) {
+            w[tid] = data_w[w_offset + tid * nbw0 + it * nbw1];
+        }
+        barrier();
    }
-    barrier();

    // After the barrier, so every invocation reaches it.
    const uint i0 = gl_WorkGroupID.x * BLOCK_SIZE + tid;
@@ -48,12 +57,19 @@ void main() {
    }

    const uint xb = x_offset + i0 * nbx0 + it * nbx2;
+    const uint wb = w_offset + i0 * nbw0 + it * nbw2;

    float result = 0.0f;
    [[unroll]]
    for (uint ih = 0; ih < hc; ++ih) {
-        result = fma(data_x[xb + ih * nbx1], w[ih], result);
+        float wv;
+        if (GATED == 1) {
+            wv = 1.0f / (1.0f + exp(-data_w[wb + ih * nbw1]));
+        } else {
+            wv = w[ih];
+        }
+        result = fma(data_x[xb + ih * nbx1], wv, result);
    }

-    data_d[d_offset + i0 * nbd0 + it * nbd1] = result;
+    data_d[d_offset + i0 * nbd0 + it * nbd1] = scale * result;
 }
@@ -31,7 +31,7 @@ layout (binding = 0) readonly buffer X {A_TYPE data_a[];};
 layout (binding = 1) writeonly buffer D {D_TYPE data_d[];};

 #if BDA
-layout (buffer_reference) buffer D_ptr {D_TYPE d;};
+layout (buffer_reference, buffer_reference_align = D_SIZE) buffer D_ptr {D_TYPE d;};
 #endif

 void im2col(const uint ow, const uint z_idx) {
@@ -50,7 +50,7 @@ layout (binding = 0) readonly buffer X {A_TYPE data_a[];};
 layout (binding = 1) writeonly buffer D {D_TYPE data_d[];};

 #if BDA
-layout (buffer_reference) buffer D_ptr {D_TYPE d;};
+layout (buffer_reference, buffer_reference_align = D_SIZE) buffer D_ptr {D_TYPE d;};
 #endif

 void main() {
@@ -313,6 +313,9 @@ void main() {

    // Workgroup has no work
    if (ic * BN >= _ne1) return;
+
+    uint required_work_items = (_ne1 - ic * BN) * BK / LOAD_VEC_B_EFF / LOAD_VEC_BATCH_B;
+    uint required_warp_c = (_ne1 - ic * BN + WN - 1) / WN;
 #endif

 #ifdef MUL_MAT_ID
@@ -363,6 +366,9 @@ void main() {
        [[unroll]] for (uint l = 0; l < BM; l += loadstride_a) {
            load_a_to_shmem(pos_a, loadr_a, loadc_a + l, ir * BM + loadc_a + l, block, end_k);
        }
+        #ifdef MUL_MAT_ID
+        if (gl_LocalInvocationID.x < required_work_items) {
+        #endif
        [[unroll]] for (uint l = 0; l < BN; l += loadstride_b) {
 #if !defined(MUL_MAT_ID)
            load_b_to_shmem(pos_b, loadr_b, loadc_b + l, ic * BN + loadc_b + l, block, end_k);
@@ -370,6 +376,9 @@ void main() {
            load_b_to_shmem(pos_b, loadr_b, loadc_b + l, ic, _ne1, block, end_k);
 #endif
        }
+        #ifdef MUL_MAT_ID
+        }
+        #endif

        barrier();

@@ -377,6 +386,9 @@ void main() {
        pos_b += BK / LOAD_VEC_B_EFF;

 #ifdef COOPMAT
+#ifdef MUL_MAT_ID
+        if (warp_c < required_warp_c) {
+#endif
        [[unroll]] for (uint i = 0; i < BK; i += TK) {
            [[unroll]] for (uint cm_row = 0; cm_row < cms_per_row; cm_row++) {
                // Load from shared into cache
@@ -389,6 +401,9 @@ void main() {
                }
            }
        }
+#ifdef MUL_MAT_ID
+        }
+#endif
 #else
        [[unroll]] for (uint i = 0; i < BK / BK_STEP; i++) {
            // Load from shared into cache
@@ -238,6 +238,7 @@ struct gguf_reader {
        : callback(callback),
          userdata(userdata),
          max_chunk_read(max_chunk_read),
+          start_offset(data_offset),
          data_offset(data_offset),
          nbytes_remain(nbytes_remain) {
        GGML_ASSERT(max_chunk_read > 0);
@@ -366,6 +367,11 @@ struct gguf_reader {
        return data_offset;
    }

+    // position in the file where the GGUF data starts, alignment is relative to it, not to the file
+    uint64_t start() const {
+        return start_offset;
+    }
+
    bool seek(uint64_t absolute_offset) const {
        const uint64_t end_offset = uint64_t(data_offset) + nbytes_remain;
        if (absolute_offset > end_offset) {
@@ -415,6 +421,7 @@ private:
    gguf_reader_callback_t callback = nullptr;
    void * userdata = nullptr;
    size_t max_chunk_read = 0;
+    uint64_t start_offset = 0;
    mutable uint64_t data_offset = 0;
    mutable uint64_t nbytes_remain = 0;
 };
@@ -763,7 +770,7 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr
    GGML_ASSERT(int64_t(ctx->info.size()) == n_tensors);

    // we require the data section to be aligned, so take into account any padding
-    if (n_tensors > 0 && !gr.seek(GGML_PAD(gr.tell(), ctx->alignment))) {
+    if (n_tensors > 0 && !gr.seek(gr.start() + GGML_PAD(gr.tell() - gr.start(), ctx->alignment))) {
        GGML_LOG_ERROR("%s: failed to seek to beginning of data section\n", __func__);
        gguf_free(ctx);
        return nullptr;
@@ -5876,7 +5876,7 @@ GGML_QUANT_SIZES: dict[GGMLQuantizationType, tuple[int, int]] = {
    GGMLQuantizationType.Q5_0:    (32, 2 + 4 + 16),
    GGMLQuantizationType.Q5_1:    (32, 2 + 2 + 4 + 16),
    GGMLQuantizationType.Q8_0:    (32, 2 + 32),
-    GGMLQuantizationType.Q8_1:    (32, 4 + 4 + 32),
+    GGMLQuantizationType.Q8_1:    (32, 2 + 2 + 32),
    GGMLQuantizationType.Q2_K:    (256, 2 + 2 + QK_K // 16 + QK_K // 4),
    GGMLQuantizationType.Q3_K:    (256, 2 + QK_K // 4 + QK_K // 8 + 12),
    GGMLQuantizationType.Q4_K:    (256, 2 + 2 + QK_K // 2 + 12),
@@ -518,6 +518,8 @@ extern "C" {
              struct llama_model_params   params);

    // Load a model from an open FILE pointer
+    // The GGUF is read from the current position, so it can be embedded in a larger file
+    // mmap needs the GGUF data section at a file offset to be aligned to the CPU tensor alignment (32 bytes)
    LLAMA_API struct llama_model * llama_model_load_from_file_ptr(
                                   FILE * file,
              struct llama_model_params   params);
@@ -681,6 +683,11 @@ extern "C" {
            struct llama_model * model,
            const char * path_lora);

+    // Load a LoRA adapter from an open FILE pointer, reading from its current position
+    LLAMA_API struct llama_adapter_lora * llama_adapter_lora_init_from_file_ptr(
+            struct llama_model * model,
+            FILE * file);
+
    // Functions to access the adapter's GGUF metadata scalar values
    // - The functions return the length of the string on success, or -1 on failure
    // - The output string is always null-terminated and cleared on failure
@@ -8,7 +8,7 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR})

 if (EMSCRIPTEN)
 else()
-    if (NOT GGML_BACKEND_DL)
+    if (NOT GGML_BACKEND_DL AND GGML_CPU)
        add_subdirectory(vdot)
    endif()
 endif()
@@ -55,7 +55,9 @@ fi
 # Some generic functions
 usage() {
    echo "Usage: $0 [ --include-path <dir> ] --generate <build-dir> libXXX [ libYYY ... ]" >&2
-    echo "       $0 --check <old-build-dir> <new-build-dir>" >&2
+    echo "       $0 [ --strict ] --check <old-build-dir> <new-build-dir>" >&2
+    echo "" >&2
+    echo "  --strict: fail on any API/ABI change, including backwards-compatible additions" >&2
 }

 get_cmake_project_name() {
@@ -72,6 +74,7 @@ get_cmake_version() {
 # Option parsing and validation
 DO_GEN=0
 DO_CHECK=0
+STRICT=0
 BUILD_DIR=
 BUILD_DIR_NEW=
 INCLUDE_PATHS=
@@ -119,6 +122,10 @@ while [ "$#" -gt 0 ]; do
        INCLUDE_PATHS="$INCLUDE_PATHS $2"
        shift 2
        ;;
+    --strict)
+        STRICT=1
+        shift
+        ;;
    -h | --help)
        usage
        exit 1
@@ -258,13 +265,24 @@ elif [ "$DO_CHECK" -eq 1 ]; then
        abidiff "$xml_file" "$xml_file_new"
        res=$?
        [ "$((res & 8))" -ne 0 ] && ABI_RESULT=1
+        # check bit 2 for compatible ABI changes (like new symbols) and if
+        # STRICT is set then handle this as an error.
+        [ "$STRICT" -eq 1 ] && [ "$((res & 4))" -ne 0 ] && ABI_RESULT=1
    done

    if [ "$API_RESULT" -gt 0 ]; then
-        echo "ERROR: API changed with possible backwards-compatibility problems." >&2
+        if [ "$STRICT" -eq 1 ]; then
+            echo "ERROR: API changed — a minor version bump is required." >&2
+        else
+            echo "ERROR: API changed with backwards-incompatible problems — a major version bump is required." >&2
+        fi
    fi
    if [ "$ABI_RESULT" -gt 0 ]; then
-        echo "ERROR: ABI changed with possible backwards-compatibility problems." >&2
+        if [ "$STRICT" -eq 1 ]; then
+            echo "ERROR: ABI changed — a minor version bump is required." >&2
+        else
+            echo "ERROR: ABI changed with backwards-incompatible problems — a major version bump is required." >&2
+        fi
    fi
    if [ "$((API_RESULT + ABI_RESULT))" -gt 0 ]; then
        exit 1
@@ -0,0 +1,131 @@
+#!/bin/bash
+# Check API/ABI compatibility between the previous release tag and current HEAD.
+#
+# Finds the most recent vX.Y.Z tag, checks it out in a temporary git worktree,
+# builds both versions with shared libs enabled, and uses check-apiabi-compat.sh
+# to compare the results.
+#
+# Exit codes:
+#   0: compatible, or check was skipped
+#   1: backwards-incompatible changes found, or build failed
+#
+# Options:
+#   --tag <version>: compare against this tag instead of the latest release
+#
+# Environment:
+#   SKIP_APIABI_CHECK:    set to 1 or true to skip
+#   APIABI_COMPARE_TAG:   equivalent to --tag (used by CI)
+
+set -euo pipefail
+
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)"
+
+usage() {
+    echo "Usage: $0 [--tag <version>]" >&2
+    echo "  --tag <version>  Compare against this release tag (default: latest)" >&2
+}
+
+COMPARE_TAG="${APIABI_COMPARE_TAG:-}"
+while [[ "$#" -gt 0 ]]; do
+    case "$1" in
+    --tag)
+        if [[ -z "${2:-}" ]]; then usage; exit 1; fi
+        COMPARE_TAG="$2"
+        shift 2
+        ;;
+    --tag=*)
+        COMPARE_TAG="${1#*=}"
+        shift
+        ;;
+    -h | --help)
+        usage; exit 0
+        ;;
+    *)
+        usage; exit 1
+        ;;
+    esac
+done
+
+if [[ "${SKIP_APIABI_CHECK:-}" == "1" || "${SKIP_APIABI_CHECK:-}" == "true" ]]; then
+    echo "SKIP_APIABI_CHECK is set - skipping API/ABI compatibility check"
+    exit 0
+fi
+
+if ! command -v abi-compliance-checker >/dev/null 2>&1 || ! command -v abidw >/dev/null 2>&1; then
+    echo "Warning: abi-compliance-checker or abigail-tools not installed - skipping API/ABI check"
+    exit 0
+fi
+
+discover_libs() {
+    local build_dir="$1"
+    local libs=()
+    for dir in "$build_dir/src" "$build_dir/bin"; do
+        [[ -d "$dir" ]] || continue
+        for f in "$dir"/lib*.so; do
+            [[ -f "$f" ]] && libs+=("$(basename "$f" .so)")
+        done
+    done
+    echo "${libs[@]}"
+}
+
+if [[ -n "${COMPARE_TAG}" ]]; then
+    PREV_TAG="${COMPARE_TAG}"
+    if ! git -C "$REPO_ROOT" rev-parse --verify "${PREV_TAG}^{}" >/dev/null 2>&1; then
+        echo "Error: tag '${PREV_TAG}' not found in repository." >&2
+        exit 1
+    fi
+else
+    PREV_TAG=$(git -C "$REPO_ROOT" tag --sort=-v:refname | grep -E '^v[0-9]+\.[0-9]+\.[0-9]+$' | head -n 1 || true)
+    if [[ -z "${PREV_TAG}" ]]; then
+        echo "Warning: no previous release tag found - skipping API/ABI check"
+        exit 0
+    fi
+fi
+OLD_VERSION="${PREV_TAG#v}"
+OLD_MAJOR="${OLD_VERSION%%.*}"
+OLD_MINOR="${OLD_VERSION#*.}"; OLD_MINOR="${OLD_MINOR%%.*}"
+
+NEW_MAJOR=$(grep "set(LLAMA_VERSION_MAJOR" "$REPO_ROOT/CMakeLists.txt" | sed 's/.*MAJOR \([0-9]*\).*/\1/')
+NEW_MINOR=$(grep "set(LLAMA_VERSION_MINOR" "$REPO_ROOT/CMakeLists.txt" | sed 's/.*MINOR \([0-9]*\).*/\1/')
+
+if [[ "$NEW_MAJOR" -gt "$OLD_MAJOR" ]]; then
+    echo "Major version increment ($OLD_MAJOR -> $NEW_MAJOR): API/ABI breaking changes are expected, skipping compatibility check."
+    exit 0
+fi
+
+CHECK_FLAGS=()
+if [[ "$NEW_MINOR" -eq "$OLD_MINOR" ]]; then
+    echo "Patch version bump detected: checking for any API/ABI changes (a minor bump is required if any are found)..."
+    CHECK_FLAGS+=(--strict)
+else
+    echo "Minor version bump detected: checking for backwards-incompatible API/ABI changes..."
+fi
+
+echo "Checking API/ABI compatibility against ${PREV_TAG}..."
+
+WORKTREE_DIR=$(mktemp -d)
+BUILD_OLD=$(mktemp -d)
+BUILD_NEW=$(mktemp -d)
+
+cleanup() {
+    git -C "$REPO_ROOT" worktree remove --force "$WORKTREE_DIR" 2>/dev/null || true
+    rm -rf "$WORKTREE_DIR" "$BUILD_OLD" "$BUILD_NEW"
+}
+trap cleanup EXIT
+
+git -C "$REPO_ROOT" worktree add "$WORKTREE_DIR" "$PREV_TAG"
+
+cmake -S "$WORKTREE_DIR" -B "$BUILD_OLD" -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=RelWithDebInfo
+cmake --build "$BUILD_OLD" --parallel "$(nproc)"
+OLD_LIBS=($(discover_libs "$BUILD_OLD"))
+echo "Libraries found in old build: ${OLD_LIBS[*]}"
+
+cmake -S "$REPO_ROOT"    -B "$BUILD_NEW" -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=RelWithDebInfo
+cmake --build "$BUILD_NEW" --parallel "$(nproc)"
+NEW_LIBS=($(discover_libs "$BUILD_NEW"))
+echo "Libraries found in new build: ${NEW_LIBS[*]}"
+
+(cd "$WORKTREE_DIR" && "$SCRIPT_DIR/check-apiabi-compat.sh" --include-path ggml/include --generate "$BUILD_OLD" "${OLD_LIBS[@]}")
+(cd "$REPO_ROOT"    && "$SCRIPT_DIR/check-apiabi-compat.sh" --include-path ggml/include --generate "$BUILD_NEW" "${NEW_LIBS[@]}")
+(cd "$REPO_ROOT"    && "$SCRIPT_DIR/check-apiabi-compat.sh" "${CHECK_FLAGS[@]}" --check "$BUILD_OLD" "$BUILD_NEW")
@@ -165,6 +165,23 @@ else
    fi
 fi

+echo "Checking API/ABI compatibility..."
+set +e
+bash "$SCRIPT_DIR/check-release-apiabi.sh"
+APIABI_RESULT=$?
+set -e
+if [[ $APIABI_RESULT -ne 0 ]]; then
+    if [[ "$DRY_RUN" == "true" ]]; then
+        echo "Warning: API/ABI check found backwards-incompatible changes (dry run, continuing)."
+        CHECKS_PASSED=false
+    else
+        echo "Error: API/ABI check found backwards-incompatible changes."
+        exit 1
+    fi
+else
+    echo "API/ABI compatibility check passed - OK"
+fi
+
 if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
    echo "checks_passed=${CHECKS_PASSED}" >> "$GITHUB_OUTPUT"
 fi
@@ -6,6 +6,8 @@

 #include <map>
 #include <cassert>
+#include <cerrno>
+#include <cstring>
 #include <sstream>
 #include <stdexcept>

@@ -146,22 +148,23 @@ llama_adapter_lora_weight * llama_adapter_lora::get_weight(ggml_tensor * w) {
    return nullptr;
 }

-static void llama_adapter_lora_init_impl(llama_model & model, const char * path_lora, llama_adapter_lora & adapter) {
-    LLAMA_LOG_INFO("%s: loading lora adapter from '%s' ...\n", __func__, path_lora);
-
+static void llama_adapter_lora_init_impl(llama_model & model, FILE * file, llama_adapter_lora & adapter) {
    ggml_context * ctx_init;
    gguf_init_params meta_gguf_params = {
        /* .no_alloc = */ true,
        /* .ctx      = */ &ctx_init,
    };

-    gguf_context_ptr ctx_gguf { gguf_init_from_file(path_lora, meta_gguf_params) };
+    gguf_context_ptr ctx_gguf { gguf_init_from_file_ptr(file, meta_gguf_params) };
    if (!ctx_gguf) {
-        throw std::runtime_error("failed to load lora adapter file from " + std::string(path_lora));
+        throw std::runtime_error("failed to load lora adapter from file");
    }

    ggml_context_ptr ctx { ctx_init };

+    // must come after gguf_init_from_file_ptr, the llama_file constructor moves the file position
+    llama_file gguf_file(file);
+
    // check metadata
    {
        const gguf_context * gguf_ctx = ctx_gguf.get();
@@ -393,7 +396,6 @@ static void llama_adapter_lora_init_impl(llama_model & model, const char * path_

    // set tensor data
    {
-        llama_file gguf_file(path_lora, "rb");
        std::vector<uint8_t> read_buf;
        auto set_tensor = [&](ggml_tensor * orig, ggml_tensor * dev) {
            const size_t offs = gguf_get_data_offset(ctx_gguf.get()) + gguf_get_tensor_offset(ctx_gguf.get(), gguf_find_tensor(ctx_gguf.get(), orig->name));
@@ -421,10 +423,30 @@ static void llama_adapter_lora_init_impl(llama_model & model, const char * path_
 }

 llama_adapter_lora * llama_adapter_lora_init(llama_model * model, const char * path_lora) {
+    LLAMA_LOG_INFO("%s: loading lora adapter from '%s' ...\n", __func__, path_lora);
+
+    FILE * file = ggml_fopen(path_lora, "rb");
+    if (!file) {
+        LLAMA_LOG_ERROR("%s: failed to open '%s': %s\n", __func__, path_lora, strerror(errno));
+        return nullptr;
+    }
+
+    llama_adapter_lora * adapter = llama_adapter_lora_init_from_file_ptr(model, file);
+    fclose(file);
+
+    return adapter;
+}
+
+llama_adapter_lora * llama_adapter_lora_init_from_file_ptr(llama_model * model, FILE * file) {
+    if (!file) {
+        LLAMA_LOG_ERROR("%s: file is NULL\n", __func__);
+        return nullptr;
+    }
+
    llama_adapter_lora * adapter = new llama_adapter_lora(model);

    try {
-        llama_adapter_lora_init_impl(*model, path_lora, *adapter);
+        llama_adapter_lora_init_impl(*model, file, *adapter);
        return adapter;
    } catch (const std::exception & err) {
        LLAMA_LOG_ERROR("%s: failed to apply lora adapter: %s\n", __func__, err.what());
@@ -599,8 +599,11 @@ void llama_context::sched_reserve() {

    LLAMA_LOG_DEBUG("%s: max_nodes = %zu\n", __func__, max_nodes);

-    gf_res_prev.reset(new llm_graph_result(max_nodes));
+    for (auto & res : gf_res_prev) {
+        res.reset();
+    }
    gf_res_reserve.reset(new llm_graph_result(max_nodes));
+    gf_res_prev_active = nullptr;

    sched.reset(ggml_backend_sched_new(backend_ptrs.data(), backend_buft.data(), backend_ptrs.size(), max_nodes, cparams.pipeline_parallel, cparams.op_offload));

@@ -816,10 +819,14 @@ bool llama_context::memory_update(bool optimize) {
                }
        }

-        // reset the previous graph result to make sure that it won't be reused
-        // TODO: change the mctx->apply() to return information if a graph reserve is needed
-        //       reset the graph result only if the memory module did reset the scheduler
-        gf_res_prev->reset();
+        // reset the previous graph results to make sure that they won't be reused
+        // TODO: make mctx->apply() report if a graph reserve is needed, then reset graph results only if the memory module reset the scheduler
+        for (auto & res : gf_res_prev) {
+            if (res) {
+                res->reset();
+            }
+        }
+        gf_res_prev_active = nullptr;

        if (!mctx->apply()) {
            LLAMA_LOG_ERROR("%s: failed to apply memory update\n", __func__);
@@ -1340,14 +1347,14 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll
        return nullptr;
    }

-    auto * res = gf_res_prev.get();
+    auto * res = get_gf_res_prev();
    auto * gf  = res->get_gf();

    // the new graph parameters
    // in order to correctly reuse a graph, it's full topology has to be uniquely determined by these parameters
    const auto gparams = graph_params(res, ubatch, mctx, gtype);

-    if (!graph_reuse_disable && res->can_reuse(gparams)) {
+    if (!graph_reuse_disable && gf_res_prev_active == res && res->can_reuse(gparams)) {
        //LLAMA_LOG_DEBUG("%s: reusing previous graph\n", __func__);

        // with pipeline parallelism, the previous graph_compute_async may still be running
@@ -1359,6 +1366,7 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll

        n_reused++;
    } else {
+        gf_res_prev_active = nullptr;
        res->reset();

        ggml_backend_sched_reset(sched.get());
@@ -1381,6 +1389,8 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll
            ret = GGML_STATUS_ALLOC_FAILED;
            return nullptr;
        }
+
+        gf_res_prev_active = res;
    }

    // set the input data for the input tensors
@@ -2361,6 +2371,14 @@ llm_graph_result * llama_context::get_gf_res_reserve() const {
    return static_cast<llm_graph_result *>(gf_res_reserve.get());
 }

+llm_graph_result * llama_context::get_gf_res_prev() {
+    auto & res = gf_res_prev[n_outputs > 0];
+    if (!res) {
+        res.reset(new llm_graph_result(gf_res_reserve->get_max_nodes()));
+    }
+    return res.get();
+}
+
 // pack sampler outputs into as few sequences as possible before using sequences without samplers
 static void ubatch_prepare_reserve(
              llama_ubatch                            & ubatch,
@@ -2430,8 +2448,13 @@ ggml_cgraph * llama_context::graph_reserve(

    ggml_backend_sched_reset(sched.get());

-    // when the scheduler is reset, we cannot reuse the old graph, so we reset the previous graph result to prevent that
-    gf_res_prev->reset();
+    // when the scheduler is reset, we cannot reuse old graphs, so we reset the previous graph results
+    for (auto & res : gf_res_prev) {
+        if (res) {
+            res->reset();
+        }
+    }
+    gf_res_prev_active = nullptr;

    // store the n_outputs as it is, and restore it afterwards
    // TODO: not sure if needed, might simplify in the future by removing this
@@ -3521,10 +3544,12 @@ void llama_context::opt_epoch_iter(
                break;
            }

-            auto * res = gf_res_prev.get();
+            auto * res = get_gf_res_prev();

            const auto gparams = graph_params(res, ubatch, mctx.get(), ctx_type_to_graph_type(cparams.ctx_type));

+            // the optimizer graph is allocated outside sched, so the next decode must rebuild
+            gf_res_prev_active = nullptr;
            res->reset();

            auto * gf = model.build_graph(gparams);
@@ -11,6 +11,7 @@
 #include "ggml-cpp.h"
 #include "ggml-opt.h"

+#include <array>
 #include <map>
 #include <vector>

@@ -254,6 +255,8 @@ public:
    bool set_sampler(llama_seq_id seq_id, llama_sampler * sampler);

 private:
+    llm_graph_result * get_gf_res_prev();
+
    llm_graph_params graph_params(
                        llm_graph_result * res,
                      const llama_ubatch & ubatch,
@@ -364,9 +367,12 @@ private:
    std::vector<ggml_backend_buffer_type_t> backend_buft;
    std::vector<size_t>                     backend_buf_exp_size; // expected buffer sizes

-    llm_graph_result_ptr gf_res_prev;
+    // Separate arenas give batches with and without outputs distinct CUDA graph cache keys.
+    std::array<llm_graph_result_ptr, 2> gf_res_prev;
    llm_graph_result_ptr gf_res_reserve;

+    llm_graph_result * gf_res_prev_active = nullptr;
+
    // host buffer for the model output (logits and embeddings)
    ggml_backend_buffer_ptr buf_output;

@@ -685,6 +685,13 @@ llama_model_loader::llama_model_loader(
            throw std::runtime_error(format("%s: failed to load model from file pointer", __func__));
        }

+        // mmap places tensors at their file offsets, so an embedded GGUF must be aligned in the file too
+        const size_t tensor_align = ggml_backend_buft_get_alignment(ggml_backend_cpu_buffer_type());
+        if (use_mmap && gguf_get_data_offset(metadata) % tensor_align != 0) {
+            throw std::runtime_error(format("%s: GGUF data section at file offset %zu is not %zu byte aligned, cannot mmap",
+                __func__, gguf_get_data_offset(metadata), tensor_align));
+        }
+
        get_key(llm_kv(LLM_KV_GENERAL_ARCHITECTURE), arch_name, false);
        llm_kv = LLM_KV(llm_arch_from_string(arch_name));

@@ -603,8 +603,20 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
    };

    auto get_split_segments = [&](int axis, uint32_t il) -> std::vector<std::pair<int64_t, uint32_t>> {
+        // TODO: clarify why this is necessary specifically for these models
+        // TODO: deduplicate condition [TAG_SPLIT_QGATE_QWEN]
        if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE ||
                ud->model->arch == LLM_ARCH_QWEN4EXP) {
+
+            // fused full attention layers with Q gate tensors that need n_embd doubled:
+            if (!hparams.is_recr(il) && (std::regex_match(tensor_name, pattern_qkv_weight) || std::regex_match(tensor_name, pattern_qkv_bias))) {
+                const int64_t n_embd      = hparams.n_head(il) * hparams.n_embd_head_k(il) * 2;
+                const int64_t n_embd_gqa  = hparams.n_embd_v_gqa(il);
+                GGML_ASSERT(hparams.n_embd_k_gqa(il) == n_embd_gqa);
+                GGML_ASSERT(tensor->ne[axis] == n_embd + 2*n_embd_gqa);
+                return {{n_embd, 1}, {n_embd_gqa, 2}};
+            }
+
            const int64_t head_k_dim = hparams.ssm_d_state;
            const int64_t head_v_dim = hparams.ssm_d_state;
            const int64_t n_k_heads  = hparams.ssm_n_group;
@@ -654,9 +666,9 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
        }

        if (std::regex_match(tensor_name, pattern_qkv_weight) || std::regex_match(tensor_name, pattern_qkv_bias)) {
-            const int64_t n_embd      = hparams.n_embd;
+            const int64_t n_embd      = hparams.n_head(il) * hparams.n_embd_head_k(il);
            const int64_t n_embd_gqa  = hparams.n_embd_v_gqa(il);
-            GGML_ASSERT(hparams.n_embd_k_gqa() == n_embd_gqa);
+            GGML_ASSERT(hparams.n_embd_k_gqa(il) == n_embd_gqa);
            GGML_ASSERT(tensor->ne[axis] == n_embd + 2*n_embd_gqa);
            return {{n_embd, 1}, {n_embd_gqa, 2}};
        }
@@ -742,6 +754,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
            if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_q_bias)) {
                GGML_ASSERT(segments.size() == 1);
                // some models have Q gate tensors, for those cases the granularity needs to be doubled:
+                // TODO: deduplicate condition [TAG_SPLIT_QGATE_QWEN]
                if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE ||
                        ud->model->arch == LLM_ARCH_QWEN4EXP) {
                    return {std::lcm(2*n_embd_q, blck_size_perf)};
@@ -769,6 +782,12 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
            }
            if (std::regex_match(tensor_name, pattern_qkv_weight) || std::regex_match(tensor_name, pattern_qkv_bias)) {
                GGML_ASSERT(segments.size() == 2);
+                // fused full attention layers need Q gate tensors handled like above:
+                // TODO: deduplicate condition [TAG_SPLIT_QGATE_QWEN]
+                if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE ||
+                        ud->model->arch == LLM_ARCH_QWEN4EXP) {
+                    return {std::lcm(2*n_embd_q, blck_size_perf), granularity_kv};
+                }
                return {granularity_q, granularity_kv};
            }
        }
@@ -3232,6 +3251,15 @@ ggml_tensor * llama_model_base::create_tensor(const LLM_TN_IMPL & tn, const std:
 }

 void llama_model_base::create_tensor_gate_up_exps(llama_layer & layer, int bid, int64_t n_embd_, int64_t n_ff_, int64_t n_expert_, int flags) {
+    if (flags & TENSOR_SKIP) {
+        const int skip = TENSOR_NOT_REQUIRED | TENSOR_SKIP;
+
+        create_tensor(tn(LLM_TENSOR_FFN_GATE_UP_EXPS, "weight", bid), {n_embd_, n_ff_ * 2, n_expert_}, skip | TENSOR_SKIP_IF_VIRTUAL);
+        create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS,    "weight", bid), {n_embd_, n_ff_,     n_expert_}, skip);
+        create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,      "weight", bid), {n_embd_, n_ff_,     n_expert_}, skip);
+        return;
+    }
+
    layer.ffn_gate_up_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_UP_EXPS, "weight", bid), {n_embd_, n_ff_ * 2, n_expert_}, TENSOR_NOT_REQUIRED);
    if (layer.ffn_gate_up_exps == nullptr) {
        layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", bid), {n_embd_, n_ff_, n_expert_}, flags);
@@ -100,6 +100,18 @@ llama_model_nemotron_h_moe::graph_mtp::graph_mtp(const llama_model & model, cons
        ggml_tensor * router_logits = build_lora_mm(layer.ffn_gate_inp, cur);
        cb(router_logits, "mtp_ffn_moe_logits", il);

+        ggml_tensor * ffn_shexp = build_ffn(cur,
+                layer.ffn_up_shexp,   NULL, layer.ffn_up_shexp_s,
+                NULL,                 NULL, NULL,
+                layer.ffn_down_shexp, NULL, layer.ffn_down_shexp_s,
+                NULL,
+                LLM_FFN_RELU_SQR, LLM_FFN_PAR, il);
+        cb(ffn_shexp, "mtp_ffn_shexp", il);
+
+        if (layer.ffn_latent_down) {
+            cur = ggml_mul_mat(ctx0, layer.ffn_latent_down, cur);
+        }
+
        ggml_tensor * moe_out =
            build_moe_ffn(cur,
                layer.ffn_gate_inp,
@@ -118,13 +130,9 @@ llama_model_nemotron_h_moe::graph_mtp::graph_mtp(const llama_model & model, cons
                layer.ffn_down_exps_s);
        cb(moe_out, "mtp_ffn_moe_out", il);

-        ggml_tensor * ffn_shexp = build_ffn(cur,
-                layer.ffn_up_shexp,   NULL, layer.ffn_up_shexp_s,
-                NULL,                 NULL, NULL,
-                layer.ffn_down_shexp, NULL, layer.ffn_down_shexp_s,
-                NULL,
-                LLM_FFN_RELU_SQR, LLM_FFN_PAR, il);
-        cb(ffn_shexp, "mtp_ffn_shexp", il);
+        if (layer.ffn_latent_up) {
+            moe_out = ggml_mul_mat(ctx0, layer.ffn_latent_up, moe_out);
+        }

        cur = ggml_add(ctx0, moe_out, ffn_shexp);
        cb(cur, "mtp_ffn_out", il);
@@ -172,6 +172,8 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
        layer.attn_post_norm  = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM,  "weight", i), {n_embd}, mtp_flags);
        layer.ffn_gate_inp    = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,    "weight", i), {n_embd, n_expert}, mtp_flags);
        layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias",   i), {n_expert}, mtp_flags);
+        layer.ffn_latent_down = create_tensor(tn(LLM_TENSOR_FFN_LATENT_DOWN, "weight", i), {n_embd, moe_n_embd}, mtp_flags | TENSOR_NOT_REQUIRED);
+        layer.ffn_latent_up   = create_tensor(tn(LLM_TENSOR_FFN_LATENT_UP,   "weight", i), {moe_n_embd, n_embd}, mtp_flags | TENSOR_NOT_REQUIRED);
        layer.ffn_down_exps   = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS,   "weight", i), {n_ff_exp,   moe_n_embd, n_expert}, mtp_flags);
        layer.ffn_up_exps     = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,     "weight", i), {moe_n_embd, n_ff_exp,   n_expert}, mtp_flags);
        layer.ffn_down_shexp  = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP,  "weight", i), {n_ff_shexp, n_embd}, mtp_flags);
@@ -322,7 +322,7 @@ if (APPLE)
    llama_build(test-rset-release.cpp)
 endif()

-if (NOT GGML_BACKEND_DL)
+if (NOT GGML_BACKEND_DL AND GGML_CPU)
    # these tests use the backends directly and cannot be built with dynamic loading
    llama_build_and_test(test-barrier.cpp)
    llama_build_and_test(test-quantize-fns.cpp)
@@ -10077,6 +10077,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {

    // gpt-oss issue with Vulkan mmq_id
    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_MXFP4, GGML_TYPE_F32, 32, 2, false, 2880, 32, 2880));
+    // more than 256 experts (hoisted row-id path): 512 as in Qwen3.8-Flash-Next,
+    // and 1024 at the LLAMA_MAX_EXPERTS limit
+    for (int n : {1, 5, 64, 300}) {
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_IQ3_S,  GGML_TYPE_F32,  512, 10, false, 128, n, 512));
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_0,   GGML_TYPE_F32,  512, 10, false, 256, n, 128));
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_IQ3_S,  GGML_TYPE_F32, 1024, 10, false, 128, n, 512));
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_0,   GGML_TYPE_F32, 1024, 10, false, 256, n, 128));
+    }
    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_0, GGML_TYPE_F32, 32, 2, false, 2880, 32, 2880));

    // multiple blocks per row: exercises the block-stride loop and the
@@ -1167,15 +1167,17 @@ static bool same_tensor_data(const struct ggml_context * orig, const struct ggml

 enum roundtrip_read_mode {
    ROUNDTRIP_READ_MODE_FILE,
+    ROUNDTRIP_READ_MODE_FILE_OFFSET, // GGUF embedded after some bytes of a bigger file
    ROUNDTRIP_READ_MODE_BUFFER,
    ROUNDTRIP_READ_MODE_CALLBACK,
 };

 static const char * roundtrip_read_mode_name(const roundtrip_read_mode mode) {
    switch (mode) {
-        case ROUNDTRIP_READ_MODE_FILE:     return "file";
-        case ROUNDTRIP_READ_MODE_BUFFER:   return "buffer";
-        case ROUNDTRIP_READ_MODE_CALLBACK: return "callback";
+        case ROUNDTRIP_READ_MODE_FILE:        return "file";
+        case ROUNDTRIP_READ_MODE_FILE_OFFSET: return "file_offset";
+        case ROUNDTRIP_READ_MODE_BUFFER:      return "buffer";
+        case ROUNDTRIP_READ_MODE_CALLBACK:    return "callback";
    }

    GGML_ABORT("fatal error");
@@ -1214,6 +1216,12 @@ static std::pair<int, int> test_roundtrip(
    GGML_ASSERT(file);
 #endif // _WIN32

+    // not a multiple of any alignment, so the data section padding must be relative to the GGUF start
+    const long prefix = read_mode == ROUNDTRIP_READ_MODE_FILE_OFFSET ? 7 : 0;
+    for (long i = 0; i < prefix; ++i) {
+        fputc(0xAB, file);
+    }
+
    gguf_write_to_file_ptr(gguf_ctx_0, file, only_meta);
    rewind(file);

@@ -1236,6 +1244,7 @@ static std::pair<int, int> test_roundtrip(
        };
        gguf_ctx_1 = gguf_init_from_callback(read_buffer_callback, &reader, 4096, 4ull << 30 /* 4GB */, gguf_params);
    } else {
+        GGML_ASSERT(fseek(file, prefix, SEEK_SET) == 0);
        gguf_ctx_1 = gguf_init_from_file_ptr(file, gguf_params);
    }

@@ -1451,6 +1460,11 @@ int main(int argc, char ** argv) {
            npass += result.first;
            ntest += result.second;
        }
+        {
+            std::pair<int, int> result = test_roundtrip(dev, seed, /*only_meta=*/false, ROUNDTRIP_READ_MODE_FILE_OFFSET);
+            npass += result.first;
+            ntest += result.second;
+        }
        {
            std::pair<int, int> result = test_roundtrip(dev, seed, /*only_meta=*/false, ROUNDTRIP_READ_MODE_BUFFER);
            npass += result.first;
@@ -1,3 +1,4 @@
+#include "ggml-backend.h"
 #include "ggml-rpc.h"
 #ifdef _WIN32
 #  define NOMINMAX
@@ -270,7 +271,8 @@ static std::vector<ggml_backend_dev_t> get_devices(const rpc_server_params & par
    if (devices.empty()) {
        for (size_t i = 0; i < ggml_backend_dev_count(); i++) {
            ggml_backend_dev_t dev = ggml_backend_dev_get(i);
-            if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) {
+            enum ggml_backend_dev_type dev_type = ggml_backend_dev_type(dev);
+            if (dev_type != GGML_BACKEND_DEVICE_TYPE_CPU && dev_type != GGML_BACKEND_DEVICE_TYPE_ACCEL) {
                devices.push_back(dev);
            }
        }
@@ -1,6 +1,10 @@
 <script lang="ts">
 	import { File, Image, MessageSquare, Mic, Plus, Video } from '@lucide/svelte';
-	import { ChatFormActionAddToolsSubmenu, McpLogo } from '$lib/components/app';
+	import {
+		ChatFormActionAddReasoningSubmenu,
+		ChatFormActionAddToolsSubmenu,
+		McpLogo
+	} from '$lib/components/app';
 	import { buttonVariants } from '$lib/components/ui/button';
 	import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
 	import * as Tooltip from '$lib/components/ui/tooltip';
@@ -13,6 +17,7 @@
 	import { getChatFormActionsContext } from '$lib/contexts';
 	import { AttachmentAction, AttachmentItemEnabledWhen } from '$lib/enums';
 	import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
+	import { serverStore } from '$lib/stores';

 	interface Props {
 		class?: string;
@@ -92,6 +97,13 @@
 				}
 			}}
 		>
+			<!-- in router mode the models selector owns the reasoning submenu -->
+			{#if !serverStore.isRouterMode}
+				<ChatFormActionAddReasoningSubmenu />
+
+				<DropdownMenu.Separator />
+			{/if}
+
 			<DropdownMenu.Item
 				class="flex cursor-pointer items-center gap-2"
 				onclick={() => attachmentMenu.callbacks[AttachmentAction.FILE_UPLOAD]()}
@@ -45,7 +45,7 @@
 		await screen.findByRole('menu');

 		await waitFor(() => {
-			expect(document.activeElement).toHaveTextContent('Add files');
+			expect(document.activeElement).toHaveTextContent('Reasoning');
 		});
 	}}
 />
Author	SHA1	Message	Date
Nikita Gordeev	8ed1a55efc	cmake : fix build when GGML_CPU=OFF and GGML_CUDA=ON (#29026 ) * fix: build fails when GGML_CPU=OFF and GGML_CUDA=ON * fix: eol in examples/convert-llama2c-to-ggml/CMakeLists.txt file	2026-09-18 11:44:03 +03:00
yanghongandSigbjørn Skjæret	bb11ebb682	gguf-py: fix Q8_1 block size in GGML_QUANT_SIZES (2+2+32) (#29036 ) * gguf-py: fix Q8_1 block size in GGML_QUANT_SIZES * --whitespace --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>	2026-09-18 11:43:06 +03:00
Sigbjørn Skjæret	f03cf3e9b8	ci : disable GHA cache for copilot (#29068 )	2026-09-18 10:20:53 +02:00
Sigbjørn Skjæret	bdcbaaf6e7	ci : bump android-actions/setup-android to 4.0.4 (#29065 )	2026-09-18 09:04:14 +02:00
drluoto	5c53396b89	vulkan: raise the hoisted row-id limit for mul_mat_id from 256 to 512 experts (#28501 ) * vulkan: raise the hoisted row-id limit for mul_mat_id to 512 experts The expert-count shader (count_experts.comp) sizes its shared arrays with BLOCK_SIZE, which is 256. Because of that, row-id hoisting is switched off for any model with more than 256 experts, and every mul_mat_id workgroup has to rescan the whole ids tensor on its own. Qwen3.8-Flash-Next has 512 experts and was quietly running on that slow path. This change sizes the arrays with a separate MAX_EXPERTS constant (512), clears them in a loop instead of one entry per thread, and raises the matching limit on the host side. On Strix Halo at batch 2048 the expert matmuls drop from 12.5 to 9.5 ms (iq3_s) and from 14.0 to 7.5 ms (iq4_nl) per op, and prompt processing gets about 19 % faster at 8k tokens. test-backend-ops MUL_MAT_ID passes (891/891) with new 512-expert test cases. Assisted-by: Claude Fable 5.1 * vulkan: raise the hoisted row-id limit for mul_mat_id to 1024 experts Follow-up to review feedback: 1024 matches LLAMA_MAX_EXPERTS instead of stopping at 512. The three shared arrays in count_experts.comp grow to 3 * 1024 * 4 = 12 KiB, which fits the 16 KiB that Vulkan guarantees for maxComputeSharedMemorySize. Adds mul_mat_id test cases at 1024 experts alongside the existing 512 ones. test-backend-ops MUL_MAT_ID passes on Vulkan (RADV, Strix Halo, Radeon 8060S): 889/889.	2026-09-18 09:00:15 +02:00
Sigbjørn Skjæret	972d2313bc	ci : add missing evict-old-files (#29041 )	2026-09-17 19:05:40 +02:00
Pedro Cuenca	c77ae695c9	rpc : skip ACCEL devices (#29020 )	2026-09-17 18:36:53 +03:00
David Friehs	b49650adb3	model : skip gate_up_exps if TENSOR_SKIP is set (#29014 ) required for qwen35moe if MTP tensors are fused but not loaded	2026-09-17 13:53:54 +02:00
Kartik Gulia	7076180486	model : extend Nemotron MTP support (#29018 ) * first fix * removed unnecessary declarations	2026-09-17 13:52:36 +02:00
Ravi PanchumarthyandMostafa Faheem	ebbb185227	openvino : Update OpenVINO to 2026.4;fix clangd,MSVC warnings; (#29009 ) * Update to openvino-2026.4 * Update OV docs * ggml-openvino : fix clangd and MSVC warnings * fix int to ptr cast, more internal linkage enforcement, and avoiding duplicate switch case --------- Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>	2026-09-17 12:46:14 +02:00
Benjamin Babik Ben Babik Pascal	4ff829ec2e	ui: fix removed reasoning menu in single model mode on desktop (#27985 ) * ui: fix accidentally removed reasoning menu in single model mode on desktop * ui: formatting task run to fix storybook test * ui: mount the add menu reasoning submenu outside router mode only The models selector already owns the reasoning submenu in router mode, so the add menu only mounts it in single model mode. The first enabled item of the add menu is now the reasoning submenu, the accessibility story expects it. --------- Co-authored-by: Ben Babik <work@benjaminbabik.com> Co-authored-by: Pascal <admin@serveurperso.com>	2026-09-17 11:18:13 +02:00
Ruben Ortlam	f172be756a	vulkan: split buffers and debug code into separate files, add shared headers (#28732 )	2026-09-17 11:16:35 +03:00
Daniel Bevenius	87f9c82f2d	ci : add API/ABI check to make-release workflow [no ci] (#28947 ) * ci : add API/ABI check to make-release workflow [no ci] This commit adds an API/ABI compatibility check to the make-release workflow. The motivation for this to allow us to detect any potential breaking changes in API/ABI compatibility between releases and fail the the release if there are any. The workflow can be triggered manually as before and this check can be skipped if needed as it does take some time which might be useful when doing a dry-run and not specifically interested in the API/ABI check. By default this will check the current release against the latest release, but this can also be configured in the workflow, or in the script run on the command line, to check a different tag. * add check for minor version bumps [no ci] This commit also changes the build type to be RelWithDebInfo so that the reported information is more useful.	2026-09-17 09:55:16 +02:00
midagedevandSigbjørn Skjæret	7f6f0c2a9d	chat : add message delimiters to the DeepSeek V3.2/V4 parser (#29008 ) * chat : add message delimiters to the DeepSeek V3.2/V4 parser Assisted-by: Claude Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>	2026-09-17 09:41:08 +02:00
Yuri KhrustalevandJohannes Gäßler	81aeaeb74b	gguf : align the data section relative to the GGUF start, not the file (#28993 ) * gguf : align the data section relative to the GGUF start, not the file gguf_init_from_file_ptr reads a GGUF from the current file position, but padded the data section from file offset 0, so a GGUF embedded at an offset that is not a multiple of the alignment loaded without error and returned wrong tensor data. Also adds llama_adapter_lora_init_from_file_ptr, and disables mmap with a warning when an embedded data section is not aligned, instead of asserting in ggml. Assisted-by: Claude Opus 5 * llama : load lora from path through the FILE* variant The test now checks that mmap is disabled only for an unaligned offset. Assisted-by: Claude Fable 5.1 * Update ggml/src/gguf.cpp Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * Update include/llama.h Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * llama : error on unaligned mmap of an embedded GGUF, drop test-load-file-ptr --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de>	2026-09-17 09:19:44 +02:00
Neo Zhang	c9a5eeeb34	sycl : fix the B70 mem allocate error when >19.3GB (#28953 )	2026-09-17 09:56:03 +03:00
Jiang, Fish	7490357f22	vulkan: skip unneeded MoE work in mul_mm coopmat1 path (#25483 )	2026-09-17 09:53:07 +03:00
Titaniumtown	817e5f83eb	sycl: ssm_conv: fuse the SiLU epilogue into the ssm_conv kernel (#28929 )	2026-09-17 09:51:50 +03:00
lhez	c57da6fd81	opencl: fix various warnings (#28984 ) * opencl: fix warnings * opencl: fix warnings for non adreno	2026-09-17 09:48:05 +03:00
Johannes Gäßler	79bfc1d43a	docs: remove JG as CODEOWNER for test-llama-archs (#29003 )	2026-09-17 08:14:58 +02:00
Abir Deol	05f2dcfdba	vulkan: fix buffer_reference alignment in im2col shaders (#28996 ) Both im2col.comp and im2col_3d.comp declare D_ptr without an explicit buffer_reference_align, so glslang emits writes through it as Aligned 16. The shaders advance the pointer by D_SIZE, a per-variant define set to 4 for float and 2 for float16_t, so most write addresses are not 16-byte aligned. This triggers VUID-RuntimeSpirv-PhysicalStorageBuffer64-06315 under GPU-AV. Declaring buffer_reference_align = D_SIZE matches the alignment to the actual write stride and takes validation hits from 20 to 0 for both IM2COL and IM2COL_3D. Fixes #28960	2026-09-17 07:14:33 +02:00
Ruben Ortlam	35822afe58	vulkan: support qwen4exp hc ops (#28988 ) * vulkan: support qwen4exp hc ops * fix stale comment [no-ci]	2026-09-17 06:34:23 +02:00
Michael Taylor	aa39d7a3e1	[SYCL] Fix function signature for `ggml_backend_sycl_split_buffer_type` (#28981 )	2026-09-16 22:15:47 -04:00
Jeff Bolz	4bc272fd72	vulkan: work around NV bug with argsort_large.comp (#28975 )	2026-09-16 18:40:07 -05:00
David Friehs	fb27a525d2	TP: fix split state and granularity for fused QKV gemma4, qwen35 (#28965 ) * model: calculate split states for attn_qkv from n_head * n_embd_head_k required for gemma4 with --fuse-qkv, where n_embd is 5376 but Q is 8192. * model: handle fused full attention layers for qwen35/qwen35moe * model: add TODO: [TAG_SPLIT_QGATE_QWEN]	2026-09-16 22:02:12 +03:00
Eve	c6824a9e42	ci: switch fast jobs back to github (#28959 ) * switch jobs to ubuntu-slim * ubuntu slim almost takes 15 minutes for check requirements so use something faster	2026-09-16 16:50:26 +00:00
Gaurav Garg	2f3fd02526	Enable CUDA graph for MTP draft (#28549 ) * Improve CUDA graph usage for MTP * Rename field * Address review feedback	2026-09-16 21:46:54 +05:30