mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-10-30 08:27:34 -05:00
Compare commits
19
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Kartik Gulia
|
7076180486 | ||
|
Ravi PanchumarthyandMostafa Faheem
|
ebbb185227 | ||
|
|
4ff829ec2e | ||
|
Ruben Ortlam
|
f172be756a | ||
|
Daniel Bevenius
|
87f9c82f2d | ||
|
midagedevandSigbjørn Skjæret
|
7f6f0c2a9d | ||
|
Yuri KhrustalevandJohannes Gäßler
|
81aeaeb74b | ||
|
Neo Zhang
|
c9a5eeeb34 | ||
|
Jiang, Fish
|
7490357f22 | ||
|
Titaniumtown
|
817e5f83eb | ||
|
lhez
|
c57da6fd81 | ||
|
Johannes Gäßler
|
79bfc1d43a | ||
|
Abir Deol
|
05f2dcfdba | ||
|
Ruben Ortlam
|
35822afe58 | ||
|
Michael Taylor
|
aa39d7a3e1 | ||
|
Jeff Bolz
|
4bc272fd72 | ||
|
David Friehs
|
fb27a525d2 | ||
|
Eve
|
c6824a9e42 | ||
|
Gaurav Garg
|
2f3fd02526 |
@@ -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; \
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -555,8 +555,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
|
||||
@@ -669,8 +669,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
|
||||
|
||||
@@ -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(
|
||||
|
||||
+14
-12
@@ -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>
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
+461
-473
File diff suppressed because it is too large
Load Diff
@@ -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;
|
||||
|
||||
+264
-39
@@ -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);
|
||||
}
|
||||
}
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
+455
-5406
File diff suppressed because it is too large
Load Diff
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
+8
-1
@@ -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;
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
Executable
+131
@@ -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
|
||||
|
||||
+29
-7
@@ -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());
|
||||
|
||||
+35
-10
@@ -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);
|
||||
|
||||
+7
-1
@@ -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));
|
||||
|
||||
|
||||
+21
-2
@@ -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};
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
|
||||
+17
-3
@@ -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;
|
||||
|
||||
+13
-1
@@ -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');
|
||||
});
|
||||
}}
|
||||
/>
|
||||
|
||||
Reference in New Issue
Block a user