Compare commits

..

8 Commits

Author SHA1 Message Date
forforever73 5522498343 metal : port new kernels into the split sources
Move the kernels added on master after the split (lightning indexer,
DSv4 hyper-connections, silu_back, f16 bin ops) into the corresponding
kernels/*.metal sources. Copied verbatim, no functional change.
2026-08-05 16:22:32 +08:00
Georgi Gerganov db0a8ce844 ggml-metal: FWHT kernel for metal backend (#25924) 2026-08-05 16:22:32 +08:00
Georgi Gerganov b2d79a7e7f metal: fuse snake activation (mul, sin, sqr, mul, add) (#25459) 2026-08-05 16:22:32 +08:00
Georgi Gerganov 0b6cbb572c metal : add Q2_0 support (#25419) 2026-08-05 16:22:32 +08:00
Georgi Gerganov 7649856ea2 metal : add CONV_2D_DW (depthwise convolution) support (#21565) 2026-08-05 16:22:32 +08:00
Georgi Gerganov f348d51391 metal : add set_rows with src0 f16 (#25434) 2026-08-05 16:22:32 +08:00
Georgi Gerganov d2112aa30e metal: add col2im_1d op (f32/f16/bf16) (#25176) 2026-08-05 16:22:32 +08:00
YiChen Lv ff8169ce01 metal : per-op source split + parallel compile (#24021)
* preliminary extract common header

* op source split

* split metallib into 8 libs && load in parallel

* derive kernel->library routing from functionNames

* x-macro lib list + underscore filenames, dedup QK_NL, MRC fixes

* op source split 8 to 20

* improve robustness of source fallback

* clean up

* change bool -> atomic_bool

* only prepend headers that source actually includes

* no semaphore, use GCD global queue

* dedup library compile path, fix NSError lifetime, rename gla

* relocate upstream concat/rope_back/repeat kernel changes into split files

* move ggml-common.h from common.h into dequantize.h to shrink binary size

---------

Co-authored-by: lvyichen <lvyichen@stepfun.com>
2026-08-05 16:22:32 +08:00
928 changed files with 24965 additions and 66465 deletions
+1
View File
@@ -57,6 +57,7 @@ COPY --from=web /app/tools/ui/dist tools/ui/dist
RUN HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
cmake -S . -B build \
-DGGML_HIP=ON \
-DGGML_HIP_ROCWMMA_FATTN=ON \
-DAMDGPU_TARGETS="$ROCM_DOCKER_ARCH" \
-DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON \
-DCMAKE_BUILD_TYPE=Release -DLLAMA_BUILD_TESTS=OFF \
@@ -4,10 +4,6 @@ inputs:
cuda_version:
description: "CUDA toolkit version"
required: true
cuda_arch:
description: "CUDA target architecture"
required: false
default: "x64"
runs:
using: "composite"
@@ -131,26 +127,3 @@ runs:
echo "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3\bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
echo "CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3" | Out-File -FilePath $env:GITHUB_ENV -Append -Encoding utf8
echo "CUDA_PATH_V13_3=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3" | Out-File -FilePath $env:GITHUB_ENV -Append -Encoding utf8
- name: Install Cuda Toolkit 13.4 for ARM64
if: ${{ inputs.cuda_version == '13.4' && inputs.cuda_arch == 'arm64' }}
shell: pwsh
run: |
mkdir -p "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4"
choco install unzip -y
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-x86_64/5B515474-7E78-11F1-8656-C51E4F4B317F/cccl-windows-x86_64-13.3.4.1.2-archive.zip"
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-x86_64/5B515474-7E78-11F1-8656-C51E4F4B317F/cuda_crt-windows-x86_64-13.4.46-archive.zip"
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-x86_64/5B515474-7E78-11F1-8656-C51E4F4B317F/cuda_nvcc-windows-x86_64-13.4.46-archive.zip"
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-x86_64/5B515474-7E78-11F1-8656-C51E4F4B317F/libnvvm-windows-x86_64-13.4.46-archive.zip"
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-arm64/5B515474-7E78-11F1-8656-C51E4F4B317F/cuda_cudart-windows-arm64-13.4.46-archive.zip"
curl -O "https://packages.nvidia.com/bin-archive/pool/windows-arm64/5B515474-7E78-11F1-8656-C51E4F4B317F/libcublas-windows-arm64-13.7.0.10-archive.zip"
unzip '*.zip' -d "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4"
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\cccl-windows-x86_64-13.3.4.1.2-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\cuda_crt-windows-x86_64-13.4.46-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\cuda_nvcc-windows-x86_64-13.4.46-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\libnvvm-windows-x86_64-13.4.46-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\cuda_cudart-windows-arm64-13.4.46-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
xcopy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\libcublas-windows-arm64-13.7.0.10-archive\*" "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" /E /I /H /Y
echo "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4\bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
echo "CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" | Out-File -FilePath $env:GITHUB_ENV -Append -Encoding utf8
echo "CUDA_PATH_V13_4=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.4" | Out-File -FilePath $env:GITHUB_ENV -Append -Encoding utf8
+5 -23
View File
@@ -8,26 +8,8 @@ inputs:
runs:
using: "composite"
steps:
- name: Install ROCm with Wheels
shell: pwsh
run: |
$ErrorActionPreference = "Stop"
write-host "Setting up Python virtual environment"
# Create the venv directly at the cache location to avoid relocation issues
New-Item -Path "C:\TheRock\build" -ItemType Directory -Force | Out-Null
python -m venv C:\TheRock\build\.venv
& C:\TheRock\build\.venv\Scripts\Activate.ps1
write-host "Upgrading pip"
python -m pip install --upgrade pip
write-host "Installing ROCm wheels for multi-arch support"
# Install ROCm wheels for multi-arch support (this may take several minutes)
python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ inputs.version }}"
# Pre-expand the devel tree so it is included in the cache
write-host "Initializing ROCm devel tree"
rocm-sdk init
if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
write-host "Completed ROCm wheel installation to C:\TheRock\build"
- name: Setup ROCm
uses: ./.github/actions/install-exe
with:
url: https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ inputs.version }}-Win11-For-HIP.exe
args: -install
+5
View File
@@ -60,6 +60,7 @@ jobs:
-DCMAKE_BUILD_RPATH="@loader_path" \
-DLLAMA_FATAL_WARNINGS=ON \
-DLLAMA_BUILD_BORINGSSL=ON \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=OFF \
-DGGML_METAL_SHADER_DEBUG=ON \
-DGGML_RPC=ON \
@@ -126,6 +127,7 @@ jobs:
run: |
sysctl -a
cmake -B build -G Xcode \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=ON \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_APP=OFF \
@@ -176,6 +178,7 @@ jobs:
run: |
sysctl -a
cmake -B build -G Xcode \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=ON \
-DLLAMA_BUILD_COMMON=OFF \
-DLLAMA_BUILD_APP=OFF \
@@ -209,6 +212,7 @@ jobs:
run: |
sysctl -a
cmake -B build -G Xcode \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=ON \
-DLLAMA_BUILD_COMMON=OFF \
-DLLAMA_BUILD_APP=OFF \
@@ -253,6 +257,7 @@ jobs:
run: |
sysctl -a
cmake -B build -G Xcode \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=ON \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_APP=OFF \
+20 -20
View File
@@ -119,27 +119,27 @@ jobs:
version_major: ${{ env.OPENVINO_VERSION_MAJOR }}
version_full: ${{ env.OPENVINO_VERSION_FULL }}
# windows-2022-rocm-cache:
# runs-on: windows-2022
windows-2022-rocm-cache:
runs-on: windows-2022
# env:
# # Make sure this is in sync with release.yml and build-cuda-windows.yml
# ROCM_VERSION: "7.14.0"
env:
# Make sure this is in sync with build.yml
HIPSDK_INSTALLER_VERSION: "26.Q1"
# steps:
# - name: Clone
# id: checkout
# uses: actions/checkout@v6
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
# - name: Setup Cache
# uses: actions/cache@v5
# id: cache-rocm
# with:
# path: C:\TheRock\build
# key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
- name: Setup Cache
uses: actions/cache@v5
id: cache-rocm
with:
path: C:\Program Files\AMD\ROCm
key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
# - name: Setup ROCm
# if: steps.cache-rocm.outputs.cache-hit != 'true'
# uses: ./.github/actions/windows-setup-rocm
# with:
# version: ${{ env.ROCM_VERSION }}
- name: Setup ROCm
if: steps.cache-rocm.outputs.cache-hit != 'true'
uses: ./.github/actions/windows-setup-rocm
with:
version: ${{ env.HIPSDK_INSTALLER_VERSION }}
+4 -10
View File
@@ -5,7 +5,7 @@ on:
jobs:
linux:
runs-on: [self-hosted, Linux]
runs-on: [self-hosted, Linux, CPU]
steps:
- uses: actions/checkout@v6
with:
@@ -21,21 +21,15 @@ jobs:
-DLLAMA_BUILD_TOOLS=OFF \
-DLLAMA_BUILD_EXAMPLES=OFF \
-DLLAMA_BUILD_APP=OFF \
-DLLAMA_BUILD_IS_DEV=OFF \
-DCMAKE_BUILD_TYPE=Release
cmake --build build --config Release -j $(nproc)
cmake --build build --config Release
cmake --install build --prefix "$PREFIX" --config Release
export LLAMA_CONFIG="$PREFIX"/lib/cmake/llama/llama-config.cmake
tclsh <<'EOF'
set build(commit) [string trim [exec git rev-parse --short HEAD]]
set build(number) [string trim [exec git rev-list --count HEAD]]
set cmakelists [read [open "CMakeLists.txt" r]]
regexp {set\(LLAMA_VERSION_MAJOR\s+(\d+)\)} $cmakelists -> major
regexp {set\(LLAMA_VERSION_MINOR\s+(\d+)\)} $cmakelists -> minor
regexp {set\(LLAMA_VERSION_PATCH\s+(\d+)\)} $cmakelists -> patch
set build(version) "$major.$minor.$patch"
set build(version) "0.0.$build(number)"
set llamaconfig [read [open "$env(LLAMA_CONFIG)" r]]
set checks [list "set\\(LLAMA_VERSION \\s+$build(version)\\)" \
@@ -54,4 +48,4 @@ jobs:
cd examples/simple-cmake-pkg
cmake -S . -B build -DCMAKE_PREFIX_PATH="$PREFIX"/lib/cmake
cmake --build build -j $(nproc)
cmake --build build
+1 -3
View File
@@ -94,10 +94,8 @@ jobs:
id: cmake_build
run: |
cmake -B build \
-DGGML_NATIVE=OFF \
-DLLAMA_FATAL_WARNINGS=ON \
-DGGML_RPC=ON \
-DGGML_NATIVE=OFF
-DGGML_RPC=ON
time cmake --build build --config Release -j $(nproc)
- name: Test
+1
View File
@@ -99,6 +99,7 @@ jobs:
run: |
cmake -B build -S . \
-DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
-DGGML_HIP_ROCWMMA_FATTN=ON \
-DGPU_TARGETS="gfx1030" \
-DGGML_HIP=ON
cmake --build build --config Release -j $(nproc)
+35 -50
View File
@@ -83,7 +83,7 @@ jobs:
env:
# Make sure this is in sync with build-cache.yml
ROCM_VERSION: "7.14.0"
HIPSDK_INSTALLER_VERSION: "26.Q1"
strategy:
matrix:
@@ -97,53 +97,36 @@ jobs:
id: checkout
uses: actions/checkout@v6
# - name: Cache ROCm Installation
# uses: actions/cache@v5
# id: cache-rocm
# with:
# path: C:\TheRock\build
# key: rocm-wheels-${{ env.ROCM_VERSION }}-multi-arch-${{ runner.os }}
- name: Grab rocWMMA package
id: grab_rocwmma
run: |
curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2.1/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70201-81~24.04_amd64.deb"
7z x rocwmma.deb
7z x data.tar
- name: Use ROCm Installation Cache
uses: actions/cache@v5
id: cache-rocm
with:
path: C:\Program Files\AMD\ROCm
key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
- name: Setup ROCm
# if: steps.cache-rocm.outputs.cache-hit != 'true'
if: steps.cache-rocm.outputs.cache-hit != 'true'
uses: ./.github/actions/windows-setup-rocm
with:
version: ${{ env.ROCM_VERSION }}
- name: Setup ROCm Environment
run: |
$ErrorActionPreference = "Stop"
# Activate venv from cache or fresh install
& C:\TheRock\build\.venv\Scripts\Activate.ps1
# Expand the devel tree (idempotent; no-op if already done during install)
rocm-sdk init
if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
# Get ROCm installation paths using the rocm-sdk CLI tool
$rocmPath = (rocm-sdk path --root)
if (-not $rocmPath) { throw "rocm-sdk path --root returned empty - devel package may not be installed" }
$rocmPath = $rocmPath.Trim()
$cmakePath = (rocm-sdk path --cmake).Trim()
$binPath = (rocm-sdk path --bin).Trim()
write-host "ROCm root: $rocmPath"
echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
echo "CMAKE_PREFIX_PATH=$cmakePath" >> $env:GITHUB_ENV
echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
echo "$binPath" >> $env:GITHUB_PATH
# Keep venv in PATH for subsequent steps
echo "C:\TheRock\build\.venv\Scripts" >> $env:GITHUB_PATH
version: ${{ env.HIPSDK_INSTALLER_VERSION }}
- name: Verify ROCm
id: verify
run: |
# Test the ROCm clang shipped in the installed wheel
& "${env:HIP_PATH}\lib\llvm\bin\clang.exe" --version
# Find and test ROCm installation
$clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1
if (-not $clangPath) {
Write-Error "ROCm installation not found"
exit 1
}
& $clangPath.FullName --version
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
@@ -151,27 +134,29 @@ jobs:
# TODO: this build does not match the build in release.yml, so we use a different cache key
# ideally, the builds should match, similar to the CUDA build above so that we would be able
# to populate the ccache for the release with manual runs of this workflow
#key: release-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
key: cuda-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
#key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
key: cuda-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
- name: Build
id: cmake_build
run: |
$env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path)
$env:CMAKE_PREFIX_PATH="${env:HIP_PATH}"
cmake -G "Unix Makefiles" -B build -S . `
-DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
-DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
-DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
-DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" `
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" `
-DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.1/include/" `
-DCMAKE_BUILD_TYPE=Release `
-DLLAMA_BUILD_BORINGSSL=ON `
-DHIP_PATH="${env:HIP_PATH}" `
-DROCM_DIR="${env:HIP_PATH}" `
-DGGML_HIP=ON `
-DGPU_TARGETS="gfx1100" `
-DGGML_HIP_ROCWMMA_FATTN=ON `
-DGPU_TARGETS="gfx1100" `
-DGGML_RPC=ON
cmake --build build -j ${env:NUMBER_OF_PROCESSORS}
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
#key: release-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
key: cuda-windows-2022-x64-hip-${{ env.ROCM_VERSION }}-${{ matrix.name }}
#key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
key: cuda-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
+3 -25
View File
@@ -15,12 +15,6 @@ on:
'**/*.cpp'
]
pull_request:
types: [opened, synchronize, reopened]
paths: [
'.github/workflows/build-sanitize.yml'
]
concurrency:
group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
cancel-in-progress: true
@@ -34,35 +28,19 @@ env:
jobs:
ctest:
runs-on: [self-hosted, X64, CPU, Linux]
continue-on-error: true
strategy:
matrix:
include:
# thread and address doesn't run properly on some self hosted machines, so run it on Github instead
- sanitizer: ADDRESS
machine: ubuntu-24.04
- sanitizer: THREAD
machine: ubuntu-24.04
- sanitizer: UNDEFINED
machine: [self-hosted, X64, Linux]
runs-on: ${{ matrix.machine }}
sanitizer: [ADDRESS, THREAD, UNDEFINED]
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
# - name: ccache
# uses: ggml-org/ccache-action@v1.2.21
# if: ${{ matrix.sanitizer != 'UNDEFINED' }}
# with:
# key: ctest-${{ matrix.sanitizer }}-ubuntu-24.04
# variant: ccache
# evict-old-files: 1d
# save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
# with UNDEFINED sanitizer, we have to build in Debug to avoid GCC 13 false-positive warnings
- name: Build (undefined)
id: cmake_build_undefined
-20
View File
@@ -71,26 +71,6 @@ jobs:
nvidia-smi
GG_BUILD_CUDA=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
gpu-rocm:
runs-on: [self-hosted, Linux, AMD]
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
- name: Test
id: ggml-ci
# HIP_LAUNCH_BLOCKING=1: workaround for an async-execution correctness
# issue on integrated RDNA3.5 (gfx1151) where batched inference returns
# incorrect output (perplexity ~88 vs ~9.4). Serializing kernel launches
# restores correctness. Remove once the underlying ROCm/HIP issue is fixed.
env:
HIP_LAUNCH_BLOCKING: "1"
run: |
rocminfo
GG_BUILD_ROCM=1 GG_BUILD_AMDGPU_TARGETS=gfx1151 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
gpu-vulkan-nvidia-cm:
runs-on: [self-hosted, Linux, NVIDIA]
-46
View File
@@ -1,46 +0,0 @@
name: Make Release
on:
workflow_dispatch:
inputs:
dry_run:
description: 'Dry run - validate without creating the tag'
required: true
type: boolean
default: true
env:
GH_TOKEN: ${{ github.token }}
permissions:
contents: write
jobs:
make-release:
runs-on: ubuntu-latest
steps:
- name: Checkout
uses: actions/checkout@v6
- name: Run release checks
id: checks
run: bash scripts/make-release-checks.sh ${{ github.event.inputs.dry_run == 'true' && '--dry-run' || '' }}
env:
GITHUB_REPOSITORY: ${{ github.repository }}
- name: Create release tag
if: ${{ github.event.inputs.dry_run == 'false' }}
run: |
VERSION="${{ steps.checks.outputs.version }}"
git config user.name "github-actions[bot]"
git config user.email "github-actions[bot]@users.noreply.github.com"
git tag -a "${VERSION}" -m "Release ${VERSION}"
git push origin "${VERSION}"
echo "Created and pushed tag ${VERSION}"
- name: Dry run summary
if: ${{ github.event.inputs.dry_run == 'true' }}
run: |
echo "Dry run complete - all checks passed."
echo "Would have created tag: ${{ steps.checks.outputs.version }}"
-23
View File
@@ -1,23 +0,0 @@
name: Convert PR to draft
on:
pull_request_target:
types: [labeled]
permissions:
pull-requests: write
issues: write
contents: write # required for "gh pr ready" command, see https://github.com/cli/cli/issues/8910
jobs:
convert-to-draft:
if: github.event.label.name == 'draft' && github.event.pull_request.draft == false
runs-on: ubuntu-slim
steps:
- name: Convert PR to draft
env:
GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
PR_URL: ${{ github.event.pull_request.html_url }}
run: |
gh pr ready --undo "$PR_URL"
gh pr edit "$PR_URL" --remove-label draft
+245 -266
View File
@@ -93,13 +93,13 @@ jobs:
- build: 'arm64'
arch: 'arm64'
os: macos-26
defines: "-DGGML_METAL_EMBED_LIBRARY=ON -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3"
defines: "-DGGML_METAL_USE_BF16=ON -DGGML_METAL_EMBED_LIBRARY=ON -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3"
# TODO: this build is disabled to save Github Actions resources (https://github.com/ggml-org/llama.cpp/pull/23780)
# in order to enable it again, we have to provision dedicated runners to run it
#- build: 'arm64-kleidiai'
# arch: 'arm64'
# os: macos-14
# defines: "-DGGML_METAL_EMBED_LIBRARY=ON -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3 -DGGML_CPU_KLEIDIAI=ON"
# defines: "-DGGML_METAL_USE_BF16=ON -DGGML_METAL_EMBED_LIBRARY=ON -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3 -DGGML_CPU_KLEIDIAI=ON"
- build: 'x64'
arch: 'x64'
os: macos-15-intel
@@ -748,135 +748,6 @@ jobs:
path: llama-bin-win-cpu-${{ matrix.arch }}.zip
name: llama-bin-win-cpu-${{ matrix.arch }}.zip
windows-rocm:
needs: [check-release]
if: ${{ needs.check-release.outputs.should_release == 'true' }}
runs-on: windows-2022
strategy:
matrix:
include:
- ROCM_VERSION: "7.14.0"
gpu_targets: "gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
build: x64
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
with:
fetch-depth: 0
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
evict-old-files: 1d
# - name: Cache ROCm Installation
# id: cache-rocm
# uses: actions/cache@v5
# with:
# path: C:\TheRock\build
# key: rocm-wheels-${{ matrix.ROCM_VERSION }}-multi-arch-${{ runner.os }}
- name: Setup ROCm
# if: steps.cache-rocm.outputs.cache-hit != 'true'
uses: ./.github/actions/windows-setup-rocm
with:
version: ${{ matrix.ROCM_VERSION }}
- name: Setup ROCm Environment
run: |
$ErrorActionPreference = "Stop"
# Activate venv from cache or fresh install
& C:\TheRock\build\.venv\Scripts\Activate.ps1
# Expand the devel tree (idempotent; no-op if already done during install)
rocm-sdk init
if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
# Get ROCm installation paths using the rocm-sdk CLI tool
$rocmPath = (rocm-sdk path --root)
if (-not $rocmPath) { throw "rocm-sdk path --root returned empty - devel package may not be installed" }
$rocmPath = $rocmPath.Trim()
$cmakePath = (rocm-sdk path --cmake).Trim()
$binPath = (rocm-sdk path --bin).Trim()
write-host "ROCm root: $rocmPath"
write-host "CMake path: $cmakePath"
write-host "Bin path: $binPath"
echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
echo "CMAKE_PREFIX_PATH=$cmakePath" >> $env:GITHUB_ENV
echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
echo "$binPath" >> $env:GITHUB_PATH
# Keep venv in PATH for subsequent steps
echo "C:\TheRock\build\.venv\Scripts" >> $env:GITHUB_PATH
- name: Build
run: |
mkdir build
cd build
cmake .. `
-G "Unix Makefiles" `
-DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
-DCMAKE_BUILD_TYPE=Release `
-DGGML_BACKEND_DL=ON `
-DGGML_NATIVE=OFF `
-DGGML_CPU=ON `
-DGGML_CPU_ALL_VARIANTS=ON `
-DGGML_HIP=ON `
-DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
-DCMAKE_C_FLAGS="-Wno-error=incompatible-pointer-types" `
-DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
-DHIP_PATH="${env:HIP_PATH}" `
-DGGML_HIP_ROCWMMA_FATTN=ON `
-DAMDGPU_TARGETS="${{ matrix.gpu_targets }}"
cmake --build . --config Release --parallel ${env:NUMBER_OF_PROCESSORS}
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
- name: Verify HIP backend was built
run: |
$hipDll = Get-ChildItem -Path build\bin -Filter "ggml-hip*.dll" -ErrorAction SilentlyContinue
if (-not $hipDll) {
Write-Host "##[error]ggml-hip*.dll was NOT produced. The HIP backend silently failed to build."
Write-Host "Contents of build\bin:"
Get-ChildItem build\bin | Format-Table -AutoSize
exit 1
}
Write-Host "HIP backend artifact found:"
$hipDll | Format-Table FullName, Length -AutoSize
- name: Determine tag name
id: tag
uses: ./.github/actions/get-tag-name
- name: Get ROCm short version
run: |
$rocmVersionShort = ('${{ matrix.ROCM_VERSION }}'.Split('.')[0..1] -join '.')
echo "ROCM_VERSION_SHORT=$rocmVersionShort" >> $env:GITHUB_ENV
- name: Pack artifacts
run: |
cp "LICENSE" "build\bin\"
7z a -snl llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip .\build\bin\*
- name: Upload artifacts
uses: actions/upload-artifact@v6
with:
path: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip
name: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip
windows:
needs: [check-release]
if: ${{ needs.check-release.outputs.should_release == 'true' }}
@@ -977,7 +848,6 @@ jobs:
name: llama-bin-win-${{ matrix.backend }}-${{ matrix.arch }}.zip
windows-cuda:
name: windows-cuda (${{ matrix.cuda }}, ${{ matrix.arch }})
needs: [check-release]
if: ${{ needs.check-release.outputs.should_release == 'true' }}
@@ -988,16 +858,7 @@ jobs:
strategy:
matrix:
include:
- cuda: '12.4'
arch: x64
defines: '-DGGML_CUDA_CUB_3DOT2=ON'
- cuda: '13.3'
arch: x64
defines: ''
- cuda: '13.4'
arch: arm64
defines: '-DCMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-msvc-cuda.cmake'
cuda: ['12.4', '13.3']
steps:
- name: Clone
@@ -1015,7 +876,6 @@ jobs:
uses: ./.github/actions/windows-setup-cuda
with:
cuda_version: ${{ matrix.cuda }}
cuda_arch: ${{ matrix.arch }}
- name: Install Ninja
id: install_ninja
@@ -1025,62 +885,54 @@ jobs:
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
key: release-windows-2022-x64-cuda-${{ matrix.cuda }}
- name: Build
id: cmake_build
shell: cmd
# TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project
run: |
call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" ${{ matrix.arch == 'x64' && 'x64' || 'amd64_arm64' }}
call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64
cmake -S . -B build -G "Ninja Multi-Config" ^
-DGGML_BACKEND_DL=ON ^
-DGGML_NATIVE=OFF ^
-DGGML_CPU=OFF ^
-DGGML_CUDA=ON ^
-DLLAMA_BUILD_BORINGSSL=ON ${{ matrix.defines }}
-DLLAMA_BUILD_BORINGSSL=ON ^
-DGGML_CUDA_CUB_3DOT2=ON
set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1
cmake --build build --config Release -j %NINJA_JOBS% --target ggml-cuda
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
key: release-windows-2022-x64-cuda-${{ matrix.cuda }}
- name: Pack artifacts
id: pack_artifacts
run: |
7z a -snl llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip .\build\bin\Release\ggml-cuda.dll
7z a -snl llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip .\build\bin\Release\ggml-cuda.dll
- name: Upload artifacts
uses: actions/upload-artifact@v6
with:
path: llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip
name: llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip
path: llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip
name: llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip
- name: Copy and pack Cuda runtime (x64)
if: ${{ matrix.arch == 'x64' }}
- name: Copy and pack Cuda runtime
run: |
echo "Cuda install location: ${{ env.CUDA_PATH }}"
$dst='.\build\bin\cudart\'
robocopy "${{env.CUDA_PATH}}\bin" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll
robocopy "${{env.CUDA_PATH}}\lib" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll
robocopy "${{env.CUDA_PATH}}\bin\x64" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll
7z a cudart-llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip $dst\*
- name: Copy and pack Cuda runtime (ARM64)
if: ${{ matrix.arch == 'arm64' }}
run: |
echo "Cuda install location: ${{ env.CUDA_PATH }}"
$dst='.\build\bin\cudart\'
robocopy "${{env.CUDA_PATH}}\bin\arm64" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll
7z a cudart-llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip $dst\*
7z a cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip $dst\*
- name: Upload Cuda runtime
uses: actions/upload-artifact@v6
with:
path: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip
name: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-${{ matrix.arch }}.zip
path: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip
name: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip
windows-sycl:
needs: [check-release]
@@ -1285,123 +1137,250 @@ jobs:
path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz
name: llama-bin-ubuntu-sycl-${{ matrix.build }}-x64.tar.gz
# ubuntu-22-rocm:
# needs: [check-release, get-version]
# if: ${{ needs.check-release.outputs.should_release == 'true' }}
ubuntu-22-rocm:
needs: [check-release, get-version]
if: ${{ needs.check-release.outputs.should_release == 'true' }}
# runs-on: ubuntu-22.04
runs-on: ubuntu-22.04
# permissions:
# actions: write
permissions:
actions: write
# strategy:
# matrix:
# include:
# - ROCM_VERSION: "7.14.0"
# gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
# build: 'x64'
strategy:
matrix:
include:
- ROCM_VERSION: "7.2.1"
gpu_targets: "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1150;gfx1200;gfx1201"
build: 'x64'
# steps:
# - name: Clone
# id: checkout
# uses: actions/checkout@v6
# with:
# fetch-depth: 0
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
with:
fetch-depth: 0
# - name: Setup Node.js
# uses: actions/setup-node@v6
# with:
# node-version: "24"
# cache: "npm"
# cache-dependency-path: "tools/ui/package-lock.json"
- name: Setup Node.js
uses: actions/setup-node@v6
with:
node-version: "24"
cache: "npm"
cache-dependency-path: "tools/ui/package-lock.json"
# - name: Free up disk space
# uses: ggml-org/free-disk-space@v1.3.1
# with:
# tool-cache: true
- name: Free up disk space
uses: ggml-org/free-disk-space@v1.3.1
with:
tool-cache: true
# # - name: ccache
# # uses: ggml-org/ccache-action@v1.2.21
# # with:
# # key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
# - name: Dependencies
# id: depends
# run: |
# sudo apt install -y build-essential git cmake wget
- name: Dependencies
id: depends
run: |
sudo apt install -y build-essential git cmake wget
# - name: Setup TheRock with Wheels
# id: therock_env
# run: |
# # Create Python virtual environment
# python3 -m venv .venv
# source .venv/bin/activate
- name: Setup Legacy ROCm
if: matrix.ROCM_VERSION == '7.2.1'
id: legacy_env
run: |
sudo mkdir --parents --mode=0755 /etc/apt/keyrings
wget https://repo.radeon.com/rocm/rocm.gpg.key -O - | \
gpg --dearmor | sudo tee /etc/apt/keyrings/rocm.gpg > /dev/null
# # Install ROCm wheels for build
# # libraries = HIP runtime and CMake configs needed for linking
# # devel = compilers, headers, static libs
# python -m pip install --upgrade pip
# python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"
sudo tee /etc/apt/sources.list.d/rocm.list << EOF
deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/${{ matrix.ROCM_VERSION }} jammy main
EOF
# # Get ROCm installation paths using the rocm-sdk CLI tool
# ROCM_PATH=$(rocm-sdk path --root)
# CMAKE_PATH=$(rocm-sdk path --cmake)
# BIN_PATH=$(rocm-sdk path --bin)
# echo "ROCM_PATH=$ROCM_PATH"
# echo "CMAKE_PATH=$CMAKE_PATH"
# echo "BIN_PATH=$BIN_PATH"
sudo tee /etc/apt/preferences.d/rocm-pin-600 << EOF
Package: *
Pin: release o=repo.radeon.com
Pin-Priority: 600
EOF
# # Set environment variables
# echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
# echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
# echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
# echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
# echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
sudo apt update
sudo apt-get install -y libssl-dev rocm-hip-sdk
# # Keep venv activated for subsequent steps
# echo "$(pwd)/.venv/bin" >> $GITHUB_PATH
- name: Setup TheRock
if: matrix.ROCM_VERSION != '7.2.1'
id: therock_env
run: |
wget https://repo.amd.com/rocm/tarball/therock-dist-linux-gfx1151-${{ matrix.ROCM_VERSION }}.tar.gz
mkdir install
tar -xf *.tar.gz -C install
export ROCM_PATH=$(pwd)/install
echo ROCM_PATH=$ROCM_PATH >> $GITHUB_ENV
echo PATH=$PATH:$ROCM_PATH/bin >> $GITHUB_ENV
echo LD_LIBRARY_PATH=$ROCM_PATH/lib:$ROCM_PATH/llvm/lib:$ROCM_PATH/lib/rocprofiler-systems >> $GITHUB_ENV
# - name: Build with native CMake HIP support
# id: cmake_build
# run: |
# cmake -B build -S . \
# -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
# -DCMAKE_BUILD_TYPE=Release \
# -DGGML_BACKEND_DL=ON \
# -DGGML_NATIVE=OFF \
# -DCMAKE_INSTALL_RPATH='$ORIGIN' \
# -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
# -DGGML_CPU_ALL_VARIANTS=ON \
# -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
# -DGGML_HIP=ON \
# -DHIP_PLATFORM=amd \
# -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
# ${{ env.CMAKE_ARGS }}
# cmake --build build --config Release -j $(nproc)
- name: Build with native CMake HIP support
id: cmake_build
run: |
cmake -B build -S . \
-DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
-DCMAKE_BUILD_TYPE=Release \
-DGGML_BACKEND_DL=ON \
-DGGML_NATIVE=OFF \
-DCMAKE_INSTALL_RPATH='$ORIGIN' \
-DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
-DGGML_CPU_ALL_VARIANTS=ON \
-DGPU_TARGETS="${{ matrix.gpu_targets }}" \
-DGGML_HIP=ON \
-DHIP_PLATFORM=amd \
-DGGML_HIP_ROCWMMA_FATTN=ON \
-DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
${{ env.CMAKE_ARGS }}
cmake --build build --config Release -j $(nproc)
# # - name: ccache-clear
# # uses: ./.github/actions/ccache-clear
# # with:
# # key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
# - name: Determine tag name
# id: tag
# uses: ./.github/actions/get-tag-name
- name: Determine tag name
id: tag
uses: ./.github/actions/get-tag-name
# - name: Get ROCm short version
# run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
- name: Get ROCm short version
run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
# - name: Pack artifacts
# id: pack_artifacts
# run: |
# cp LICENSE ./build/bin/
# tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
- name: Pack artifacts
id: pack_artifacts
run: |
cp LICENSE ./build/bin/
tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
# - name: Upload artifacts
# uses: actions/upload-artifact@v6
# with:
# path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
# name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
- name: Upload artifacts
uses: actions/upload-artifact@v6
with:
path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
windows-hip:
needs: [check-release, get-version]
if: ${{ needs.check-release.outputs.should_release == 'true' }}
runs-on: windows-2022
permissions:
actions: write
env:
HIPSDK_INSTALLER_VERSION: "26.Q1"
strategy:
matrix:
include:
- name: "radeon"
gpu_targets: "gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032"
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
- name: Setup Node.js
uses: actions/setup-node@v6
with:
node-version: "24"
cache: "npm"
cache-dependency-path: "tools/ui/package-lock.json"
- name: Grab rocWMMA package
id: grab_rocwmma
run: |
curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2.1/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70201-81~24.04_amd64.deb"
7z x rocwmma.deb
7z x data.tar
- name: Cache ROCm Installation
id: cache-rocm
uses: actions/cache@v5
with:
path: C:\Program Files\AMD\ROCm
key: cache-gha-rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
- name: Install ROCm
if: steps.cache-rocm.outputs.cache-hit != 'true'
id: depends
run: |
$ErrorActionPreference = "Stop"
write-host "Downloading AMD HIP SDK Installer"
Invoke-WebRequest -Uri "https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ env.HIPSDK_INSTALLER_VERSION }}-Win11-For-HIP.exe" -OutFile "${env:RUNNER_TEMP}\rocm-install.exe"
write-host "Installing AMD HIP SDK"
$proc = Start-Process "${env:RUNNER_TEMP}\rocm-install.exe" -ArgumentList '-install' -NoNewWindow -PassThru
$completed = $proc.WaitForExit(600000)
if (-not $completed) {
Write-Error "ROCm installation timed out after 10 minutes. Killing the process"
$proc.Kill()
exit 1
}
if ($proc.ExitCode -ne 0) {
Write-Error "ROCm installation failed with exit code $($proc.ExitCode)"
exit 1
}
write-host "Completed AMD HIP SDK installation"
- name: Verify ROCm
id: verify
run: |
# Find and test ROCm installation
$clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1
if (-not $clangPath) {
Write-Error "ROCm installation not found"
exit 1
}
& $clangPath.FullName --version
- name: Build
id: cmake_build
run: |
$env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path)
$env:CMAKE_PREFIX_PATH="${env:HIP_PATH}"
cmake -G "Unix Makefiles" -B build -S . `
-DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" `
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" `
-DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.1/include/ -Wno-ignored-attributes -Wno-nested-anon-types" `
-DCMAKE_BUILD_TYPE=Release `
-DGGML_BACKEND_DL=ON `
-DGGML_NATIVE=OFF `
-DGGML_CPU=OFF `
-DGPU_TARGETS="${{ matrix.gpu_targets }}" `
-DGGML_HIP_ROCWMMA_FATTN=ON `
-DGGML_HIP=ON `
-DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} `
-DLLAMA_BUILD_BORINGSSL=ON
cmake --build build --target ggml-hip -j ${env:NUMBER_OF_PROCESSORS}
md "build\bin\rocblas\library\"
md "build\bin\hipblaslt\library"
cp "${env:HIP_PATH}\bin\libhipblas.dll" "build\bin\"
cp "${env:HIP_PATH}\bin\libhipblaslt.dll" "build\bin\"
cp "${env:HIP_PATH}\bin\rocblas.dll" "build\bin\"
cp "${env:HIP_PATH}\bin\rocblas\library\*" "build\bin\rocblas\library\"
cp "${env:HIP_PATH}\bin\hipblaslt\library\*" "build\bin\hipblaslt\library\"
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
key: release-windows-2022-x64-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ matrix.name }}
- name: Pack artifacts
id: pack_artifacts
run: |
7z a -snl llama-bin-win-hip-${{ matrix.name }}-x64.zip .\build\bin\*
- name: Upload artifacts
uses: actions/upload-artifact@v6
with:
path: llama-bin-win-hip-${{ matrix.name }}-x64.zip
name: llama-bin-win-hip-${{ matrix.name }}-x64.zip
ios-xcode:
needs: [check-release, get-version]
@@ -1423,6 +1402,7 @@ jobs:
run: |
sysctl -a
cmake -B build -G Xcode \
-DGGML_METAL_USE_BF16=ON \
-DGGML_METAL_EMBED_LIBRARY=ON \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_APP=OFF \
@@ -1576,9 +1556,9 @@ jobs:
- windows-cpu
- windows-cuda
#- windows-sycl
- windows-rocm
- windows-hip
- windows-openvino
#- ubuntu-22-rocm
- ubuntu-22-rocm
- ubuntu-cpu
- ubuntu-vulkan
- ubuntu-24-openvino
@@ -1688,7 +1668,7 @@ jobs:
- [Ubuntu s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-s390x.tar.gz)
- [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-x64.tar.gz)
- [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-arm64.tar.gz)
- Ubuntu x64 (ROCm 7.14)[DISABLED](https://github.com/ggml-org/llama.cpp/pull/26969)
- [Ubuntu x64 (ROCm 7.2)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-7.2-x64.tar.gz)
- [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-openvino-${{ needs.ubuntu-24-openvino.outputs.openvino_version }}-x64.tar.gz)
- [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp32-x64.tar.gz)
- [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp16-x64.tar.gz)
@@ -1702,11 +1682,10 @@ jobs:
- [Windows arm64 (OpenCL Adreno)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-opencl-adreno-arm64.zip)
- [Windows x64 (CUDA 12)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-cuda-12.4-x64.zip) - [CUDA 12.4 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/cudart-llama-bin-win-cuda-12.4-x64.zip)
- [Windows x64 (CUDA 13)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-cuda-13.3-x64.zip) - [CUDA 13.3 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/cudart-llama-bin-win-cuda-13.3-x64.zip)
- [Windows arm64 (CUDA 13) (preview)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-cuda-13.4-arm64.zip) - [CUDA 13.4 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/cudart-llama-bin-win-cuda-13.4-arm64.zip)
- [Windows x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-vulkan-x64.zip)
- [Windows x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-openvino-${{ needs.windows-openvino.outputs.openvino_version }}-x64.zip)
- [Windows x64 (SYCL)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-sycl-x64.zip)
- [Windows x64 (ROCm 7.14)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-rocm-7.14-x64.zip)
- [Windows x64 (HIP)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-hip-radeon-x64.zip)
**openEuler:**
- [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23705)
+6 -16
View File
@@ -25,12 +25,6 @@ on:
'tools/server/**.*'
]
pull_request:
types: [opened, synchronize, reopened]
paths: [
'.github/workflows/server-sanitize.yml'
]
env:
LLAMA_ARG_LOG_COLORS: 1
LLAMA_ARG_LOG_PREFIX: 1
@@ -96,27 +90,23 @@ jobs:
- name: Python setup
id: setup_python
uses: actions/setup-python@v7
- name: Install Python dependencies
run: |
python3 -m venv .venv
.venv/bin/pip install -r tools/server/tests/requirements.txt
uses: actions/setup-python@v6
with:
python-version: '3.11'
pip-install: -r tools/server/tests/requirements.txt
- name: Tests
id: server_integration_tests
if: ${{ (!matrix.disabled_on_pr || !github.event.pull_request) }}
run: |
source .venv/bin/activate
cd tools/server/tests
export ${{ matrix.extra_args }}
./tests.sh
pytest -v -x -m "not slow"
- name: Slow tests
id: server_integration_tests_slow
if: ${{ (github.event.schedule || github.event.inputs.slow_tests == 'true') && matrix.build_type == 'Release' }}
run: |
source .venv/bin/activate
cd tools/server/tests
export ${{ matrix.extra_args }}
SLOW_TESTS=1 ./tests.sh
SLOW_TESTS=1 pytest -v -x
+9 -9
View File
@@ -72,7 +72,7 @@ jobs:
run: |
cd tools/server/tests
source venv/bin/activate
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx1, backend-sampling)
id: server_integration_tests_backend_sampling
@@ -81,7 +81,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export LLAMA_ARG_BACKEND_SAMPLING=1
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx2)
id: server_integration_tests_gpu2
@@ -90,7 +90,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export GGML_METAL_DEVICES=2
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx2, backend-sampling)
id: server_integration_tests_gpu2_backend_sampling
@@ -99,7 +99,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export GGML_METAL_DEVICES=2 LLAMA_ARG_BACKEND_SAMPLING=1
./tests.sh
pytest -v -x -m "not slow"
server-cuda:
runs-on: [self-hosted, llama-server, Linux, NVIDIA]
@@ -132,7 +132,7 @@ jobs:
run: |
cd tools/server/tests
source venv/bin/activate
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx1, backend-sampling)
id: server_integration_tests_backend_sampling
@@ -141,7 +141,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export LLAMA_ARG_BACKEND_SAMPLING=1
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx2)
id: server_integration_tests_gpu2
@@ -150,7 +150,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export GGML_CUDA_DEVICES=2
./tests.sh
pytest -v -x -m "not slow"
- name: Tests (GPUx2, backend-sampling)
id: server_integration_tests_gpu2_backend_sampling
@@ -159,7 +159,7 @@ jobs:
cd tools/server/tests
source venv/bin/activate
export GGML_CUDA_DEVICES=2 LLAMA_ARG_BACKEND_SAMPLING=1
./tests.sh
pytest -v -x -m "not slow"
server-kleidiai:
runs-on: ah-ubuntu_22_04-c8g_8x
@@ -219,4 +219,4 @@ jobs:
run: |
cd tools/server/tests
source venv/bin/activate
./tests.sh
pytest -v -x -m "not slow"
+8 -10
View File
@@ -104,21 +104,21 @@ jobs:
id: server_integration_tests
run: |
cd tools/server/tests
./tests.sh
pytest -v -x -m "not slow"
- name: Slow tests
id: server_integration_tests_slow
if: ${{ github.event.schedule || github.event.inputs.slow_tests == 'true' }}
run: |
cd tools/server/tests
SLOW_TESTS=1 ./tests.sh
SLOW_TESTS=1 pytest -v -x
- name: Tests (Backend sampling)
id: server_integration_tests_backend_sampling
run: |
cd tools/server/tests
export LLAMA_ARG_BACKEND_SAMPLING=1
./tests.sh
pytest -v -x -m "not slow"
- name: Slow tests (Backend sampling)
id: server_integration_tests_slow_backend_sampling
@@ -126,7 +126,7 @@ jobs:
run: |
cd tools/server/tests
export LLAMA_ARG_BACKEND_SAMPLING=1
SLOW_TESTS=1 ./tests.sh
SLOW_TESTS=1 pytest -v -x
windows:
runs-on: windows-2025
@@ -167,17 +167,15 @@ jobs:
- name: Tests
id: server_integration_tests
shell: bash
run: |
cd tools/server/tests
export PYTHONIOENCODING=":replace"
./tests.sh
$env:PYTHONIOENCODING = ":replace"
pytest -v -x -m "not slow"
- name: Slow tests
id: server_integration_tests_slow
if: ${{ github.event.schedule || github.event.inputs.slow_tests == 'true' }}
shell: bash
run: |
cd tools/server/tests
export SLOW_TESTS="1"
./tests.sh
$env:SLOW_TESTS = "1"
pytest -v -x
-2
View File
@@ -19,8 +19,6 @@ jobs:
run: |
cargo binstall komac@2.16.0 -y
# TODO: This should later be updated to publish releases instead of
# development release builds.
- name: Find latest release
id: find_latest_release
uses: actions/github-script@v8
+7 -23
View File
@@ -2,26 +2,6 @@ cmake_minimum_required(VERSION 3.14...3.28) # for add_link_options and implicit
project("llama.cpp" C CXX)
include(CheckIncludeFileCXX)
### llama.cpp version
set(LLAMA_VERSION_MAJOR 0)
set(LLAMA_VERSION_MINOR 1)
set(LLAMA_VERSION_PATCH 0)
set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
# whether this is a development/nightly build
# set this to OFF when making a release from a release tag (vX.Y.Z)
# ref: https://github.com/ggml-org/ggml/discussions/1579
option(LLAMA_BUILD_IS_DEV "llama: dev build" ON)
if (LLAMA_BUILD_IS_DEV)
set(LLAMA_VERSION "${LLAMA_VERSION_BASE}-dev")
else()
# TODO: check that the current commit is tagged correctly according to the version specified above
set(LLAMA_VERSION "${LLAMA_VERSION_BASE}")
endif()
message(STATUS "llama.cpp version: ${LLAMA_VERSION}")
#set(CMAKE_WARN_DEPRECATED YES)
set(CMAKE_WARN_UNUSED_CLI YES)
@@ -44,6 +24,9 @@ if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
set(LLAMA_STANDALONE ON)
include(git-vars)
# configure project version
# TODO
else()
set(LLAMA_STANDALONE OFF)
endif()
@@ -156,6 +139,7 @@ endif()
if (NOT DEFINED LLAMA_BUILD_COMMIT)
set(LLAMA_BUILD_COMMIT ${BUILD_COMMIT})
endif()
set(LLAMA_INSTALL_VERSION 0.0.${LLAMA_BUILD_NUMBER})
# override ggml options
set(GGML_ALL_WARNINGS ${LLAMA_ALL_WARNINGS})
@@ -291,12 +275,12 @@ configure_package_config_file(
LLAMA_BIN_INSTALL_DIR )
write_basic_package_version_file(
${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
VERSION ${LLAMA_VERSION}
${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
VERSION ${LLAMA_INSTALL_VERSION}
COMPATIBILITY SameMajorVersion)
install(FILES ${CMAKE_CURRENT_BINARY_DIR}/llama-config.cmake
${CMAKE_CURRENT_BINARY_DIR}/llama-config-version.cmake
${CMAKE_CURRENT_BINARY_DIR}/llama-version.cmake
DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/llama)
configure_file(cmake/llama.pc.in
+1 -2
View File
@@ -12,7 +12,7 @@
[![Docker](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
[![Winget](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
[manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3A0cc4m%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [compile times](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-compile-times.md) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
[manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3A0cc4m%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [dev branches](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-features.md) / [compile times](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-compile-times.md) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
</div>
@@ -106,7 +106,6 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
- [XCFramework](docs/xcframework.md)
- [Completions](docs/completions.md)
- [Models](docs/models.md)
- [Release process](docs/release.md)
## Contributing
-9
View File
@@ -21,18 +21,11 @@ Please disclose it as a private [security advisory](https://github.com/ggml-org/
A team of volunteers on a reasonable-effort basis maintains this project. As such, please give us at least 90 days to work on a fix before public exposure.
### AI-powered code scan
llama.cpp has an AI security scanner that scans the code periodically. The full prompts and tool set can be found in [ggml-org/security-scan-prompt](https://github.com/ggml-org/security-scan-prompt).
We greatly appreciate reports that reflect genuine research effort, and we are happy to spend our time reviewing them. Findings that an autonomous AI agent can surface on its own add little on top of the scans we already run.
### Requirements
Before submitting your report, ensure you meet the following requirements:
- You have read this policy and fully understand it.
- You have searched for existing discussions of the issue. If it has already been reported, your report will likely be rejected as a duplicate.
- AI is only permitted in an assistive capacity as stated in [AGENTS.md](AGENTS.md). We do not accept reports that are written exclusively by AI.
- Your report must include a working Proof-of-Concept in the form of a script and/or attached files.
@@ -53,8 +46,6 @@ Only vulnerabilities that fall within these parts of the project are considered
Note that none of the topics under [Using llama.cpp securely](#using-llamacpp-securely) are considered vulnerabilities in LLaMA C++.
Denial-of-Service (DoS) bugs are generally not treated as vulnerabilities. We don't reject them outright, but we look at them case-by-case and only accept those that are genuinely worth fixing.
For vulnerabilities that fall within the `vendor` directory, please report them directly to the third-party project.
## Using llama.cpp securely
+3 -5
View File
@@ -1,7 +1,5 @@
#include "build-info.h"
#include "llama.h"
#include <cstdio>
#include <cstdlib>
#include <string>
@@ -79,12 +77,12 @@ static const command cmds[] = {
#undef UPDATE_HIDDEN
static int version(int /*argc*/, char ** /*argv*/) {
llama_print_build_info(llama_version());
static int version(int argc, char ** argv) {
printf("%s\n", llama_build_info());
return 0;
}
static int licenses(int /*argc*/, char ** /*argv*/) {
static int licenses(int argc, char ** argv) {
for (int i = 0; LICENSES[i]; ++i) {
printf("%s\n", LICENSES[i]);
}
+2
View File
@@ -17,6 +17,7 @@ LLAMA_BUILD_MTMD=ON
GGML_METAL=ON
GGML_METAL_EMBED_LIBRARY=ON
GGML_BLAS_DEFAULT=ON
GGML_METAL_USE_BF16=ON
GGML_OPENMP=OFF
COMMON_C_FLAGS="-Wno-macro-redefined -Wno-shorten-64-to-32 -Wno-unused-command-line-argument -g"
@@ -43,6 +44,7 @@ COMMON_CMAKE_ARGS=(
-DGGML_METAL_EMBED_LIBRARY=${GGML_METAL_EMBED_LIBRARY}
-DGGML_BLAS_DEFAULT=${GGML_BLAS_DEFAULT}
-DGGML_METAL=${GGML_METAL}
-DGGML_METAL_USE_BF16=${GGML_METAL_USE_BF16}
-DGGML_NATIVE=OFF
-DGGML_OPENMP=${GGML_OPENMP}
)
+10 -34
View File
@@ -10,9 +10,6 @@
# # with CUDA support
# GG_BUILD_CUDA=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt
#
# # with ROCm support
# GG_BUILD_ROCM=1 GG_BUILD_AMDGPU_TARGETS=gfx1151 bash ./ci/run.sh ./tmp/results ./tmp/mnt
#
# # with SYCL support
# GG_BUILD_SYCL=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt
#
@@ -49,14 +46,6 @@ mkdir -p "$2"
OUT=$(realpath "$1")
MNT=$(realpath "$2")
# gpu-rocm self-hosted runner can't upload logs to blob; keep each run's logs in
# their own dir keyed by the GitHub run id so an Actions run URL maps to its logs.
if [ -n "${GG_BUILD_ROCM}" ] && [ -n "${GITHUB_RUN_ID}" ]; then
OUT="$OUT/run-${GITHUB_RUN_ID}-${GITHUB_RUN_ATTEMPT:-1}"
mkdir -p "$OUT"
echo "ci results dir: $OUT"
fi
rm -f $OUT/*.log
rm -f $OUT/*.exit
rm -f $OUT/*.md
@@ -100,7 +89,7 @@ if [ ! -z ${GG_BUILD_CUDA} ]; then
fi
if [ ! -z ${GG_BUILD_ROCM} ]; then
CMAKE_EXTRA="${CMAKE_EXTRA} -DCMAKE_HIP_COMPILER=$(hipconfig -l)/clang -DGGML_HIP=ON"
CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_HIP=ON"
if [ -z ${GG_BUILD_AMDGPU_TARGETS} ]; then
echo "Missing GG_BUILD_AMDGPU_TARGETS, please set it to your GPU architecture (e.g. gfx90a, gfx1100, etc.)"
exit 1
@@ -651,52 +640,39 @@ function gg_sum_rerank_tiny {
function gg_check_build_requirements {
if ! command -v git &> /dev/null; then
gg_printf 'git not found, please install\n'
exit 1
gg_printf 'git not found, please install'
fi
if ! command -v git-lfs &> /dev/null; then
gg_printf 'git-lfs not found, please install\n'
exit 1
fi
if ! git config --get filter.lfs.clean &> /dev/null; then
gg_printf 'git-lfs not initialized, please run `git lfs install`\n'
exit 1
gg_printf 'git-lfs not found, please install'
fi
if ! command -v wget &> /dev/null; then
gg_printf 'wget not found, please install\n'
exit 1
gg_printf 'wget not found, please install'
fi
if ! command -v python3 &> /dev/null; then
gg_printf 'python3 not found, please install\n'
exit 1
gg_printf 'python3 not found, please install'
fi
if ! command -v pip3 &> /dev/null; then
gg_printf 'pip3 not found, please install\n'
exit 1
gg_printf 'pip3 not found, please install'
fi
if ! python3 -m ensurepip --help &> /dev/null; then
gg_printf 'ensurepip not found, please install python3-venv package\n'
exit 1
gg_printf 'ensurepip not found, please install python3-venv package'
fi
if ! command -v cmake &> /dev/null; then
gg_printf 'cmake not found, please install\n'
exit 1
gg_printf 'cmake not found, please install'
fi
if ! command -v ccache &> /dev/null; then
gg_printf 'ccache not found, please consider installing for faster builds\n'
gg_printf 'ccache not found, please consider installing for faster builds'
fi
if ! command -v ctest &> /dev/null; then
gg_printf 'ctest not found, please install\n'
exit 1
gg_printf 'ctest not found, please install'
fi
}
-26
View File
@@ -1,26 +0,0 @@
# Used to cross-compile ggml-cuda for Windows ARM64 on an x64 Windows host.
set( CMAKE_SYSTEM_NAME Windows )
set( CMAKE_SYSTEM_PROCESSOR arm64 )
if ( DEFINED CUDAToolkit_ROOT )
file( TO_CMAKE_PATH "${CUDAToolkit_ROOT}" CUDA_ROOT )
elseif ( DEFINED ENV{CUDA_PATH} )
file( TO_CMAKE_PATH "$ENV{CUDA_PATH}" CUDA_ROOT )
else()
message( FATAL_ERROR "Set CUDAToolkit_ROOT or CUDA_PATH to a Windows CUDA Toolkit with ARM64 target libraries" )
endif()
if ( DEFINED ENV{VCToolsInstallDir} )
file( TO_CMAKE_PATH "$ENV{VCToolsInstallDir}" MSVC_TOOLS_ROOT )
set( CMAKE_CUDA_HOST_COMPILER "${MSVC_TOOLS_ROOT}/bin/Hostx64/arm64/cl.exe" CACHE FILEPATH "" )
endif()
set( CMAKE_CUDA_COMPILER "${CUDA_ROOT}/bin/nvcc.exe" CACHE FILEPATH "" )
set( CMAKE_CUDA_FLAGS_INIT "-target-dir=arm64" )
# FindCUDAToolkit selects lib/x64 from the host architecture on Windows.
set( CUDA_CUDART "${CUDA_ROOT}/lib/arm64/cudart.lib" CACHE FILEPATH "" )
set( CUDA_cudart_LIBRARY "${CUDA_ROOT}/lib/arm64/cudart.lib" CACHE FILEPATH "" )
set( CUDA_cublas_LIBRARY "${CUDA_ROOT}/lib/arm64/cublas.lib" CACHE FILEPATH "" )
set( CUDA_cublasLt_LIBRARY "${CUDA_ROOT}/lib/arm64/cublasLt.lib" CACHE FILEPATH "" )
set( CUDA_cuda_driver_LIBRARY "${CUDA_ROOT}/lib/arm64/cuda.lib" CACHE FILEPATH "" )
+1 -1
View File
@@ -1,4 +1,4 @@
set(LLAMA_VERSION @LLAMA_VERSION@)
set(LLAMA_VERSION @LLAMA_INSTALL_VERSION@)
set(LLAMA_BUILD_COMMIT @LLAMA_BUILD_COMMIT@)
set(LLAMA_BUILD_NUMBER @LLAMA_BUILD_NUMBER@)
set(LLAMA_SHARED_LIB @BUILD_SHARED_LIBS@)
+1 -1
View File
@@ -5,6 +5,6 @@ includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@
Name: llama
Description: Port of Facebook's LLaMA model in C/C++
Version: @LLAMA_VERSION@
Version: @LLAMA_INSTALL_VERSION@
Libs: -L${libdir} -lggml -lggml-base -lllama
Cflags: -I${includedir}
+2 -2
View File
@@ -121,8 +121,8 @@ add_library(${TARGET}
)
set_target_properties(${TARGET} PROPERTIES
VERSION ${LLAMA_VERSION_BASE}
SOVERSION ${LLAMA_VERSION_MAJOR}
VERSION ${LLAMA_INSTALL_VERSION}
SOVERSION 0
MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
)
+4 -66
View File
@@ -35,7 +35,6 @@
#include <regex>
#include <set>
#include <string>
#include <system_error>
#include <thread> // for hardware_concurrency
#include <vector>
@@ -705,61 +704,12 @@ void common_models_handler_apply(common_models_handler & handler, common_params
// CLI argument parsing functions
//
// apply config files (if present), a later file overrides an earlier one:
// 1. system-wide: /etc/llama.cpp/config.ini (%PROGRAMDATA%\llama.cpp\config.ini on windows)
// 2. user-level: ${XDG_CONFIG_HOME:-~/.config}/llama.cpp/config.ini (%APPDATA%\llama.cpp\config.ini on windows)
static void common_params_apply_system_config(common_params & params, llama_example ex) {
std::vector<std::string> paths;
#if defined(_WIN32)
const std::string program_data = common_get_env("PROGRAMDATA");
if (!program_data.empty()) {
paths.push_back(program_data + "\\llama.cpp\\config.ini");
}
#else
paths.push_back("/etc/llama.cpp/config.ini");
#endif
try {
paths.push_back(fs_get_config_directory() + "config.ini");
} catch (const std::exception & e) {
LOG_DBG("cannot read user-level config file, skipping: %s\n", e.what());
}
std::vector<std::string> found;
for (const auto & path : paths) {
std::error_code ec;
if (std::filesystem::exists(path, ec)) {
found.push_back(path);
}
}
if (found.empty()) {
return;
}
common_preset_context ctx(ex);
ctx.ignore_unknown_keys = true; // the same config file is shared by all programs
for (const auto & path : found) {
LOG_INF("using config file: %s\n", path.c_str());
common_preset global;
common_presets presets = ctx.load_from_ini(path, global);
global.apply_to_params(params);
auto it = presets.find(COMMON_PRESET_DEFAULT_NAME);
if (it != presets.end()) {
it->second.apply_to_params(params);
}
}
}
static bool common_params_parse_ex(int argc, char ** argv, common_params_context & ctx_arg) {
common_params & params = ctx_arg.params;
// setup log directly from params.verbosity: see tools/cli/cli.cpp
common_log_set_verbosity_thold(params.verbosity);
// config file applies first, so env variables and CLI arguments override it
common_params_apply_system_config(params, ctx_arg.ex);
std::unordered_map<std::string, std::pair<common_arg *, bool>> arg_to_options;
for (auto & opt : ctx_arg.options) {
for (const auto & arg : opt.args) {
@@ -1440,7 +1390,8 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
{"--version"},
"show version and build info",
[](common_params &) {
llama_print_build_info(llama_version());
fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit());
fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
exit(0);
}
));
@@ -2654,16 +2605,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
).set_env("LLAMA_ARG_DIO"));
add_opt(common_arg(
{"-lm", "--load-mode"}, "MODE",
"model loading mode (default: auto)\n"
"- auto: mmap, unless a device does not support it\n"
"model loading mode (default: mmap)\n"
"- none: no special loading mode\n"
"- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)\n"
"- mlock: force system to keep model in RAM rather than swapping or compressing\n"
"- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing\n"
"- dio: use DirectIO if available\n",
[](common_params & params, const std::string & value) {
/**/ if (value == "auto") { params.load_mode = LLAMA_LOAD_MODE_AUTO; }
else if (value == "none") { params.load_mode = LLAMA_LOAD_MODE_NONE; }
/**/ if (value == "none") { params.load_mode = LLAMA_LOAD_MODE_NONE; }
else if (value == "mmap") { params.load_mode = LLAMA_LOAD_MODE_MMAP; }
else if (value == "mlock") { params.load_mode = LLAMA_LOAD_MODE_MLOCK; }
else if (value == "mmap+mlock") { params.load_mode = LLAMA_LOAD_MODE_MMAP_MLOCK; }
@@ -3359,17 +3308,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.server_tools = parse_csv_row(value);
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_TOOLS"));
add_opt(common_arg(
{"--tools-runtime"}, "OPTION",
"experimental: run tools in a separate runtime environment (default: none, use host environment)\n"
"available options:\n"
" 'docker:<image>', 'podman:<image>': spin up a new container and reuse it for all invocations, clean up on server exit\n"
" 'docker-container:<id>', 'podman-container:<id>': use an existing container by ID, won't stop on server exit\n"
" 'ssh:<target>': run tools on a remote POSIX host over SSH, key-based auth and a trusted host key are required\n",
[](common_params & params, const std::string & value) {
params.server_tools_runtime = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_TOOLS_RUNTIME"));
add_opt(common_arg(
{"--mcp-servers-config"}, "PATH",
"experimental: path to JSON file with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)\n"
+3 -3
View File
@@ -29,7 +29,7 @@ const char * llama_build_info(void) {
return s.c_str();
}
void llama_print_build_info(const char * llama_version) {
fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version, llama_build_number(), llama_commit());
fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
void llama_print_build_info(void) {
fprintf(stderr, "%s: build = %d (%s)\n", __func__, llama_build_number(), llama_commit());
fprintf(stderr, "%s: built with %s for %s\n", __func__, llama_compiler(), llama_build_target());
}
+1 -1
View File
@@ -8,4 +8,4 @@ const char * llama_compiler(void);
const char * llama_build_target(void);
const char * llama_build_info(void);
void llama_print_build_info(const char *);
void llama_print_build_info(void);
+7 -167
View File
@@ -1166,16 +1166,6 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
data.prompt += data.generation_prompt;
}
std::vector<std::string> tool_call_starts = { "<tool_call>" };
// Match complete <function=name> opener for Qwen3-Coder models that occasionally omit the
// starting <tool_call>. The model may hallucinate a tool name, but it is preferable over
// constraining on <function which may occur in valid content generation, e.g. #include <functional>
foreach_function(inputs.tools, [&](const json & tool) {
const std::string name = tool.at("function").at("name");
tool_call_starts.push_back("<function=" + name + ">");
});
auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) {
auto generation_prompt = p.literal(GEN_PREFIX);
@@ -1248,7 +1238,7 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
auto tool_calls = p.trigger_rule("tool-call-root", p.repeat(calls, min_calls, 1));
return generation_prompt +
(reasoning << p.content(p.until_one_of(tool_call_starts)) << tool_calls);
(reasoning << p.content(p.until_one_of({ "<tool_call>", "<function=" })) << tool_calls);
}
// Content only parser
@@ -1274,9 +1264,12 @@ static common_chat_params common_chat_params_init_qwen3_coder(const common_chat_
});
if (data.grammar_lazy) {
for (const auto & start : tool_call_starts) {
data.grammar_triggers.push_back({ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, start });
}
data.grammar_triggers = {
{ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "<tool_call>" },
// Trigger on "<function" and not "<function=" because the trailing "=" is part of
// the token with the function name e.g. "=read"
{ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "<function" },
};
}
}
@@ -3093,153 +3086,6 @@ static common_chat_params common_chat_params_init_minicpm5(const common_chat_tem
return data;
}
// An assistant turn is rendered as one or more messages, each
// "<|start|>assistant to=<recipient><|message|>{content}{END}" where END is
// <|eom|> (more messages follow) or <|eot|> (end of turn):
// - chain-of-thought: to=self, terminated by <|eom|>
// - final answer: to=user, terminated by <|eot|>
// The generation prompt is just "<|start|>assistant"; the model emits its own
// " to=...<|message|>".
static common_chat_params common_chat_params_init_muse_glimmer(const common_chat_template & tmpl,
const autoparser::generation_params & inputs) {
common_chat_params data;
data.prompt = common_chat_template_direct_apply_impl(tmpl, inputs);
data.generation_prompt = "<|start|>assistant";
data.format = COMMON_CHAT_FORMAT_PEG_NATIVE;
data.supports_thinking = true;
data.preserved_tokens = {
"<|start|>", "<|message|>", "<|eom|>", "<|eot|>",
// ATEM tool-call markup emitted on " to=<tool>" turns.
"<atem:function_calls>", "<atem:invoke", "<atem:parameter", "</atem:parameter>",
"</atem:invoke>", "</atem:function_calls>",
};
data.message_delimiters = {
{ COMMON_CHAT_ROLE_ASSISTANT, "<|start|>assistant" },
{ COMMON_CHAT_ROLE_USER, "<|start|>user" },
{ COMMON_CHAT_ROLE_SYSTEM, "<|start|>system" },
{ COMMON_CHAT_ROLE_TOOL, "<|start|>tool" },
};
if (inputs.has_continuation()) {
const auto & msg = inputs.continue_msg;
data.generation_prompt = "<|start|>assistant to=self<|message|>" + msg.reasoning_content;
if (inputs.continue_final_message == COMMON_CHAT_CONTINUATION_CONTENT) {
data.generation_prompt += "<|eom|><|start|>assistant to=user<|message|>" + msg.render_content();
}
data.prompt += data.generation_prompt;
}
auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE;
auto has_tools = inputs.tools.is_array() && !inputs.tools.empty();
// Constrained grammar whenever tools are offered.
auto include_grammar = has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE;
auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) {
auto start = p.rule("start", p.literal("<|start|>assistant"));
if (!extract_reasoning && !include_grammar) {
return start + p.content(p.rest());
}
if (extract_reasoning) {
p.rule("analysis", p.literal(" to=self<|message|>") + p.reasoning(p.until("<|eom|>")) + p.literal("<|eom|>"));
} else {
p.rule("analysis", p.literal(" to=self<|message|>") + p.content(p.until("<|eom|>")) + p.literal("<|eom|>"));
}
auto analysis = p.ref("analysis");
auto recipient = p.optional(p.literal(" to=user"));
auto final_msg = p.rule("final", recipient + p.literal("<|message|>") +
p.content(p.until_one_of({ "<|eot|>", "<|eom|>" })));
if (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE) {
auto string_value = p.ac(
p.tool_arg_string_value(p.until("</atem:parameter>")) + p.tool_arg_close(p.literal("</atem:parameter>")),
"</atem:parameter>");
auto tool_choice = p.choice();
foreach_function(inputs.tools, [&](const json & tool) {
const auto & function = tool.at("function");
const std::string name = function.at("name");
auto params = function.contains("parameters") ? function.at("parameters") : json::object();
auto args = p.eps();
if (params.contains("properties") && params.at("properties").is_object() && !params.at("properties").empty()) {
auto schema_info = common_schema_info();
schema_info.resolve_refs(params);
auto arg_choice = p.choice();
for (const auto & [prop_name, prop_schema] : params.at("properties").items()) {
auto value_parser = p.eps();
if (schema_info.resolves_to_string(prop_schema)) {
value_parser = string_value;
} else {
value_parser = p.tool_arg_json_value(
p.schema(p.json(), "tool-" + name + "-arg-" + prop_name + "-schema", prop_schema, false))
+ p.tool_arg_close(p.literal("</atem:parameter>"));
}
auto arg_rule = p.tool_arg(
p.tool_arg_open(p.literal("<atem:parameter name=\"") + p.tool_arg_name(p.literal(prop_name)) + p.literal("\">")) +
value_parser);
arg_choice |= arg_rule;
}
args = p.zero_or_more(arg_choice + p.space());
}
auto tool_parser = p.tool(
p.tool_open(p.literal(" to=") + p.until("<|message|>") +
p.literal("<|message|><atem:function_calls>") + p.space() +
p.literal("<atem:invoke name=\"") + p.tool_name(p.literal(name)) + p.literal("\">") + p.space())
<< p.tool_args(args)
<< p.tool_close(p.literal("</atem:invoke>") + p.space() + p.literal("</atem:function_calls>")));
tool_choice |= p.rule("tool-" + name, tool_parser);
});
auto tool_calls = inputs.parallel_tool_calls
? p.trigger_rule("tool-call", tool_choice + p.zero_or_more(p.literal("<|eom|>") + start + tool_choice))
: p.trigger_rule("tool-call", tool_choice);
if (inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED) {
return p.zero_or_more(start + analysis) + start + tool_calls;
}
auto trailing_calls = p.optional(p.literal("<|eom|>") + start + tool_calls);
return p.zero_or_more(start + analysis) + start + (tool_calls | (final_msg + trailing_calls));
}
return p.zero_or_more(start + analysis) + start + final_msg;
});
data.parser = parser.save();
if (include_grammar) {
data.grammar_lazy = inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_REQUIRED;
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
foreach_function(inputs.tools, [&](const json & tool) {
const auto & function = tool.at("function");
auto schema = function.contains("parameters") ? function.at("parameters") : json::object();
builder.resolve_refs(schema);
});
parser.build_grammar(builder, data.grammar_lazy);
});
data.grammar_triggers = {
{ COMMON_GRAMMAR_TRIGGER_TYPE_PATTERN,
"<\\|start\\|>assistant( to=(?!self<\\|message\\|>)(?!user<\\|message\\|>)[^<]*?<\\|message\\|>)" },
};
}
return data;
}
static json common_chat_extra_context() {
json ctx = json::object();
std::chrono::system_clock::time_point now = std::chrono::system_clock::now();
@@ -3268,12 +3114,6 @@ std::optional<common_chat_params> common_chat_try_specialized_template(
return common_chat_params_init_gpt_oss(tmpl, params);
}
// Muse Glimmer format using " to=<recipient>" recipients and <|eom|>/<|eot|> message terminators.
if (src.find("<atem:function_calls>") != std::string::npos && src.find("<|eom|>") != std::string::npos) {
LOG_DBG("Using specialized template: Muse Glimmer\n");
return common_chat_params_init_muse_glimmer(tmpl, params);
}
// Functionary v3.2 - uses recipient-based format with >>>recipient\n{content}
// Detection: template has ">>>all" for content and ">>>" prefix for tool calls
if (src.find(">>>all") != std::string::npos && src.find(">>>${recipient}") != std::string::npos) {
+10 -64
View File
@@ -1019,21 +1019,20 @@ std::string fs_get_cache_directory() {
std::string cache_directory = "";
auto ensure_trailing_slash = [](std::string p) {
// Make sure to add trailing slash
if (p.empty() || p.back() != DIRECTORY_SEPARATOR) {
if (p.back() != DIRECTORY_SEPARATOR) {
p += DIRECTORY_SEPARATOR;
}
return p;
};
cache_directory = common_get_env("LLAMA_CACHE");
if (cache_directory.empty()) {
if (getenv("LLAMA_CACHE")) {
cache_directory = std::getenv("LLAMA_CACHE");
} else {
#if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \
defined(__OpenBSD__) || defined(__NetBSD__)
const std::string xdg_cache_home = common_get_env("XDG_CACHE_HOME");
const std::string home = common_get_env("HOME");
if (!xdg_cache_home.empty()) {
cache_directory = xdg_cache_home;
} else if (!home.empty()) {
cache_directory = home + "/.cache/";
if (std::getenv("XDG_CACHE_HOME")) {
cache_directory = std::getenv("XDG_CACHE_HOME");
} else if (std::getenv("HOME")) {
cache_directory = std::getenv("HOME") + std::string("/.cache/");
} else {
#if defined(__linux__)
/* no $HOME is defined, fallback to getpwuid */
@@ -1048,16 +1047,9 @@ std::string fs_get_cache_directory() {
#endif /* defined(__linux__) */
}
#elif defined(__APPLE__)
cache_directory = common_get_env("HOME");
if (cache_directory.empty()) {
throw std::runtime_error("Failed to find $HOME directory");
}
cache_directory += "/Library/Caches/";
cache_directory = std::getenv("HOME") + std::string("/Library/Caches/");
#elif defined(_WIN32)
cache_directory = common_get_env("LOCALAPPDATA");
if (cache_directory.empty()) {
throw std::runtime_error("Failed to find %LOCALAPPDATA% directory");
}
cache_directory = std::getenv("LOCALAPPDATA");
#elif defined(__EMSCRIPTEN__)
GGML_ABORT("not implemented on this platform");
#else
@@ -1069,51 +1061,6 @@ std::string fs_get_cache_directory() {
return ensure_trailing_slash(cache_directory);
}
std::string fs_get_config_directory() {
std::string config_directory = "";
auto ensure_trailing_slash = [](std::string p) {
if (p.empty() || p.back() != DIRECTORY_SEPARATOR) {
p += DIRECTORY_SEPARATOR;
}
return p;
};
#if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \
defined(__OpenBSD__) || defined(__NetBSD__) || defined(__APPLE__)
const std::string xdg_config_home = common_get_env("XDG_CONFIG_HOME");
const std::string home = common_get_env("HOME");
if (!xdg_config_home.empty()) {
config_directory = xdg_config_home;
} else if (!home.empty()) {
config_directory = home + "/.config/";
} else {
#if defined(__linux__)
/* no $HOME is defined, fallback to getpwuid */
struct passwd *pw = getpwuid(getuid());
if ((!pw) || (!pw->pw_dir)) {
throw std::runtime_error("Failed to find $HOME directory");
}
config_directory = std::string(pw->pw_dir) + std::string("/.config/");
#else
throw std::runtime_error("Failed to find $HOME directory");
#endif
}
#elif defined(_WIN32)
config_directory = common_get_env("APPDATA");
if (config_directory.empty()) {
throw std::runtime_error("Failed to find %APPDATA% directory");
}
#elif defined(__EMSCRIPTEN__)
// caller decides what to do when there is no config directory
throw std::runtime_error("not implemented on this platform");
#else
# error Unknown architecture
#endif
config_directory = ensure_trailing_slash(config_directory);
config_directory += "llama.cpp";
return ensure_trailing_slash(config_directory);
}
std::string fs_get_cache_file(const std::string & filename) {
GGML_ASSERT(filename.find(DIRECTORY_SEPARATOR) == std::string::npos);
std::string cache_directory = fs_get_cache_directory();
@@ -1692,7 +1639,6 @@ struct llama_context_params common_context_params_to_llama(const common_params &
cparams.n_seq_max = params.n_parallel;
cparams.n_rs_seq = params.speculative.need_n_rs_seq();
cparams.n_outputs_max = std::max(params.n_outputs_max, 0);
cparams.n_outputs_max_per_seq = std::max(params.n_outputs_max_per_seq, 0);
cparams.n_batch = params.n_batch;
cparams.n_ubatch = params.n_ubatch;
cparams.n_threads = params.cpuparams.n_threads;
+1 -4
View File
@@ -447,7 +447,6 @@ struct common_params {
int32_t n_parallel = 1; // number of parallel sequences to decode
int32_t n_sequences = 1; // number of sequences to decode
int32_t n_outputs_max = 0; // max outputs in a batch (0 = n_batch)
int32_t n_outputs_max_per_seq = 1; // max outputs per sequence
int32_t grp_attn_n = 1; // group-attention factor
int32_t grp_attn_w = 512; // group-attention width
int32_t n_print = -1; // print token count every n tokens (-1 = disabled)
@@ -473,7 +472,7 @@ struct common_params {
std::vector<size_t> fit_params_target = std::vector<size_t>(llama_max_devices(), 1024 * 1024*1024);
enum llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER; // how to split the model across GPUs
enum llama_load_mode load_mode = LLAMA_LOAD_MODE_AUTO; // how to load the model
enum llama_load_mode load_mode = LLAMA_LOAD_MODE_MMAP; // how to load the model
common_cpu_params cpuparams;
common_cpu_params cpuparams_batch;
@@ -656,7 +655,6 @@ struct common_params {
// enable built-in tools
std::vector<std::string> server_tools;
std::string server_tools_runtime;
// MCP server configs (Cursor-compatible JSON)
std::string mcp_servers_config; // path to JSON file with MCP server definitions
@@ -881,7 +879,6 @@ bool fs_is_directory(const std::string & path);
std::string fs_get_cache_directory();
std::string fs_get_cache_file(const std::string & filename);
std::string fs_get_config_directory();
struct common_file_info {
std::string path;
+1 -4
View File
@@ -136,10 +136,7 @@ static std::vector<llama_device_memory_data> common_get_device_memory_data_impl(
devs.push_back(llama_model_get_device(model, i));
}
hp_ngl = llama_model_n_layer(model);
if (mparams->load_mtp) {
hp_ngl += llama_model_n_layer_nextn(model);
}
hp_ngl = llama_model_n_layer(model) + llama_model_n_layer_nextn(model);
hp_n_ctx_train = llama_model_n_ctx_train(model);
hp_n_expert = llama_model_n_expert(model);
-2
View File
@@ -116,8 +116,6 @@ static llama_sampler_i llama_sampler_llg_i = {
/* .backend_accept = */ NULL,
/* .backend_apply = */ NULL,
/* .backend_set_input = */ NULL,
/* .backend_reset = */ NULL,
/* .copy_state = */ NULL,
};
static size_t llama_sampler_llg_tokenize_fn(const void * user_data, const uint8_t * bytes, size_t bytes_len,
+4 -15
View File
@@ -570,34 +570,23 @@ struct parser_executor {
}
static common_peg_parse_result handle_escape_sequence(common_peg_parse_context & ctx, size_t start, size_t & pos, const char delimiter) {
auto save = pos;
++pos; // consume '\'
if (pos >= ctx.input.size()) {
if (!ctx.is_lenient()) {
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
}
pos = save; // suppress unmatched '\'
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_NEED_MORE_INPUT, start, pos);
}
char c = ctx.input[pos];
if (c == delimiter || c == '\\' || c == '/' || c == 'b' || c == 'f' || c == 'n' || c == 'r' || c == 't') {
++pos;
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_SUCCESS, start, pos);
} else if (c == 'u') {
return handle_unicode_escape(ctx, start, pos);
} else {
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
}
if (c == 'u') {
auto result = handle_unicode_escape(ctx, start, pos);
if (result.need_more_input()) {
pos = save; // suppress incomplete sequence
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_NEED_MORE_INPUT, start, pos);
}
return result;
}
return common_peg_parse_result(COMMON_PEG_PARSE_RESULT_FAIL, start);
}
static common_peg_parse_result handle_unicode_escape(common_peg_parse_context & ctx, size_t start, size_t & pos) {
-2
View File
@@ -322,8 +322,6 @@ common_presets common_preset_context::load_from_ini(const std::string & path, co
preset.options[opt] = value;
}
LOG_DBG("accepted option: %s = %s\n", key.c_str(), preset.options[opt].c_str());
} else if (ignore_unknown_keys) {
LOG_WRN("ignoring option '%s' from %s: not supported by this program\n", key.c_str(), path.c_str());
} else {
throw std::runtime_error(string_format(
"option '%s' not recognized in preset '%s'",
-4
View File
@@ -59,10 +59,6 @@ struct common_preset_context {
bool filter_allowed_keys = false;
std::set<std::string> allowed_keys;
// if true, options unknown to the current example are skipped instead of being an error
// used for config files shared by all binaries, where each binary only knows a subset of options
bool ignore_unknown_keys = false;
// if only_remote_allowed is true, only accept whitelisted keys
common_preset_context(llama_example ex);
-2
View File
@@ -217,8 +217,6 @@ static struct llama_sampler_i common_reasoning_budget_i = {
/* .backend_accept = */ nullptr,
/* .backend_apply = */ nullptr,
/* .backend_set_input = */ nullptr,
/* .backend_reset = */ nullptr,
/* .copy_state = */ nullptr,
};
static struct llama_sampler * common_reasoning_budget_clone(const struct llama_sampler * smpl) {
-20
View File
@@ -518,26 +518,6 @@ struct common_sampler * common_sampler_clone(common_sampler * gsmpl) {
};
}
void common_sampler_copy(const common_sampler * src, common_sampler * dst) {
if (!src || !dst || src == dst) {
return;
}
GGML_ASSERT((src->grmr == nullptr) == (dst->grmr == nullptr));
GGML_ASSERT((src->rbudget == nullptr) == (dst->rbudget == nullptr));
llama_sampler_copy(src->grmr, dst->grmr);
llama_sampler_copy(src->rbudget, dst->rbudget);
llama_sampler_copy(src->chain, dst->chain);
dst->params = src->params;
dst->prev = src->prev;
dst->cur = src->cur;
dst->cur_p = src->cur_p;
dst->cur_p.data = src->cur_p.data ? dst->cur.data() : nullptr; // re-point to dst's buffer
dst->t_total_us = src->t_total_us;
}
void common_perf_print(const struct llama_context * ctx, const struct common_sampler * gsmpl) {
// TODO: measure grammar performance
-1
View File
@@ -47,7 +47,6 @@ void common_sampler_free(struct common_sampler * gsmpl);
void common_sampler_accept(struct common_sampler * gsmpl, llama_token token, bool is_generated);
void common_sampler_reset (struct common_sampler * gsmpl);
struct common_sampler * common_sampler_clone (struct common_sampler * gsmpl);
void common_sampler_copy (const struct common_sampler * src, struct common_sampler * dst);
// arguments can be nullptr to skip printing
void common_perf_print(const struct llama_context * ctx, const struct common_sampler * gsmpl);
+73 -28
View File
@@ -171,6 +171,12 @@ struct common_speculative_impl {
// (optional) serialize/restore per-seq internal state (e.g. eagle3's deferred boundary).
virtual bool get_state(llama_seq_id /*seq_id*/, std::vector<uint8_t> & /*data*/) const { return false; }
virtual void set_state(llama_seq_id /*seq_id*/, const std::vector<uint8_t> & /*data*/) {}
// true if this implementation requires the target context to extract post-norm embeddings
virtual bool need_embd() const = 0;
// true if this implementation requires the target context to extract pre-norm embeddings
virtual bool need_embd_nextn() const { return false; }
};
struct common_speculative_impl_draft_simple : public common_speculative_impl {
@@ -187,10 +193,6 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
auto * ctx_dft = this->params.ctx_dft;
auto * ctx_tgt = this->params.ctx_tgt;
if (!ctx_dft) {
throw std::runtime_error("draft-simple requires a draft context");
}
SPC_TRC("%s", "adding speculative implementation 'draft-simple'\n");
SPC_TRC("- n_max=%d, n_min=%d, p_min=%f\n", this->params.n_max, this->params.n_min, this->params.p_min);
SPC_TRC("- gpu_layers=%d, cache_k=%s, cache_v=%s, ctx_tgt=%s, ctx_dft=%s, devices=[%s]\n",
@@ -383,6 +385,10 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
// noop
}
bool need_embd() const override {
return false;
}
};
@@ -901,6 +907,10 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
pending_g_last[seq_id].resize(n_embd_dec);
std::memcpy(pending_g_last[seq_id].data(), data.data() + sizeof(llama_pos), (size_t) n_embd_dec * sizeof(float));
}
bool need_embd() const override {
return false;
}
};
// DFlash: block-diffusion drafting with a draft-side KV cache injection
@@ -1022,14 +1032,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
return true;
}
// Target prefill may contain token IDs or multimodal embeddings. Both
// produce the target-layer features used to seed the draft KV cache, so
// skipping the embedding batches leaves a hole in the draft's cache and
// the next injection fails to initialize.
// TODO: revisit after https://github.com/ggml-org/llama.cpp/pull/24669 is merged
const bool has_tokens = batch_in.token != nullptr;
const bool has_embeddings = batch_in.embd != nullptr;
if (has_tokens == has_embeddings) {
if (batch_in.token == nullptr || batch_in.embd != nullptr) {
return true;
}
@@ -1237,6 +1240,10 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
// noop
}
bool need_embd() const override {
return false;
}
};
struct common_speculative_impl_draft_mtp : public common_speculative_impl {
@@ -1675,6 +1682,14 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
const size_t row_bytes = (size_t) n_embd * sizeof(float);
std::memcpy(pending_h[seq_id].data(), verify_h[seq_id].data() + (size_t) i_h * n_embd, row_bytes);
}
bool need_embd() const override {
return false;
}
bool need_embd_nextn() const override {
return true;
}
};
// state of self-speculation (simple implementation, not ngram-map)
@@ -1721,6 +1736,10 @@ struct common_speculative_impl_ngram_simple : public common_speculative_impl {
void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
// noop
}
bool need_embd() const override {
return false;
}
};
struct common_speculative_impl_ngram_map_k : public common_speculative_impl {
@@ -1775,6 +1794,10 @@ struct common_speculative_impl_ngram_map_k : public common_speculative_impl {
common_ngram_map_accept(config[seq_id], n_accepted);
}
bool need_embd() const override {
return false;
}
};
struct common_speculative_impl_ngram_mod : public common_speculative_impl {
@@ -1950,6 +1973,10 @@ struct common_speculative_impl_ngram_mod : public common_speculative_impl {
}
}
}
bool need_embd() const override {
return false;
}
};
struct common_speculative_impl_ngram_cache : public common_speculative_impl {
@@ -2089,6 +2116,10 @@ struct common_speculative_impl_ngram_cache : public common_speculative_impl {
void accept(llama_seq_id /*seq_id*/, uint16_t /*n_accepted*/, bool /*is_other*/) override {
// noop
}
bool need_embd() const override {
return false;
}
};
struct common_speculative {
@@ -2261,7 +2292,6 @@ common_params common_base_params_to_speculative(const common_params & params) {
result.cache_type_k = params_spec.cache_type_k;
result.cache_type_v = params_spec.cache_type_v;
result.n_outputs_max = params.n_parallel;
result.n_outputs_max_per_seq = 1;
return result;
}
@@ -2284,6 +2314,7 @@ common_speculative_init_result::common_speculative_init_result(
const bool spec_mtp = std::find(params.speculative.types.begin(),
params.speculative.types.end(),
COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end();
GGML_ASSERT(has_draft || spec_mtp);
auto mparams = common_model_params_to_llama(params);
auto cparams = common_context_params_to_llama(params);
@@ -2346,17 +2377,6 @@ common_speculative_init_result_ptr common_speculative_init_from_params(common_pa
return std::make_unique<common_speculative_init_result>(params, model_tgt, ctx_tgt);
}
common_speculative_output_limits common_speculative_get_output_limits(
int32_t n_batch, int32_t n_parallel, int32_t n_draft) {
const int64_t per_seq = 1 + (int64_t) std::max(0, n_draft);
const int64_t total = (int64_t) n_parallel * per_seq;
return {
/* .total = */ (int32_t) std::min<int64_t>(n_batch, total),
/* .per_seq = */ (int32_t) std::min<int64_t>(n_batch, per_seq),
};
}
// initialization of the speculative decoding system
//
common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq) {
@@ -2521,6 +2541,34 @@ bool common_speculative_process(common_speculative * spec, const llama_batch & b
return result;
}
bool common_speculative_need_embd(common_speculative * spec) {
if (spec == nullptr) {
return false;
}
for (auto & impl : spec->impls) {
if (impl->need_embd()) {
return true;
}
}
return false;
}
bool common_speculative_need_embd_nextn(common_speculative * spec) {
if (spec == nullptr) {
return false;
}
for (auto & impl : spec->impls) {
if (impl->need_embd_nextn()) {
return true;
}
}
return false;
}
void common_speculative_draft(common_speculative * spec) {
if (spec == nullptr) {
return;
@@ -2605,10 +2653,7 @@ void common_speculative_draft(common_speculative * spec) {
void common_speculative_accept(common_speculative * spec, llama_seq_id seq_id, uint16_t n_accepted) {
common_speculative_impl * impl = spec->impl_last[seq_id];
if (impl == nullptr) {
GGML_ASSERT(n_accepted == 0);
return;
}
GGML_ASSERT(impl);
{
common_time_meas tm(impl->t_accept_us, !impl->gen_perf);
+6 -9
View File
@@ -25,15 +25,6 @@ int32_t common_speculative_n_max(const common_params_speculative * spec);
common_params common_base_params_to_speculative(const common_params & params);
struct common_speculative_output_limits {
int32_t total;
int32_t per_seq;
};
// return the output limits needed for speculative decoding
common_speculative_output_limits common_speculative_get_output_limits(
int32_t n_batch, int32_t n_parallel, int32_t n_draft);
common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq);
void common_speculative_free(common_speculative * spec);
@@ -67,6 +58,12 @@ void common_speculative_begin(common_speculative * spec, llama_seq_id seq_id, co
// process the batch and update the internal state of the speculative context
bool common_speculative_process(common_speculative * spec, const llama_batch & batch);
// true if any implementation requires target post-norm embeddings to be extracted
bool common_speculative_need_embd(common_speculative * spec);
// true if any implementation requires target nextn embeddings to be extracted
bool common_speculative_need_embd_nextn(common_speculative * spec);
// generate drafts for the sequences specified with `common_speculative_get_draft_params`
void common_speculative_draft(common_speculative * spec);
-7
View File
@@ -70,7 +70,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Exaone4ForCausalLM": "exaone",
"ExaoneForCausalLM": "exaone",
"ExaoneMoEForCausalLM": "exaone",
"ExaoneMoeForCausalLM": "exaone",
"FalconForCausalLM": "falcon",
"FalconH1ForCausalLM": "falcon_h1",
"FalconMambaForCausalLM": "mamba",
@@ -103,7 +102,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"GraniteMoeForCausalLM": "granite",
"GraniteMoeHybridForCausalLM": "granite",
"GraniteMoeSharedForCausalLM": "granite",
"GraniteSwitchForCausalLM": "granite",
"GraniteSpeechForConditionalGeneration": "granite",
"GraniteSpeechPlusForConditionalGeneration": "granite",
"Grok1ForCausalLM": "grok",
@@ -183,8 +181,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Olmo3ForCausalLM": "olmo",
"OlmoForCausalLM": "olmo",
"OlmoeForCausalLM": "olmo",
"MuseGlimmerAssistantModel": "muse_glimmer",
"MuseGlimmerForConditionalGeneration": "muse_glimmer",
"OpenELMForCausalLM": "openelm",
"OrionForCausalLM": "orion",
"PLMForCausalLM": "plm",
@@ -214,7 +210,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Qwen3MoeForCausalLM": "qwen",
"Qwen3NextForCausalLM": "qwen",
"Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
"PocketTTSModel": "pockettts",
"Qwen3TTSForConditionalGeneration": "qwen3tts",
"Qwen3VLForConditionalGeneration": "qwen3vl",
"Qwen3VLMoeForConditionalGeneration": "qwen3vl",
@@ -301,7 +296,6 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
"MiniCPMV4_6ForConditionalGeneration": "minicpm",
"Mistral3ForConditionalGeneration": "llava",
"NemotronH_Nano_VL_V2": "nemotron",
"MuseGlimmerForConditionalGeneration": "muse_glimmer",
"PaddleOCRVisionModel": "ernie",
"Phi4ForCausalLMV": "phi",
"Qwen2AudioForConditionalGeneration": "ultravox",
@@ -311,7 +305,6 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
"Qwen2_5_VLForConditionalGeneration": "qwenvl",
"Qwen3ASRForConditionalGeneration": "qwen3vl",
"Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
"PocketTTSModel": "pockettts",
"Qwen3TTSForConditionalGeneration": "qwen3tts",
"Qwen3VLForConditionalGeneration": "qwen3vl",
"Qwen3VLMoeForConditionalGeneration": "qwen3vl",
+1 -29
View File
@@ -58,11 +58,6 @@ logger = logging.getLogger("hf-to-gguf")
AnyModel = TypeVar("AnyModel", bound="type[ModelBase]")
# for checkpoints that ship no config.json, we will try to provide a synthetic one
HparamsMatcher = Callable[[Path], bool]
HparamsLoader = Callable[[Path], dict[str, Any]]
class SentencePieceTokenTypes(IntEnum):
NORMAL = 1
UNKNOWN = 2
@@ -82,7 +77,6 @@ class ModelBase:
ModelType.TEXT: {},
ModelType.MMPROJ: {},
}
_hparams_loaders: list[tuple[HparamsMatcher, HparamsLoader]] = []
dir_model: Path
ftype: gguf.LlamaFileType
@@ -829,7 +823,7 @@ class ModelBase:
elif any(str(v.get("quant_algo")).endswith("NVFP4") for v in quant_layers.values() if isinstance(v, dict)):
quant_algo = "NVFP4"
self._is_nvfp4 = quant_algo in ("NVFP4", "W4A16_NVFP4")
self._is_nvfp4 = quant_algo == "NVFP4"
self._is_mxfp4 = quant_method == "mxfp4"
# NVFP4 weights are repacked and written directly to gguf_writer.
@@ -1046,24 +1040,6 @@ class ModelBase:
return part_names
@staticmethod
def load_hparams_guess(dir_model: Path) -> dict[str, Any] | None:
# some models ship no config.json, will try to guess them
from conversion import load_all_models
load_all_models()
for matcher, loader in ModelBase._hparams_loaders:
if matcher(dir_model):
return loader(dir_model)
return None
@classmethod
def register_hparams_loader(cls, matcher: HparamsMatcher) -> Callable[[HparamsLoader], HparamsLoader]:
def inner(loader: HparamsLoader) -> HparamsLoader:
cls._hparams_loaders.append((matcher, loader))
return loader
return inner
@staticmethod
def load_hparams(dir_model: Path, is_mistral_format: bool):
if is_mistral_format:
@@ -1077,10 +1053,6 @@ class ModelBase:
config = AutoConfig.from_pretrained(dir_model, trust_remote_code=False).to_dict()
except Exception as e:
logger.warning(f"Failed to load model config from {dir_model}: {e}")
if not (dir_model / "config.json").is_file():
config = ModelBase.load_hparams_guess(dir_model)
if config is not None:
return config
logger.warning("Trying to load config.json instead")
with open(dir_model / "config.json", "r", encoding="utf-8") as f:
config = json.load(f)
+1 -21
View File
@@ -17,11 +17,8 @@ from .base import LazyTorchTensor, MmprojModel, ModelBase, TextModel, gguf, logg
from .qwen import QwenModel
@ModelBase.register("DeepseekOCRForCausalLM")
@ModelBase.register("DeepseekOCRForCausalLM", "UnlimitedOCRForCausalLM")
class DeepseekOCRVisionModel(MmprojModel):
# HF dynamic_preprocess() max_num, which differs per model
preproc_max_tiles = 9
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
self.clip_projector_type = gguf.VisionProjectorType.DEEPSEEKOCR
@@ -46,9 +43,6 @@ class DeepseekOCRVisionModel(MmprojModel):
# @bluebread: there's no window_size in config but just add it here anyway
self.gguf_writer.add_vision_window_size(self.hparams.get("window_size", 14))
self.gguf_writer.add_vision_preproc_min_tiles(2)
self.gguf_writer.add_vision_preproc_max_tiles(self.preproc_max_tiles)
# SAM configuration
sam_hparams = hparams['sam']
self.gguf_writer.add_vision_sam_layers_count(sam_hparams['layers'])
@@ -99,15 +93,8 @@ class DeepseekOCRVisionModel(MmprojModel):
return super().filter_tensors((name, gen))
@ModelBase.register("UnlimitedOCRForCausalLM")
class UnlimitedOCRVisionModel(DeepseekOCRVisionModel):
preproc_max_tiles = 32
@ModelBase.register("DeepseekOCR2ForCausalLM")
class DeepseekOCR2VisionModel(DeepseekOCRVisionModel):
preproc_max_tiles = 6
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
self.clip_projector_type = gguf.VisionProjectorType.DEEPSEEKOCR2
@@ -533,13 +520,6 @@ class DeepseekV4Model(TextModel):
for key, value in raw_hparams.items():
self.hparams.setdefault(key, value)
# workaround for special rope_parameters (main/compress) in transformers 5.x
if self.rope_parameters.get("full_attention", self.rope_parameters).get("rope_type") is None:
if (rope_scaling := raw_hparams.get("rope_scaling")) is not None:
if "rope_type" not in rope_scaling and (rope_type := rope_scaling.get("type")) is not None:
rope_scaling["rope_type"] = rope_type
self.rope_parameters.update(**rope_scaling)
self.block_count = self.hparams["num_hidden_layers"]
if self.mtp_only:
self.block_count += self.hparams.get("num_nextn_predict_layers", 0)
+1 -3
View File
@@ -123,9 +123,7 @@ class Exaone4Model(TextModel):
yield (self.format_tensor_name(gguf.MODEL_TENSOR.ROPE_FREQS), torch.tensor(rope_factors, dtype=torch.float32))
# note: transformers >= 5.1 renamed the class to "ExaoneMoeForCausalLM" (lowercase 'e'),
# so accept both spellings - LG AI have updated the configs of already-released models
@ModelBase.register("ExaoneMoEForCausalLM", "ExaoneMoeForCausalLM")
@ModelBase.register("ExaoneMoEForCausalLM")
class ExaoneMoEModel(Exaone4Model):
model_arch = gguf.MODEL_ARCH.EXAONE_MOE
+4 -33
View File
@@ -665,18 +665,7 @@ class Gemma4Model(Gemma3Model):
swa_layers = [t == "sliding_attention" for t in self.hparams["layer_types"]]
self.gguf_writer.add_sliding_window_pattern(swa_layers)
per_layer_config = self.hparams.get("per_layer_config")
layer_types = self.hparams.get("layer_types", [])
if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
for layer_idx, layer_config in per_layer_config.items():
layer_idx = int(layer_idx)
if layer_idx < len(layer_types):
if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
head_dim_full = layer_config["head_dim"]
break
assert head_dim_full is not None
head_dim_full = self.hparams["global_head_dim"]
head_dim_swa = self.hparams["head_dim"]
# correct the head dim for global/swa layers
self.gguf_writer.add_key_length(head_dim_full)
@@ -696,14 +685,8 @@ class Gemma4Model(Gemma3Model):
n_ff_arr = [n_ff if il < first_kv_shared_layer_idx else n_ff * 2 for il in range(self.block_count)]
self.gguf_writer.add_feed_forward_length(n_ff_arr)
if (num_key_value_heads_full := self.hparams.get("num_global_key_value_heads")) is None and per_layer_config is not None:
for layer_idx, layer_config in per_layer_config.items():
layer_idx = int(layer_idx)
if layer_idx < len(layer_types):
if layer_types[layer_idx] == "full_attention" and "num_key_value_heads" in layer_config:
num_key_value_heads_full = layer_config["num_key_value_heads"]
break
# handle num_global_key_value_heads
num_key_value_heads_full = self.hparams.get("num_global_key_value_heads")
num_key_value_heads_swa = self.hparams.get("num_key_value_heads")
if num_key_value_heads_full is not None and num_key_value_heads_swa is not None:
value_arr = [num_key_value_heads_swa if is_swa else num_key_value_heads_full for is_swa in swa_layers]
@@ -725,19 +708,7 @@ class Gemma4Model(Gemma3Model):
# IMPORTANT: this ROPE_FREQS tensor is ONLY used by the full_attention layers
rope_params_full = self.hparams["rope_parameters"]["full_attention"]
assert rope_params_full["rope_type"] == "proportional"
per_layer_config = self.hparams.get("per_layer_config")
if (head_dim_full := self.hparams.get("global_head_dim")) is None and per_layer_config is not None:
layer_types = self.hparams.get("layer_types", [])
for layer_idx, layer_config in per_layer_config.items():
layer_idx = int(layer_idx)
if layer_idx < len(layer_types):
if layer_types[layer_idx] == "full_attention" and "head_dim" in layer_config:
head_dim_full = layer_config["head_dim"]
break
assert head_dim_full is not None
head_dim_full = (self.hparams["global_head_dim"])
partial_rotary_factor_full = rope_params_full["partial_rotary_factor"]
n_rot_full = int(head_dim_full * partial_rotary_factor_full / 2)
n_unrot_full = int(head_dim_full / 2) - n_rot_full
-160
View File
@@ -123,166 +123,6 @@ class GraniteMoeModel(GraniteModel):
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("GraniteSwitchForCausalLM")
class GraniteSwitchModel(GraniteMoeModel):
"""Dense, all-attention Granite with N per-token embedded LoRA adapters, stacked
over the adapter dim with a zero adapter at slot 0 (N = num_adapters + 1)."""
model_arch = gguf.MODEL_ARCH.GRANITE_SWITCH
# permute q/k per-slice below (NORM-rope layout), not via the parent's auto-permute
undo_permute = False
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
# the weightless switch reserves one cache slot: one fewer block than num_hidden_layers
self.block_count = self.block_count - 1
self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
self._n_adapters = int(self.hparams["num_adapters"])
self._max_lora_rank = int(self.hparams["max_lora_rank"])
self._n_slots = self._n_adapters + 1 # +1 for the zero slot at index 0
n_head = int(self.hparams["num_attention_heads"])
n_kv_head = int(self.hparams["num_key_value_heads"])
head_dim = (
self.hparams.get("projection_head_dim")
or self.hparams.get("head_dim")
or (self.hparams["hidden_size"] // n_head)
)
self._n_head = n_head
self._n_kv_head = n_kv_head
self._head_dim = int(head_dim)
self._q_size = n_head * self._head_dim
self._kv_size = n_kv_head * self._head_dim
def set_gguf_parameters(self):
super().set_gguf_parameters()
# dense: pin expert_used_count to 0 (config carries a leftover num_experts_per_tok)
if not self.hparams.get("num_local_experts"):
self.gguf_writer.add_expert_used_count(0)
self.gguf_writer.add_adapter_count(self._n_adapters)
self.gguf_writer.add_adapter_lora_rank(self._max_lora_rank)
self.gguf_writer.add_adapter_token_ids_activate(self.hparams["adapter_token_ids"])
self.gguf_writer.add_adapter_token_ids_substitute(self.hparams["adapter_substitute_token_ids"])
router_gain = float(self.hparams.get("control_token_gain", 15.0))
self.gguf_writer.add_adapter_router_gain(router_gain)
logger.info("gguf: (graniteswitch) num_adapters=%s max_lora_rank=%s n_slots=%s router_gain=%s", self._n_adapters, self._max_lora_rank, self._n_slots, router_gain)
def _lora_a(self, data: Tensor) -> Tensor:
# on-disk A: [n_adapters, 1, max_rank, in] -> [n_adapters+1, max_rank, in]
a = data.squeeze(1)
zero = torch.zeros_like(a[:1])
return torch.cat([zero, a], dim=0).contiguous()
def _lora_b(self, data: Tensor, permute_n_head: int | None = None) -> Tensor:
# on-disk B: [n_adapters, 1, out, max_rank] -> [n_adapters+1, out, max_rank]
b = data.squeeze(1)
if permute_n_head is not None:
# permute each adapter's B output rows to match the permuted q/k base
b = torch.stack([self.permute(b[i], permute_n_head, permute_n_head) for i in range(b.shape[0])], dim=0)
zero = torch.zeros_like(b[:1])
return torch.cat([zero, b], dim=0).contiguous()
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
T = gguf.MODEL_TENSOR
# skip the weightless switch + control-token buffers (rebuilt at load time)
bare = name.split(".")[-1]
if (
name.startswith("model.switch.") or name.startswith("switch.")
or bare in ("adapter_token_ids", "control_to_substitute_lut")
):
return
if "self_attn.qkv_proj" in name:
if name.endswith("base_layer.weight"):
# fused [q|k|v] rows: permute q/k row-blocks for ggml's NORM-rope layout
q, k, v = data_torch.split([self._q_size, self._kv_size, self._kv_size], dim=0)
q = self.permute(q, self._n_head, self._n_head)
k = self.permute(k, self._n_kv_head, self._n_kv_head)
fused = torch.cat([q, k, v], dim=0)
yield (self.format_tensor_name(T.ATTN_QKV, bid), fused)
return
if "lora_A_slices." in name:
slot = int(name.rsplit(".", 1)[1])
key = {0: T.ATTN_Q, 1: T.ATTN_K, 2: T.ATTN_V}[slot]
yield (self.format_tensor_name(key, bid, suffix=".lora_a"), self._lora_a(data_torch))
return
if "lora_B_slices." in name:
slot = int(name.rsplit(".", 1)[1])
key, ph = {
0: (T.ATTN_Q, self._n_head),
1: (T.ATTN_K, self._n_kv_head),
2: (T.ATTN_V, None),
}[slot]
yield (self.format_tensor_name(key, bid, suffix=".lora_b"), self._lora_b(data_torch, ph))
return
raise ValueError(f"Unexpected qkv_proj tensor: {name}")
if "self_attn.o_proj" in name:
if name.endswith("base_layer.weight"):
yield (self.format_tensor_name(T.ATTN_OUT, bid), data_torch)
return
if name.endswith("lora_A"):
yield (self.format_tensor_name(T.ATTN_OUT, bid, suffix=".lora_a"), self._lora_a(data_torch))
return
if name.endswith("lora_B"):
yield (self.format_tensor_name(T.ATTN_OUT, bid, suffix=".lora_b"), self._lora_b(data_torch))
return
raise ValueError(f"Unexpected o_proj tensor: {name}")
if "shared_mlp.input_linear" in name:
ffn = self.hparams["shared_intermediate_size"]
if name.endswith("base_layer.weight"):
gate, up = data_torch.split([ffn, ffn], dim=0)
yield (self.format_tensor_name(T.FFN_GATE, bid), gate)
yield (self.format_tensor_name(T.FFN_UP, bid), up)
return
if "lora_A_slices." in name:
slot = int(name.rsplit(".", 1)[1])
key = {0: T.FFN_GATE, 1: T.FFN_UP}[slot]
yield (self.format_tensor_name(key, bid, suffix=".lora_a"), self._lora_a(data_torch))
return
if "lora_B_slices." in name:
slot = int(name.rsplit(".", 1)[1])
key = {0: T.FFN_GATE, 1: T.FFN_UP}[slot]
yield (self.format_tensor_name(key, bid, suffix=".lora_b"), self._lora_b(data_torch))
return
raise ValueError(f"Unexpected shared_mlp.input_linear tensor: {name}")
if "shared_mlp.output_linear" in name:
if name.endswith("base_layer.weight"):
yield (self.format_tensor_name(T.FFN_DOWN, bid), data_torch)
return
if name.endswith("lora_A"):
yield (self.format_tensor_name(T.FFN_DOWN, bid, suffix=".lora_a"), self._lora_a(data_torch))
return
if name.endswith("lora_B"):
yield (self.format_tensor_name(T.FFN_DOWN, bid, suffix=".lora_b"), self._lora_b(data_torch))
return
raise ValueError(f"Unexpected shared_mlp.output_linear tensor: {name}")
if bid is not None and ".layers." in name and (
"input_layernorm" in name or "post_attention_layernorm" in name
):
key = T.ATTN_NORM if "input_layernorm" in name else T.FFN_NORM
yield (self.format_tensor_name(key, bid), data_torch)
return
if name in ("model.embed_tokens.weight", "embed_tokens.weight"):
yield (self.format_tensor_name(T.TOKEN_EMBD), data_torch)
return
if name in ("model.norm.weight", "norm.weight"):
yield (self.format_tensor_name(T.OUTPUT_NORM), data_torch)
return
if name == "lm_head.weight":
return # tied to token_embd
raise ValueError(f"graniteswitch: unhandled tensor {name!r} (bid={bid})")
@ModelBase.register("GraniteMoeHybridForCausalLM", "BambaForCausalLM")
class GraniteHybridModel(Mamba2Model, GraniteMoeModel):
"""GraniteHybrid is a hybrid SSM + Attention model that uses Mamba2 SSM
-179
View File
@@ -1,179 +0,0 @@
from __future__ import annotations
import json
from typing import Any, Iterable, TYPE_CHECKING
import torch
if TYPE_CHECKING:
from torch import Tensor
from .base import MmprojModel, ModelBase, TextModel, gguf
def _unpermute_for_rope(tensor: "Tensor", n_heads: int) -> "Tensor":
"""Invert transformers' `_permute_for_rope`: HF stores Q/K in rotate_half layout,
llama.cpp consumes the interleaved (NORM) layout."""
if tensor.ndim == 2:
dim1, dim2 = tensor.shape
return tensor.view(n_heads, 2, dim1 // n_heads // 2, dim2).transpose(1, 2).reshape(dim1, dim2)
if tensor.ndim == 1:
(dim1,) = tensor.shape
return tensor.view(n_heads, 2, dim1 // n_heads // 2).transpose(1, 2).reshape(dim1)
raise ValueError(f"_unpermute_for_rope: unexpected shape {tuple(tensor.shape)}")
@ModelBase.register("MuseGlimmerForConditionalGeneration")
class MuseGlimmerModel(TextModel):
model_arch = gguf.MODEL_ARCH.MUSE_GLIMMER
def norm_shift(self, name: str) -> float:
# All four layer norms use 1, the final norm uses 0.
return 1.0 if name.endswith("layernorm.weight") else 0.0
def set_vocab(self):
self._set_vocab_gpt2()
from transformers import AutoTokenizer
tok = AutoTokenizer.from_pretrained(self.dir_model)
eot_id = tok.convert_tokens_to_ids("<|eot|>")
if isinstance(eot_id, int) and eot_id >= 0:
self.gguf_writer.add_eot_token_id(eot_id)
def set_gguf_parameters(self):
super().set_gguf_parameters()
hparams = self.hparams
self.gguf_writer.add_final_logit_softcapping(hparams["final_logit_softcapping"])
self.gguf_writer.add_logit_scale(hparams["output_multiplier"])
self.gguf_writer.add_sliding_window(hparams["sliding_window"])
self.gguf_writer.add_sliding_window_pattern([t == "sliding_attention" for t in hparams["layer_types"]])
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
shift = self.norm_shift(name)
if shift != 0.0:
data_torch = data_torch + shift
# Invert transformers' `_permute_for_rope` on Q/K, we keep ggml's NORM (interleaved) rope
if ".self_attn.q_proj." in name:
data_torch = _unpermute_for_rope(data_torch, int(self.hparams["num_attention_heads"]))
elif ".self_attn.k_proj." in name:
data_torch = _unpermute_for_rope(data_torch, int(self.hparams["num_key_value_heads"]))
# Synthesize QK-norm weights to absorb qk_scale_factor.
# MuseGlimmer implementation: scaleless RMSNorm followed by qk_scale_factor..
if bid is not None and name.endswith(f"model.layers.{bid}.self_attn.q_proj.weight"):
head_dim = self.hparams["head_dim"]
q_scale = float(self.hparams["qk_scale_factor"])
yield (
self.map_tensor_name(f"model.layers.{bid}.self_attn.q_norm.weight"),
torch.full((head_dim,), q_scale, dtype=torch.float32),
)
yield (
self.map_tensor_name(f"model.layers.{bid}.self_attn.k_norm.weight"),
torch.ones((head_dim,), dtype=torch.float32),
)
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("MuseGlimmerForConditionalGeneration")
class MuseGlimmerVisionModel(MmprojModel):
def get_vision_config(self) -> dict[str, Any] | None:
c = self.global_config.get("vision_config")
if not c:
return None
# MuseGlimmer actually uses dynamic size, initialize with nominal size
image_size = c["pos_emb_height"] * c["patch_size"] * c["merge_size"]
return {**c, "image_size": image_size}
def set_gguf_parameters(self):
super().set_gguf_parameters()
assert self.hparams_vision is not None
c = self.hparams_vision # enriched vision_config from get_vision_config()
self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.MUSE_GLIMMER)
self.gguf_writer.add_vision_attention_layernorm_eps(float(c["layer_norm_eps"]))
self.gguf_writer.add_vision_spatial_merge_size(int(c["merge_size"]))
@classmethod
def filter_tensors(cls, item):
name, gen = item
keep = ("model.vision_tower.", "model.vision_adapter.", "model.vision_projection.")
if not any(name.startswith(k) for k in keep):
return None
return super().filter_tensors((name, gen))
# 3-layer projector MLP
_MM_MLP_MAP = {
"model.vision_adapter.fc1": (gguf.MODEL_TENSOR.V_MMPROJ, 0),
"model.vision_adapter.fc2": (gguf.MODEL_TENSOR.V_MMPROJ, 1),
"model.vision_projection": (gguf.MODEL_TENSOR.V_MMPROJ, 2),
}
def modify_tensors(self, data_torch, name, bid):
assert self.hparams_vision is not None
if ".attn.q_proj." in name or ".attn.k_proj." in name:
n_heads = int(self.hparams_vision["num_attention_heads"])
data_torch = _unpermute_for_rope(data_torch, n_heads)
# Lay out the pt=2 temporal slabs of the patch embedding as a conv2d for build_inp()
if name.endswith("patch_embedder.patch_embedding.weight"):
n_embd = data_torch.shape[0]
pt = int(self.hparams_vision["patch_temporal"])
ps = int(self.hparams_vision["patch_size"])
data_torch = data_torch.view(n_embd, pt, 3, ps, ps).sum(dim=1) # (n_embd, 3, ps, ps)
stem, _, suffix = name.rpartition(".")
if stem in self._MM_MLP_MAP:
tensor_key, idx = self._MM_MLP_MAP[stem]
yield (self.format_tensor_name(tensor_key, bid=idx, suffix="." + suffix), data_torch)
return
yield (self.map_tensor_name(name), data_torch)
@ModelBase.register("MuseGlimmerAssistantModel")
class MuseGlimmerAssistantModel(TextModel):
model_arch = gguf.MODEL_ARCH.DFLASH
def set_vocab(self):
if self.target_model_dir is None:
raise ValueError(
"MuseGlimmerAssistant (DFlash drafter) requires --target-model-dir pointing to the "
"target MuseGlimmer HF directory"
)
original_dir = self.dir_model
self.dir_model = self.target_model_dir
from . import get_model_class
with open(self.target_model_dir / "config.json", "r", encoding="utf-8") as f:
target_arch = json.load(f)["architectures"][0]
target_cls = get_model_class(target_arch)
if target_cls is not type(self):
target_cls.set_vocab(self) # ty: ignore[unresolved-attribute]
else:
super().set_vocab()
self.dir_model = original_dir
mask_token_id = self.hparams.get("mask_token_id")
if mask_token_id is not None:
self.gguf_writer.add_mask_token_id(int(mask_token_id))
def set_gguf_parameters(self):
super().set_gguf_parameters()
h = self.hparams
self.gguf_writer.add_block_size(int(h["block_size"]))
# dflash.target_layers[k] refers to the inputs going into the ith layer, which come from the (i-1)th layer's output.
# The transformers configuration refers to the outputs being recorded.
self.gguf_writer.add_target_layers([int(x) + 1 for x in h["target_layer_ids"]])
if h.get("sliding_window") and h.get("layer_types"):
self.gguf_writer.add_sliding_window(int(h["sliding_window"]))
self.gguf_writer.add_sliding_window_pattern([t == "sliding_attention" for t in h["layer_types"]])
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
# DFlash defaults to NEOX (rotate_half) rope, matching transformers HF layout for Q/K, QK-norms
# no permutation needed.
yield (self.map_tensor_name(name), data_torch)
+7 -78
View File
@@ -197,7 +197,6 @@ class NemotronHModel(GraniteHybridModel):
"""Hybrid mamba2/attention model from NVIDIA"""
model_arch = gguf.MODEL_ARCH.NEMOTRON_H
is_moe: bool = False
supports_mtp_export = True
def __init__(self, *args, **kwargs):
# We have to determine the correct model architecture (MoE vs non-MoE) before
@@ -237,25 +236,6 @@ class NemotronHModel(GraniteHybridModel):
self._ssm_layers = [i for i, val in enumerate(pattern) if val == "mamba"]
self._mlp_layers = [i for i, val in enumerate(pattern) if val == "moe"]
# `--no-mtp` drops it entirely; `--mtp` exports only the MTP head
self._mtp_bid: int | None = None
if self.is_moe and not self.no_mtp:
n_nextn = self.hparams.get("num_nextn_predict_layers", 0) or 0
if n_nextn > 0:
assert n_nextn == 1, (
"NemotronH MTP conversion currently supports num_nextn_predict_layers == 1"
)
self._mtp_bid = self.block_count
self.block_count += 1
# The folded MTP block carries both an attention sub-layer and a
# MoE sub-layer, so register it as both so the per-layer metadata arrays cover it
self._attn_layers.append(self._mtp_bid)
self._mlp_layers.append(self._mtp_bid)
self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
if self.mtp_only and self._mtp_bid is None:
raise ValueError("--mtp was requested, but this model does not contain a supported MTP head")
def get_attn_layers(self):
pattern = self.hparams.get("hybrid_override_pattern") or self.hparams.get("layers_block_type")
if pattern is None:
@@ -266,44 +246,6 @@ class NemotronHModel(GraniteHybridModel):
return [i for i, val in enumerate(pattern) if val == "attention"]
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if name.startswith("mtp."):
# --no-mtp: drop the MTP head entirely
if cls.no_mtp:
return None
elif cls.mtp_only:
# --mtp: export the MTP head plus the tensors it shares with the target model
# Include lm_head scale sidecars so NVFP4 packing sees them.
keep = name in (
"backbone.embeddings.weight",
"backbone.norm_f.weight",
"lm_head.weight",
"lm_head.weight_scale",
"lm_head.weight_scale_2",
"lm_head.weight_scale_inv",
"lm_head.input_scale",
"lm_head.input_global_scale",
"lm_head.weight_global_scale",
"lm_head.weight_packed",
)
if not keep:
return None
return super().filter_tensors((name, gen))
def prepare_metadata(self, vocab_only: bool):
from_dir = self.fname_out.is_dir()
super().prepare_metadata(vocab_only=vocab_only)
if not self.mtp_only or not from_dir:
return
output_type: str = self.ftype.name.partition("_")[2]
fname_default: str = gguf.naming_convention(
self.metadata.name, self.metadata.basename, self.metadata.finetune,
self.metadata.version, size_label=None, output_type=output_type, model_type=None)
self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf"
def set_gguf_parameters(self):
super().set_gguf_parameters()
@@ -342,10 +284,6 @@ class NemotronHModel(GraniteHybridModel):
if (latent_size := self.hparams.get("moe_latent_size")) is not None:
self.gguf_writer.add_moe_latent_size(latent_size)
# MTP head: number of trailing NextN blocks
if self._mtp_bid is not None:
self.gguf_writer.add_nextn_predict_layers(self.hparams["num_nextn_predict_layers"])
def set_vocab(self):
# The NemotronH config uses pattern characters (e.g. '-') that may not
# be supported by the installed transformers version. AutoTokenizer
@@ -412,24 +350,15 @@ class NemotronHModel(GraniteHybridModel):
if not self.is_moe:
self.gguf_writer.add_add_bos_token(True)
_MTP_SPECIAL_RENAMES = {
"mtp.layers.0.enorm.weight": "model.layers.{bid}.enorm.weight",
"mtp.layers.0.hnorm.weight": "model.layers.{bid}.hnorm.weight",
"mtp.layers.0.eh_proj.weight": "model.layers.{bid}.eh_proj.weight",
"mtp.layers.1.norm.weight": "model.layers.{bid}.post_attention_layernorm.weight",
"mtp.layers.1.final_layernorm.weight": "model.layers.{bid}.shared_head.norm.weight",
}
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
# mtp.layers.0: NextN input fusion + attention
# mtp.layers.1: MoE + final head norm
if self._mtp_bid is not None and name.startswith(("mtp.layers.0.", "mtp.layers.1.")):
suffix = name.split(".", 3)[3]
bid = self._mtp_bid
renamed = self._MTP_SPECIAL_RENAMES.get(name)
name = renamed.format(bid=bid) if renamed else f"backbone.layers.{bid}.{suffix}"
if self.is_moe and bid is not None:
# Skip Multi-Token Prediction (MTP) tensors. These are used for
# for speculative decoding but we don't include them in this model
# conversion. See https://github.com/ggml-org/llama.cpp/pull/18886
if name.startswith("mtp."):
logger.info(f"gguf: Skipping MTP (Speculative) layer: {name}")
return
if name.endswith("mixer.gate.e_score_correction.bias"):
yield from ModelBase.modify_tensors(self, data_torch, name, bid)
return
-378
View File
@@ -1,378 +0,0 @@
from __future__ import annotations
import re
from pathlib import Path
from typing import Any, Iterable, TYPE_CHECKING
import torch
if TYPE_CHECKING:
from torch import Tensor
from .base import ModelBase, MmprojModel, SentencePieceTokenTypes, TextModel, gguf, logger
# Pocket TTS is a CALM: the backbone conditions a flow-matching decoder that generates one
# continuous 32-d latent per frame. There is no codebook in this model.
# The checkpoint ships no config.json, hparams come from _load_hparams() below.
#
# Tricks being used to support this model via existing llama.cpp code paths:
# - bos_before_voice and bos_emb are learned input vectors, not tokens
# they are appended to the embedding table as extra tokens, to be looked up like any other row
# - bos_emb lives in latent space, so input_linear is folded into it here
# - the backbone has no lm_head, the embedding table is reused as output for the unused logits
#
# pipeline stage mapping:
# mimi encoder + speaker_proj --> mapped to normal mtmd audio encoder
# flow_lm.transformer --> mapped to normal libllama text model (autoregressive)
# flow_lm.flow_net + out_eos --> MTMD_GEN_PROCESS_TYPE_GEN_CODE
# mimi decoder --> MTMD_GEN_PROCESS_TYPE_GEN_WAV
# indices into mimi.encoder.model / mimi.decoder.model for stage i, see SEANetEncoder/SEANetDecoder
_ENC_RES_IDX = lambda i: 1 + 3 * i # noqa: E731
_ENC_SCALE_IDX = lambda i: 3 + 3 * i # noqa: E731
_DEC_SCALE_IDX = lambda i: 2 + 3 * i # noqa: E731
_DEC_RES_IDX = lambda i: 3 + 3 * i # noqa: E731
_N_SEANET_STAGES = 3
_SAMPLE_RATE = 24000
def _tensor_shapes(dir_model: Path) -> dict[str, tuple[int, ...]]:
part_names = ModelBase.get_model_part_names(dir_model, "model", ".safetensors")
if len(part_names) != 1:
return {}
with gguf.utility.SafetensorsLocal(dir_model / part_names[0]) as part:
return {name: tuple(part[name].shape) for name in part.keys()}
@ModelBase.register_hparams_loader(lambda dir_model: "flow_lm.bos_emb" in _tensor_shapes(dir_model))
def _load_hparams(dir_model: Path) -> dict[str, Any]:
logger.info("gguf: detected pocket-tts checkpoint, deriving hparams from tensor shapes")
shapes = _tensor_shapes(dir_model)
n_vocab, n_embd = shapes["flow_lm.conditioner.embed.weight"]
n_layer = sum(1 for name in shapes if re.fullmatch(r"flow_lm\.transformer\.layers\.\d+\.norm1\.weight", name))
n_layer_a = sum(1 for name in shapes if re.fullmatch(r"mimi\.encoder_transformer\.transformer\.layers\.\d+\.norm1\.weight", name))
n_embd_a = shapes["mimi.encoder_transformer.transformer.layers.0.norm1.weight"][0]
return {
"architectures": ["PocketTTSModel"],
"model_type": "pockettts",
"num_hidden_layers": n_layer,
"hidden_size": n_embd,
"intermediate_size": shapes["flow_lm.transformer.layers.0.linear1.weight"][0],
# the transformer is fully causal with no context limit, this only bounds the KV cache
"max_position_embeddings": 4096,
# not in the checkpoint, but every released variant uses head_dim 64
"num_attention_heads": n_embd // 64,
# extra rows for the learned input vectors, see _embd_table()
"vocab_size": n_vocab + (2 if "flow_lm.bos_before_voice" in shapes else 1),
"rope_theta": 10000.0,
"layer_norm_eps": 1e-5,
"audio_config": {
"num_hidden_layers": n_layer_a,
"hidden_size": n_embd_a,
"intermediate_size": shapes["mimi.encoder_transformer.transformer.layers.0.linear1.weight"][0],
"num_attention_heads": n_embd_a // 64,
},
}
@ModelBase.register("PocketTTSModel")
class PocketTTSModel(TextModel):
model_arch = gguf.MODEL_ARCH.POCKETTTS
_LAYER_TENSOR_MAP = {
"norm1": gguf.MODEL_TENSOR.ATTN_NORM,
"norm2": gguf.MODEL_TENSOR.FFN_NORM,
"self_attn.out_proj": gguf.MODEL_TENSOR.ATTN_OUT,
"linear1": gguf.MODEL_TENSOR.FFN_UP,
"linear2": gguf.MODEL_TENSOR.FFN_DOWN,
}
def set_vocab(self):
# this is a unigram sentencepiece model, llama.cpp's SPM tokenizer cannot do
# unigram segmentation, so use the UGM tokenizer instead
from sentencepiece import sentencepiece_model_pb2 as model
proto = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
proto.ParseFromString(open(self.dir_model / "tokenizer.model", "rb").read())
assert proto.trainer_spec.model_type == 1, "expected a unigram tokenizer"
tokens, scores, toktypes = self._create_vocab_sentencepiece()
# the last rows of the embedding table are not sentencepiece pieces
extra = self._extra_tokens()
for i, name in enumerate(extra):
tokens[len(tokens) - len(extra) + i] = name.encode("utf-8")
toktypes[len(tokens) - len(extra) + i] = SentencePieceTokenTypes.CONTROL
scores[len(tokens) - len(extra) + i] = -1000.0
self.gguf_writer.add_tokenizer_model("t5")
self.gguf_writer.add_tokenizer_pre("default")
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_scores(scores)
self.gguf_writer.add_token_types(toktypes)
self.gguf_writer.add_add_space_prefix(proto.normalizer_spec.add_dummy_prefix)
self.gguf_writer.add_remove_extra_whitespaces(proto.normalizer_spec.remove_extra_whitespaces)
if proto.normalizer_spec.precompiled_charsmap:
self.gguf_writer.add_precompiled_charsmap(proto.normalizer_spec.precompiled_charsmap)
self.gguf_writer.add_add_bos_token(False)
self.gguf_writer.add_add_eos_token(False)
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
if not name.startswith("flow_lm."):
return # mimi and the flow net go to the mmproj
if name == "flow_lm.conditioner.embed.weight":
yield (self.format_tensor_name(gguf.MODEL_TENSOR.TOKEN_EMBD), self._embd_table(data_torch))
return
if name.startswith("flow_lm.out_norm."):
suffix = "." + name.rsplit(".", 1)[1]
yield (self.format_tensor_name(gguf.MODEL_TENSOR.OUTPUT_NORM, suffix=suffix), data_torch)
return
if name.startswith("flow_lm.transformer.layers."):
assert bid is not None
key_with_suffix = name.split(f"layers.{bid}.", 1)[1]
key, suffix = key_with_suffix.rsplit(".", 1)
if key == "self_attn.in_proj":
q, k, v = data_torch.chunk(3, dim=0)
yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_Q, bid), q)
yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_K, bid), k)
yield (self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_V, bid), v)
return
tensor = self._LAYER_TENSOR_MAP.get(key)
if tensor is not None:
yield (self.format_tensor_name(tensor, bid, suffix="." + suffix), data_torch)
return
return
def _extra_tokens(self) -> list[str]:
# the conditioner's padding row, then the learned vectors appended by _embd_table().
# bos_before_voice only exists when the pack sets insert_bos_before_voice
names = ["<|pad|>"]
if "flow_lm.bos_before_voice" in self.model_tensors:
names.append("<|bos_before_voice|>")
names.append("<|audio_bos|>")
return names
def _embd_table(self, embed: Tensor) -> Tensor:
rows = [embed]
if "flow_lm.bos_before_voice" in self.model_tensors:
rows.append(self.model_tensors["flow_lm.bos_before_voice"]().reshape(1, -1).to(embed.dtype))
# bos_emb is a latent, it only enters the backbone through input_linear
bos_emb = self.model_tensors["flow_lm.bos_emb"]()
input_linear = self.model_tensors["flow_lm.input_linear.weight"]()
audio_bos = torch.nn.functional.linear(bos_emb.float(), input_linear.float()).reshape(1, -1)
rows.append(audio_bos.to(embed.dtype))
return torch.cat(rows, dim=0)
@ModelBase.register("PocketTTSModel")
class PocketTTSMmprojModel(MmprojModel):
has_audio_encoder = True
has_vision_encoder = False
_MIMI_TFM_MAP = {
"norm1": (gguf.MODEL_TENSOR.A_ENC_INPUT_NORM, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_NORM),
"norm2": (gguf.MODEL_TENSOR.A_ENC_OUTPUT_NORM, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_NORM),
"self_attn.out_proj": (gguf.MODEL_TENSOR.A_ENC_OUTPUT, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_OUT),
"linear1": (gguf.MODEL_TENSOR.A_ENC_FFN_UP, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_UP),
"linear2": (gguf.MODEL_TENSOR.A_ENC_FFN_DOWN, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_DOWN),
"layer_scale_1.scale": (gguf.MODEL_TENSOR.A_ENC_ATTN_SCALE, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_SCALE),
"layer_scale_2.scale": (gguf.MODEL_TENSOR.A_ENC_FFN_SCALE_LS, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_FFN_SCALE),
}
_MIMI_TFM_QKV = (
(gguf.MODEL_TENSOR.A_ENC_ATTN_Q, gguf.MODEL_TENSOR.A_ENC_ATTN_K, gguf.MODEL_TENSOR.A_ENC_ATTN_V),
(gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_Q, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_K, gguf.MODEL_TENSOR.A_GEN_WAV_TFM_ATTN_V),
)
def set_gguf_parameters(self):
self.gguf_writer.add_file_type(self.ftype)
assert self.hparams_audio is not None
# voice-prompt encoder: mimi encoder + speaker_proj
self.gguf_writer.add_clip_has_audio_encoder(True)
# note: the 24kHz sample rate is hardcoded on the clip.cpp side, like the other audio models
self.gguf_writer.add_clip_audio_projector_type(gguf.VisionProjectorType.POCKETTTS_SPKENC)
self.gguf_writer.add_audio_projection_dim(self.n_embd_text)
self.gguf_writer.add_audio_block_count(self.hparams_audio["num_hidden_layers"])
self.gguf_writer.add_audio_embedding_length(self.hparams_audio["hidden_size"])
self.gguf_writer.add_audio_feed_forward_length(self.hparams_audio["intermediate_size"])
self.gguf_writer.add_audio_head_count(self.hparams_audio["num_attention_heads"])
self.gguf_writer.add_audio_attention_layernorm_eps(1e-5)
# mimi convolves the waveform directly, it is passed around as a 1-row "mel"
self.gguf_writer.add_audio_num_mel_bins(1)
# generation: flow-matching decoder + mimi decoder
# the SEANet and flow net hparams are constant across the family, clip.cpp holds them
self.gguf_writer.add_clip_has_gen_audio_encoder(True)
self.gguf_writer.add_clip_gen_audio_projector_type(gguf.VisionProjectorType.POCKETTTS_GEN)
self.gguf_writer.add_gen_audio_projection_dim(self.n_embd_text)
self.gguf_writer.add_gen_audio_embedding_length(self.hparams_audio["hidden_size"])
self.gguf_writer.add_gen_audio_feed_forward_length(self.hparams_audio["intermediate_size"])
self.gguf_writer.add_gen_audio_block_count(self.hparams_audio["num_hidden_layers"])
self.gguf_writer.add_gen_audio_head_count(self.hparams_audio["num_attention_heads"])
self.gguf_writer.add_gen_audio_attention_layernorm_eps(1e-5)
self.gguf_writer.add_gen_audio_model_variant(self.dir_model.name)
def tensor_force_quant(self, name, new_name, bid, n_dims):
del name, bid, n_dims
# conv1d/conv1d_dw kernels must be F16, ggml_conv_1d(_dw) has no BF16 path
if ".seanet." in new_name or new_name in ("a.downsample.conv.weight", "a.gen.wav.upsample.weight"):
return gguf.GGMLQuantizationType.F16
return False
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
del bid # the block index of the mimi transformers is parsed here, not by the base class
T = gguf.MODEL_TENSOR
if name in ("flow_lm.bos_emb", "flow_lm.bos_before_voice", "flow_lm.conditioner.embed.weight"):
return # folded into the backbone embedding table
if name.startswith("flow_lm.transformer.") or name.startswith("flow_lm.out_norm."):
return # backbone
if name == "flow_lm.speaker_proj_weight":
yield (self.format_tensor_name(T.A_ENC_SPEAKER_PROJ), data_torch)
return
if name == "flow_lm.input_linear.weight":
yield (self.format_tensor_name(T.A_GEN_INPUT_LINEAR), data_torch)
return
if name == "flow_lm.emb_mean":
yield (self.format_tensor_name(T.A_GEN_EMB_MEAN, suffix=""), data_torch)
return
if name == "flow_lm.emb_std":
yield (self.format_tensor_name(T.A_GEN_EMB_STD, suffix=""), data_torch)
return
if name.startswith("flow_lm.out_eos."):
suffix = "." + name.rsplit(".", 1)[1]
yield (self.format_tensor_name(T.A_GEN_OUT_EOS, suffix=suffix), data_torch)
return
if name.startswith("flow_lm.flow_net."):
yield from self._flow_net_tensor(name, data_torch)
return
if name == "mimi.downsample.conv.conv.weight":
yield (self.format_tensor_name(T.A_ENC_DOWNSAMPLE_CONV), data_torch)
return
if name == "mimi.upsample.convtr.convtr.weight":
yield (self.format_tensor_name(T.A_GEN_WAV_UPSAMPLE), data_torch)
return
if name == "mimi.quantizer.output_proj.weight":
yield (self.format_tensor_name(T.A_GEN_WAV_QUANT_OUT), data_torch.squeeze(-1))
return
if "_transformer.transformer.layers." in name:
yield from self._mimi_tfm_tensor(name, data_torch)
return
if name.startswith("mimi.encoder.model.") or name.startswith("mimi.decoder.model."):
yield from self._seanet_tensor(name, data_torch)
return
return
def _flow_net_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
T = gguf.MODEL_TENSOR
key = name.split("flow_lm.flow_net.", 1)[1]
suffix = "." + key.rsplit(".", 1)[1]
simple = {
"input_proj": T.A_GEN_FLOW_INPUT_PROJ,
"cond_embed": T.A_GEN_FLOW_COND_EMBD,
"final_layer.linear": T.A_GEN_FLOW_FINAL_PROJ,
"final_layer.adaLN_modulation.1": T.A_GEN_FLOW_FINAL_ADA,
}
tensor = simple.get(key.rsplit(".", 1)[0])
if tensor is not None:
yield (self.format_tensor_name(tensor, suffix=suffix), data_torch)
return
if key.startswith("time_embed."):
bid = int(key.split(".")[1])
rest = key.split(f"time_embed.{bid}.", 1)[1]
time_map = {
"freqs": (T.A_GEN_FLOW_TIME_FREQS, ""),
"mlp.0": (T.A_GEN_FLOW_TIME_UP, suffix),
"mlp.2": (T.A_GEN_FLOW_TIME_DOWN, suffix),
"mlp.3.alpha": (T.A_GEN_FLOW_TIME_NORM, ""),
}
entry = time_map.get(rest) or time_map.get(rest.rsplit(".", 1)[0])
if entry is not None:
yield (self.format_tensor_name(entry[0], bid, suffix=entry[1]), data_torch)
return
if key.startswith("res_blocks."):
bid = int(key.split(".")[1])
rest = key.split(f"res_blocks.{bid}.", 1)[1].rsplit(".", 1)[0]
blk_map = {
"in_ln": T.A_GEN_FLOW_BLK_NORM,
"mlp.0": T.A_GEN_FLOW_BLK_UP,
"mlp.2": T.A_GEN_FLOW_BLK_DOWN,
"adaLN_modulation.1": T.A_GEN_FLOW_BLK_ADA,
}
tensor = blk_map.get(rest)
if tensor is not None:
yield (self.format_tensor_name(tensor, bid, suffix=suffix), data_torch)
return
def _mimi_tfm_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
is_decoder = name.startswith("mimi.decoder_transformer.")
bid = int(name.split("_transformer.transformer.layers.", 1)[1].split(".")[0])
key_with_suffix = name.split(f".layers.{bid}.", 1)[1]
if key_with_suffix == "self_attn.in_proj.weight":
q, k, v = data_torch.chunk(3, dim=0)
names = self._MIMI_TFM_QKV[1 if is_decoder else 0]
for tensor, part in zip(names, (q, k, v)):
yield (self.format_tensor_name(tensor, bid), part)
return
key, suffix = key_with_suffix.rsplit(".", 1)
entry = self._MIMI_TFM_MAP.get(key) or self._MIMI_TFM_MAP.get(key_with_suffix)
if entry is None:
return
tensor = entry[1 if is_decoder else 0]
suffix = ".weight" if key_with_suffix.endswith(".scale") else "." + suffix
yield (self.format_tensor_name(tensor, bid, suffix=suffix), data_torch)
def _seanet_tensor(self, name: str, data_torch: Tensor) -> Iterable[tuple[str, Tensor]]:
T = gguf.MODEL_TENSOR
is_decoder = name.startswith("mimi.decoder.")
idx = int(name.split(".model.", 1)[1].split(".")[0])
suffix = "." + name.rsplit(".", 1)[1]
conv_in, conv_out, res1, res2, scale = (
(T.A_GEN_WAV_SEANET_CONV_IN, T.A_GEN_WAV_SEANET_CONV_OUT, T.A_GEN_WAV_SEANET_RES_CONV1,
T.A_GEN_WAV_SEANET_RES_CONV2, T.A_GEN_WAV_SEANET_SCALE_CONV)
if is_decoder else
(T.A_ENC_SEANET_CONV_IN, T.A_ENC_SEANET_CONV_OUT, T.A_ENC_SEANET_RES_CONV1,
T.A_ENC_SEANET_RES_CONV2, T.A_ENC_SEANET_SCALE_CONV)
)
if idx == 0:
yield (self.format_tensor_name(conv_in, suffix=suffix), data_torch)
return
if idx == 3 * _N_SEANET_STAGES + 2:
yield (self.format_tensor_name(conv_out, suffix=suffix), data_torch)
return
for stage in range(_N_SEANET_STAGES):
res_idx = _DEC_RES_IDX(stage) if is_decoder else _ENC_RES_IDX(stage)
scale_idx = _DEC_SCALE_IDX(stage) if is_decoder else _ENC_SCALE_IDX(stage)
if idx == scale_idx:
yield (self.format_tensor_name(scale, stage, suffix=suffix), data_torch)
return
if idx == res_idx:
# block.1 is the dilated conv, block.3 the pointwise one (0 and 2 are ELU)
inner = int(name.split(".block.", 1)[1].split(".")[0])
tensor = res1 if inner == 1 else res2
yield (self.format_tensor_name(tensor, stage, suffix=suffix), data_torch)
return
+1 -10
View File
@@ -647,13 +647,10 @@ class DFlashModel(Qwen3Model):
# own tokenizer logic, not the Qwen default).
from . import get_model_class
with open(self.target_model_dir / "config.json", "r", encoding="utf-8") as f:
target_hparams = json.load(f)
target_arch = target_hparams["architectures"][0]
target_arch = json.load(f)["architectures"][0]
target_cls = get_model_class(target_arch)
if target_cls is not type(self):
if target_arch == "NemotronHForCausalLM":
setattr(self, "is_moe", "num_experts_per_tok" in target_hparams)
target_cls.set_vocab(self) # ty: ignore[unresolved-attribute]
else:
super().set_vocab()
@@ -691,12 +688,6 @@ class DFlashModel(Qwen3Model):
name = "model." + name
return super().filter_tensors((name, gen))
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
if name == "model.embed_tokens.weight" and not self.hparams.get("has_embed_tokens", True):
return
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("Qwen3DSparkModel")
class DSparkModel(DFlashModel):
-42
View File
@@ -449,8 +449,6 @@ Or
use 1 SYCL GPUs: [0] with Max compute units:512
```
User can use the device management in [docs/multi-gpu.md](https://github.com/ggml-org/llama.cpp/blob/master/docs/multi-gpu.md), like parameter `--device SYCL0,SYCL1` to assign one or more devices.
## Windows
### Install GPU driver
@@ -765,7 +763,6 @@ Or
use 1 SYCL GPUs: [0] with Max compute units:512
```
User can use the device management in [docs/multi-gpu.md](https://github.com/ggml-org/llama.cpp/blob/master/docs/multi-gpu.md), like parameter `--device SYCL0,SYCL1` to assign one or more devices.
## Environment Variable
@@ -898,45 +895,6 @@ Pass these via `CXXFLAGS` or add a one-off `#define` to enable a flag on the spo
set UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
```
- When I set `SYCL_CACHE_PERSISTENT=1` in running time, I meet crash.
`SYCL_CACHE_PERSISTENT=1` is not recommended by llama.cpp SYCL backend.
When cache is enabled, SYCL runtime will try to cache and reuse JIT-compiled binaries.
We find some AI will tell user this cmd to speed up SYCL backend. It only speeds up the startup to skip the JIT process, instead of running speed.
It will bring negative impact when the SYCL binary file is changed frequently in your running environment. The new & old codes mix will lead to crash.
Compare to the benefit, it has brought more failed cases.
If you are not familiar with the SYCL compiler principle of JIT and AOT, please don't use it.
To restore, you need to remove the local cache: `~/.cache/libsycl_cache/` and execute `unset SYCL_CACHE_PERSISTENT` in running time.
- How to use iGPU and dGPU in same time?
1. Detect the devices in your running time.
```
source /opt/intel/oneapi/setvars.sh
./build/bin/llama-server --list-devices
or
./build/bin/llama-cli --list-devices
./build/bin/llama-bench --list-devices
./build/bin/llama-completion --list-devices
Available devices:
SYCL0: Intel(R) Arc(TM) A770 Graphics (15473 MiB, 15473 MiB free)
SYCL1: Intel(R) UHD Graphics 770 (59675 MiB, 44986 MiB free)
```
The dGPU will be in the head of this list and iGPU will be the end.
If not all GPUs are listed, please check the env var: ONEAPI_DEVICE_SELECTOR and unset it.
2. Set the iGPU and dGPU
Set the iGPU and dGPU by `./build/bin/llama-server --device SYCL0,SYCL1,SYCLxxx`.
### **GitHub contribution**:
Please add the `[SYCL]` prefix/tag in issues/PRs titles to help the SYCL contributors to check/address them without delay.
+6 -6
View File
@@ -15,7 +15,7 @@ Legend:
| Operation | BLAS | CANN | CPU | CUDA | ET | MTL | OpenCL | SYCL | Vulkan | WebGPU | ZenDNN | zDNN |
|-----------|------|------|------|------|------|------|------|------|------|------|------|------|
| ABS | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| ACC | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | | ✅ | ❌ | ❌ | ❌ |
| ACC | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ |
| ADD | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| ADD1 | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
| ADD_ID | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -41,9 +41,9 @@ Legend:
| DIAG | ❌ | ❌ | ✅ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| DIAG_MASK_INF | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
| DIV | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| DSV4_HC_COMB | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_POST | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_PRE | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_COMB | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_POST | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_PRE | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| DUP | ❌ | ✅ | ✅ | 🟡 | ❌ | 🟡 | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
| ELU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| EXP | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -59,7 +59,7 @@ Legend:
| GELU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| GELU_ERF | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| GELU_QUICK | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| GET_ROWS | ❌ | 🟡 | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | ❌ | ❌ |
| GET_ROWS | ❌ | 🟡 | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | | ✅ | 🟡 | ❌ | ❌ |
| GET_ROWS_BACK | ❌ | ❌ | 🟡 | 🟡 | ❌ | ❌ | ❌ | ❌ | 🟡 | ❌ | ❌ | ❌ |
| GROUP_NORM | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ |
| HARDSIGMOID | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
@@ -68,7 +68,7 @@ Legend:
| IM2COL_3D | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| L2_NORM | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | 🟡 | ❌ | ❌ |
| LEAKY_RELU | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| LIGHTNING_INDEXER | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| LIGHTNING_INDEXER | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| LOG | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| MEAN | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ |
| MUL | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
+671 -22870
View File
File diff suppressed because it is too large Load Diff
+1 -16
View File
@@ -4,7 +4,7 @@
The INI preset feature, introduced in [PR#17859](https://github.com/ggml-org/llama.cpp/pull/17859), allows users to create reusable and shareable parameter configurations for llama.cpp.
## Using Presets with the Server
### Using Presets with the Server
When running multiple models on the server (router mode), INI preset files can be used to configure model-specific parameters. Please refer to the [server documentation](../tools/server/README.md) for more details.
@@ -93,18 +93,3 @@ llama-server -hf user/repo:gpt-oss-120b-hf
```
Please make sure to provide the correct `hf-repo` for each child preset. Otherwise, you may get error: `The specified tag is not a valid quantization scheme.`
## System-level config
The system-level config, added in PR [#26118](https://github.com/ggml-org/llama.cpp/pull/26118), allows sharing the same set of options among multiple tools and examples. Unlike the sections above, it is not limited to the server.
These files are loaded on startup if present. A later file overrides an earlier one:
1. System-wide: `/etc/llama.cpp/config.ini` (or `%PROGRAMDATA%\llama.cpp\config.ini` on Windows)
2. User-level: `$XDG_CONFIG_HOME/llama.cpp/config.ini`, `~/.config/llama.cpp/config.ini` by default (or `%APPDATA%\llama.cpp\config.ini` on Windows)
The config file is applied first, then its options are overridden by ENV variables, CLI arguments and model presets (in router mode).
Note:
- Only the `[*]` and default sections are used; options written before any section header belong to "default. Named sections are ignored
- Tool-specific options can be specified, but will be ignored (with a warning) if the example doesn't support it<br/>Example: if you specify `port = 1234`, only `llama-server` will use it, other examples will ignore it
- `model` or `hf-repo` are not recommended to be configured system-level, because it may introduce conflicts<br/>Example: a `hf-repo` in the config file still takes effect when you pass `-m` on the command line, so you may load a different model than expected
-49
View File
@@ -1,49 +0,0 @@
# Release process
llama.cpp uses [semantic versioning](https://semver.org) (`MAJOR.MINOR.PATCH`).
## Version bump guidelines
| Change type | Version component |
|---|---|
| Breaking change to the public C API (`include/llama.h`) | `MAJOR` |
| Backward-compatible features, model support, or API addition | `MINOR` |
| Bug fix with no API change | `PATCH` |
The version is set in the three variables at the top of the root `CMakeLists.txt`:
```cmake
set(LLAMA_VERSION_MAJOR 0)
set(LLAMA_VERSION_MINOR 1)
set(LLAMA_VERSION_PATCH 0)
```
_A version bump should be included in the PR that introduces the change, or in a
dedicated bump commit merged before the release is cut._
_TODO: add PR labels (`semver: patch`, `semver: minor`, `semver: major`) to help
identify which PRs require a version bump before cutting a release._
## Making a release
Releases are created by running the [make-release](.github/workflows/make-release.yml)
which is a manual workflow.
The workflow creates an annotated git tag (e.g. `v0.1.0`) and pushes it to the
remote. No GitHub Release object is created, the tag is the release artifact.
## Building a release
By default, `LLAMA_BUILD_IS_DEV=ON` which appends a `-dev` suffix to `LLAMA_VERSION`,
marking the build as a nightly/development build. Distributors building from a
release tag must pass `-DLLAMA_BUILD_IS_DEV=OFF` to produce a clean version string
(e.g. `0.1.0` instead of `0.1.0-dev`).
## How releases reach users
Currently releases are not published to github releases, only nightly/development
builds are available there. The way users can access releases are using the following
channels:
- **llama-install.sh** — downloads pre-built binaries built from the release tag.
- **Package managers** — consume the git tag directly.
- **Build from source** — users clone the repo and check out the tag.
-6
View File
@@ -202,12 +202,6 @@ Example Video:
If a draft model is combined with a draftless decoding the draftless decoding has higher precedence.
### Backend Sampling
Use `--backend-sampling` to run supported target-model samplers on the model backend. Draft-model sampling uses the backend by default and can be controlled with `--spec-draft-backend-sampling` and `--no-spec-draft-backend-sampling`.
Unsupported samplers and device layouts fall back to CPU sampling. Tensor split mode does not support backend sampling. A fixed seed produces repeatable random draws, but stochastic CPU and backend sampling can still select different tokens because floating-point operations can differ between implementations and devices. Use greedy sampling when exact output matching is required.
### General Speculative Parameters
```
-6
View File
@@ -3,11 +3,9 @@
#include "common.h"
#include "ngram-cache.h"
#include "sampling.h"
#include "speculative.h"
#include "log.h"
#include "llama.h"
#include <algorithm>
#include <clocale>
#include <cstdint>
#include <cstdio>
@@ -29,10 +27,6 @@ int main(int argc, char ** argv){
// max. number of additional tokens to draft if match is found
const int n_draft = params.speculative.draft.n_max;
const auto output_limits = common_speculative_get_output_limits(params.n_batch, params.n_parallel, n_draft);
params.n_outputs_max = output_limits.total;
params.n_outputs_max_per_seq = output_limits.per_seq;
// init llama.cpp
llama_backend_init();
llama_numa_init(params.numa);
+1 -1
View File
@@ -1,6 +1,6 @@
--extra-index-url https://download.pytorch.org/whl/cpu
torch
torchvision; platform_machine != "s390x"
torchvision
transformers
huggingface-hub
accelerate
@@ -47,7 +47,6 @@ CMD_ARGS+=("../../convert_hf_to_gguf.py" "--verbose")
CMD_ARGS+=("${MODEL_PATH}")
CMD_ARGS+=("--outfile" "${CONVERTED_MODEL}")
CMD_ARGS+=("--outtype" "${TYPE}")
CMD_ARGS+=("--model-name" "${MODEL_NAME}")
[[ -n "$METADATA_OVERRIDE" ]] && CMD_ARGS+=("--metadata" "${METADATA_OVERRIDE}")
[[ -n "$MMPROJ" ]] && CMD_ARGS+=("${MMPROJ}")
@@ -2,15 +2,12 @@
import argparse
import os
import sys
import importlib
import torch
import numpy as np
from transformers import AutoTokenizer, AutoConfig, AutoModelForCausalLM
sys.path.insert(0, os.path.join(os.path.dirname(__file__), '..'))
from utils.common import save_output_data
from pathlib import Path
unreleased_model_name = os.getenv('UNRELEASED_MODEL_NAME')
@@ -57,7 +54,6 @@ print(f"Model name: {model_name}")
prompt = "Hello world today"
input_ids = tokenizer(prompt, return_tensors="pt").input_ids # ty: ignore[call-non-callable]
token_ids = input_ids[0].cpu().tolist()
print(f"Input tokens: {input_ids}")
print(f"Input text: {repr(prompt)}")
print(f"Tokenized: {tokenizer.convert_ids_to_tokens(input_ids[0])}") # ty: ignore[unresolved-attribute]
@@ -78,8 +74,21 @@ with torch.no_grad():
print(f"Hidden dimension: {token_embeddings.shape[-1]}")
print(f"Number of tokens: {token_embeddings.shape[0]}")
# Save raw token embeddings
data_dir = Path("data")
data_dir.mkdir(exist_ok=True)
bin_filename = data_dir / f"pytorch-{model_name}-embeddings.bin"
txt_filename = data_dir / f"pytorch-{model_name}-embeddings.txt"
# Save all token embeddings as binary
print(token_embeddings)
save_output_data(token_embeddings, token_ids, prompt, model_name, type_suffix="-embeddings")
token_embeddings.astype(np.float32).tofile(bin_filename)
# Save as text for inspection
with open(txt_filename, "w") as f:
for i, embedding in enumerate(token_embeddings):
for j, val in enumerate(embedding):
f.write(f"{i} {j} {val:.6f}\n")
# Print embeddings per token in the requested format
print("\nToken embeddings:")
@@ -101,3 +110,5 @@ with torch.no_grad():
for i, token in enumerate(tokens):
print(f" Token {i}: {repr(token)}")
print(f"Saved bin logits to: {bin_filename}")
print(f"Saved txt logist to: {txt_filename}")
@@ -31,7 +31,6 @@ python ../../convert_hf_to_gguf.py --verbose \
${EMBEDDING_MODEL_PATH} \
--outfile ${CONVERTED_MODEL} \
--outtype ${TYPE} \
--model-name ${MODEL_NAME} \
${SENTENCE_TRANSFORMERS}
echo ""
+5 -42
View File
@@ -3,47 +3,10 @@
Demonstration of basic greedy speculative decoding
```bash
# spec-type draft-simple
./bin/llama-speculative-simple \
-hf ggml-org/Qwen3-8B-Base-GGUF:Q8_0 \
-hfd ggml-org/Qwen3-0.6B-Base-GGUF \
-p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-simple --spec-draft-n-max 7 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
# spec-type draft-mtp
./bin/llama-speculative-simple \
-hf ggml-org/Qwen3.6-27B-GGUF:Q8_0 \
-p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-mtp --spec-draft-n-max 3 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
# spec-type draft-mtp (with shared KV cache)
# note: this model needs a <s> token at the start to somewhat work without the chat template
./bin/llama-speculative-simple \
-hf ggml-org/Gemma-4-31B-it-GGUF:Q8_0 \
-p "<s>Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-mtp --spec-draft-n-max 3 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
# spec-type draft-eagle3
./bin/llama-speculative-simple \
-hf ggml-org/gpt-oss-20b-GGUF \
-p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-eagle3 --spec-draft-n-max 3 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
# spec-type draft-dflash
./bin/llama-speculative-simple \
-hf ggml-org/Qwen3-8B-GGUF \
-p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-dflash --spec-draft-n-max 7 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
# spec-type draft-dspark
./bin/llama-speculative-simple \
-hf ggml-org/Qwen3-8B-GGUF \
-p "Here is a quick sort implementation in C++. Just code, no comments:\n\n#include" \
--spec-type draft-dspark --spec-draft-n-max 7 -ngld 99 --color on \
-n 256 --temp 0 --top-k 1 --seed 42 -ngl 99 -lv 4
-m ../models/qwen2.5-32b-coder-instruct/ggml-model-q8_0.gguf \
-md ../models/qwen2.5-1.5b-coder-instruct/ggml-model-q4_0.gguf \
-f test.txt -c 0 -ngl 99 --color on \
--sampling-seq k --top-k 1 -fa on --temp 0.0 \
-ngld 99 --spec-draft-n-max 16 --spec-draft-n-draft-min 5 --draft-p-min 0.9
```
@@ -5,7 +5,6 @@
#include "log.h"
#include "llama.h"
#include <algorithm>
#include <clocale>
#include <cstdio>
#include <cstring>
@@ -30,11 +29,6 @@ int main(int argc, char ** argv) {
return 1;
}
const auto output_limits = common_speculative_get_output_limits(
params.n_batch, params.n_parallel, common_speculative_n_max(&params.speculative));
params.n_outputs_max = output_limits.total;
params.n_outputs_max_per_seq = output_limits.per_seq;
// init llama.cpp
llama_backend_init();
llama_numa_init(params.numa);
@@ -51,23 +45,45 @@ int main(int argc, char ** argv) {
const llama_vocab * vocab = llama_model_get_vocab(model_tgt);
// load the draft model (if any) - this also creates the MTP draft context when MTP speculation is enabled
common_speculative_init_result_ptr spec_init;
// load the draft model
llama_model_ptr model_dft;
llama_context_ptr ctx_dft;
// TODO: simplify this logic
{
common_params params_dft = common_base_params_to_speculative(params);
const auto & params_spec = params.speculative.draft;
spec_init = common_speculative_init_from_params(params_dft, model_tgt, ctx_tgt);
auto params_dft = params;
params_dft.devices = params_spec.devices;
params_dft.model = params_spec.mparams;
params_dft.n_gpu_layers = params_spec.n_gpu_layers;
if (params_spec.cpuparams.n_threads > 0) {
params_dft.cpuparams.n_threads = params.speculative.draft.cpuparams.n_threads;
params_dft.cpuparams_batch.n_threads = params.speculative.draft.cpuparams_batch.n_threads;
}
params_dft.tensor_buft_overrides = params.speculative.draft.tensor_buft_overrides;
auto mparams_dft = common_model_params_to_llama(params_dft);
model_dft.reset(llama_model_load_from_file(params_dft.model.path.c_str(), mparams_dft));
if (model_dft == nullptr) {
LOG_ERR("failed to load draft model, '%s'\n", params_dft.model.path.c_str());
return 1;
}
auto cparams = common_context_params_to_llama(params_dft);
ctx_dft.reset(llama_init_from_model(model_dft.get(), cparams));
params.speculative.draft.ctx_tgt = ctx_tgt;
params.speculative.draft.ctx_dft = spec_init->context();
params.speculative.draft.ctx_dft = ctx_dft.get();
}
llama_context * ctx_dft = params.speculative.draft.ctx_dft;
// check if the context supports partial sequence removal
const bool use_ckpt_tgt = common_context_can_seq_rm(ctx_tgt) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL;
const bool use_ckpt_dft = common_context_can_seq_rm(ctx_dft) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL;
const bool use_ckpt_tgt = (common_context_can_seq_rm(ctx_tgt) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL);
const bool use_ckpt_dft = (common_context_can_seq_rm(ctx_dft.get()) == COMMON_CONTEXT_SEQ_RM_TYPE_FULL);
if (use_ckpt_tgt) {
LOG_INF("speculative decoding will use checkpoints (context does not support partial sequence removal)\n");
@@ -113,30 +129,9 @@ int main(int argc, char ** argv) {
// target model sampling context
common_sampler_ptr smpl(common_sampler_init(model_tgt, params.sampling));
// init the speculator
const auto & params_spec = params.speculative;
struct common_speculative * spec = common_speculative_init(params.speculative, 1);
if (spec == nullptr) {
LOG_ERR("%s", "failed to initialize speculative decoding\n");
return 1;
}
// eval the prompt on the target and feed it to the speculative implementation(s)
{
llama_batch batch_prompt = llama_batch_init(inp.size(), 0, 1);
for (size_t i = 0; i < inp.size() - 1; ++i) {
common_batch_add(batch_prompt, inp[i], i, { seq_id }, false);
}
llama_decode(ctx_tgt, batch_prompt);
if (!common_speculative_process(spec, batch_prompt)) {
LOG_ERR("%s", "failed to process speculative prompt\n");
return 1;
}
}
// eval the prompt
llama_decode(ctx_tgt, llama_batch_get_one(inp.data(), inp.size() - 1));
llama_decode(ctx_dft.get(), llama_batch_get_one(inp.data(), inp.size() - 1));
// note: keep the last token separate!
llama_token id_last = inp.back();
@@ -147,12 +142,18 @@ int main(int argc, char ** argv) {
int n_past = inp.size() - 1;
// init the speculator
const auto & params_spec = params.speculative;
struct common_speculative * spec = common_speculative_init(params.speculative, 1);
common_speculative_begin(spec, seq_id, prompt_tgt);
llama_batch batch_tgt = llama_batch_init(llama_n_batch(ctx_tgt), 0, 1);
llama_tokens draft;
size_t n_draft = 0;
llama_tokens draft;
common_prompt_checkpoint ckpt;
const auto t_enc_end = ggml_time_us();
@@ -174,20 +175,13 @@ int main(int argc, char ** argv) {
llama_memory_seq_pos_max(llama_get_memory(ctx_tgt), seq_id));
if (use_ckpt_dft) {
ckpt.update_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
ckpt.update_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
}
// determine the max draft that fits the remaining context and generation budget
int n_draft_max = (int) llama_n_ctx(ctx_tgt) - n_past - 2;
if (params.n_predict >= 0) {
n_draft_max = std::min(n_draft_max, params.n_predict - n_predict - 1);
}
n_draft_max = std::max(n_draft_max, 0);
// generate a new draft
common_speculative_get_draft_params(spec, seq_id) = {
/* .drafting = */ true,
/* .n_max = */ n_draft_max,
/* .n_max = */ -1,
/* .n_past = */ n_past,
/* .id_last = */ id_last,
/* .prompt = */ &prompt_tgt,
@@ -195,6 +189,9 @@ int main(int argc, char ** argv) {
};
common_speculative_draft(spec);
// save the original draft size
n_draft = draft.size();
// save a checkpoint of the target context before evaluating the draft
// this allows us to restore the state if partial draft acceptance occurs
if (!draft.empty()) {
@@ -203,13 +200,10 @@ int main(int argc, char ** argv) {
}
}
// reset the draft context to the checkpoint before verification
if (ctx_dft) {
if (use_ckpt_dft) {
ckpt.load_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
}
{
ckpt.load_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, ckpt.pos_max + 1, -1);
llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, ckpt.pos_max + 1, -1);
}
} else {
// we have a previous (partial) draft to reuse from checkpoint restoration
@@ -233,10 +227,10 @@ int main(int argc, char ** argv) {
llama_decode(ctx_tgt, batch_tgt);
}
// feed the batch to the speculative implementation(s) - this drives the draft model, MTP, Eagle3, etc.
if (!common_speculative_process(spec, batch_tgt)) {
LOG_ERR("%s", "failed to process speculative batch\n");
break;
// evaluate the same batch with the draft model
{
// TODO: extend to support MTP, Eagle, etc. See server code for reference
llama_decode(ctx_dft.get(), batch_tgt);
}
// only save the sampler sampler state if we use checkpoints
@@ -245,9 +239,6 @@ int main(int argc, char ** argv) {
smpl_save.reset(common_sampler_clone(smpl.get()));
}
// save the size of the draft being verified
const size_t n_draft = draft.size();
// sample from the full target batch and return the accepted tokens based on the target sampler
//
// for each token to be accepted, the sampler would have to sample that same token
@@ -264,8 +255,8 @@ int main(int argc, char ** argv) {
// check for partial draft acceptance:
// if the context doesn't support partial sequence removal, restore the checkpoint
// and make the accepted tokens the new partial draft for the next iteration
if (use_ckpt_tgt && ids.size() - 1 < n_draft) {
LOG_DBG("partial acceptance: %zu < %zu, restoring checkpoint\n", ids.size() - 1, n_draft);
if (use_ckpt_tgt && ids.size() - 1 < draft.size()) {
LOG_DBG("partial acceptance: %zu < %zu, restoring checkpoint\n", ids.size() - 1, draft.size());
draft = std::move(ids);
@@ -275,10 +266,10 @@ int main(int argc, char ** argv) {
llama_memory_seq_rm(llama_get_memory(ctx_tgt), seq_id, ckpt.pos_max + 1, -1);
}
if (ctx_dft) {
ckpt.load_dft(ctx_dft, seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
{
ckpt.load_dft(ctx_dft.get(), seq_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, ckpt.pos_max + 1, -1);
llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, ckpt.pos_max + 1, -1);
}
prompt_tgt.resize(ckpt.n_tokens);
@@ -329,11 +320,8 @@ int main(int argc, char ** argv) {
{
LOG_DBG("clear kv cache from any extra tokens, n_past = %d\n", n_past);
llama_memory_seq_rm(llama_get_memory(ctx_tgt), seq_id, n_past, -1);
if (ctx_dft) {
llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, n_past, -1);
}
llama_memory_seq_rm(llama_get_memory(ctx_tgt), seq_id, n_past, -1);
llama_memory_seq_rm(llama_get_memory(ctx_dft.get()), seq_id, n_past, -1);
}
if ((params.n_predict >= 0 && n_predict > params.n_predict) || has_eos) {
@@ -359,7 +347,6 @@ int main(int argc, char ** argv) {
LOG_INF("\n");
LOG_INF("draft:\n\n");
common_speculative_print_stats(spec);
LOG_INF("\n");
LOG_INF("target:\n\n");
-8
View File
@@ -1,7 +1,6 @@
#include "arg.h"
#include "common.h"
#include "sampling.h"
#include "speculative.h"
#include "log.h"
#include "llama.h"
@@ -58,11 +57,6 @@ int main(int argc, char ** argv) {
// max number of parallel drafting sequences (i.e. tree branches)
const int n_seq_dft = params.n_parallel;
const auto output_limits = common_speculative_get_output_limits(
params.n_batch, params.n_parallel, params.speculative.draft.n_max);
params.n_outputs_max = output_limits.total;
params.n_outputs_max_per_seq = output_limits.per_seq;
// probability threshold for splitting a draft branch (only for n_seq_dft > 1)
const float p_draft_split = params.speculative.draft.p_split;
@@ -89,8 +83,6 @@ int main(int argc, char ** argv) {
params.devices = params.speculative.draft.devices;
params.model = params.speculative.draft.mparams;
params.n_gpu_layers = params.speculative.draft.n_gpu_layers;
params.n_outputs_max = params.n_parallel;
params.n_outputs_max_per_seq = 1;
if (params.speculative.draft.cpuparams.n_threads > 0) {
params.cpuparams.n_threads = params.speculative.draft.cpuparams.n_threads;
}
+7 -14
View File
@@ -12,7 +12,6 @@ This script processes files with specified options.
Options:
-h, --help Display this help message and exit.
-d, --device <value> Set SYCL devices (default: SYCL0).
-c, --context <value> Set context length. Bigger need more memory.
-p, --promote <value> Prompt to start generation with.
-m, --model <value> Full model file path.
@@ -42,16 +41,10 @@ MODEL_FILE=../models/Qwen3.5-4B-Q4_0.gguf
NGL=99
CONTEXT=4096
GGML_SYCL_DEVICE=-1
SYCL_DEVICES="SYCL0"
SPLIT_MODE=layer
LOG_VERBOSE=3
while [[ $# -gt 0 ]]; do
case "$1" in
-d|--device)
SYCL_DEVICES="$2"
shift
shift
;;
-c|--context)
CONTEXT=$2
# Shift twice to consume both the option flag and its value
@@ -102,6 +95,8 @@ while [[ $# -gt 0 ]]; do
esac
done
source /opt/intel/oneapi/setvars.sh
#export GGML_SYCL_DEBUG=1
@@ -112,19 +107,17 @@ source /opt/intel/oneapi/setvars.sh
export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
echo "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=${UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS}"
echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
if [ $GGML_SYCL_DEVICE -ne -1 ]; then
echo "Use $GGML_SYCL_DEVICE as main GPU"
#use signle GPU only
GPUS_SETTING="-mg $GGML_SYCL_DEVICE -sm ${SPLIT_MODE}"
echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
else
echo "Use Intel GPUs: ${SYCL_DEVICES}"
echo "Use all Intel GPUs, including iGPU & dGPU"
GPUS_SETTING="-sm ${SPLIT_MODE}"
fi
fi
echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap --host 0.0.0.0 --port 8000"
ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap --host 0.0.0.0 --port 8000
echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap --host 0.0.0.0 --port 8000"
ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap --host 0.0.0.0 --port 8000
+4 -12
View File
@@ -12,7 +12,6 @@ This script processes files with specified options.
Options:
-h, --help Display this help message and exit.
-d, --device <value> Set SYCL devices (default: SYCL0).
-c, --context <value> Set context length. Bigger need more memory.
-p, --promote <value> Prompt to start generation with.
-m, --model <value> Full model file path.
@@ -43,16 +42,10 @@ MODEL_FILE=../models/llama-2-7b.Q4_0.gguf
NGL=99
CONTEXT=4096
GGML_SYCL_DEVICE=-1
SYCL_DEVICES="SYCL0"
SPLIT_MODE=layer
LOG_VERBOSE=3
while [[ $# -gt 0 ]]; do
case "$1" in
-d|--device)
SYCL_DEVICES="$2"
shift
shift
;;
-c|--context)
CONTEXT=$2
# Shift twice to consume both the option flag and its value
@@ -122,17 +115,16 @@ source /opt/intel/oneapi/setvars.sh
export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
echo "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=${UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS}"
echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
if [ $GGML_SYCL_DEVICE -ne -1 ]; then
echo "Use $GGML_SYCL_DEVICE as main GPU"
#use signle GPU only
GPUS_SETTING="-mg $GGML_SYCL_DEVICE -sm ${SPLIT_MODE}"
echo "ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR}"
else
echo "Use Intel GPUs: ${SYCL_DEVICES}"
echo "Use all Intel GPUs, including iGPU & dGPU"
GPUS_SETTING="-sm ${SPLIT_MODE}"
fi
echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap "
ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --device ${SYCL_DEVICES} --mmap
echo "run cmd: ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap "
ZES_ENABLE_SYSMAN=1 ${BIN_FILE} -m ${MODEL_FILE} -no-cnv -p "${INPUT_PROMPT}" -n 200 -e -ngl ${NGL} -s ${SEED} -c ${CONTEXT} ${GPUS_SETTING} -lv ${LOG_VERBOSE} --mmap
+5 -23
View File
@@ -13,7 +13,6 @@ set "MODEL_FILE=..\models\Qwen3.5-4B-Q4_0.gguf"
set "NGL=99"
set "CONTEXT=4096"
set "GGML_SYCL_DEVICE=-1"
set "SYCL_DEVICES=SYCL0"
set "SPLIT_MODE=layer"
set "LOG_VERBOSE=3"
@@ -37,21 +36,6 @@ if /I "%~1"=="--context" (
goto parse_args
)
if /I "%~1"=="-d" (
if "%~2"=="" goto missing_value
set "SYCL_DEVICES=%~2"
shift
shift
goto parse_args
)
if /I "%~1"=="--device" (
if "%~2"=="" goto missing_value
set "SYCL_DEVICES=%~2"
shift
shift
goto parse_args
)
if /I "%~1"=="-m" (
if "%~2"=="" goto missing_value
set "MODEL_FILE=%~2"
@@ -146,7 +130,6 @@ echo This script processes files with specified options.
echo.
echo Options:
echo -h, --help Display this help message and exit.
echo -d, --device ^<value^> Set SYCL devices (default: SYCL0).
echo -c, --context ^<value^> Set context length. Bigger need more memory.
echo -m, --model ^<value^> Full model file path.
echo -mg,--main-gpu ^<value^> Set main GPU ID (0 - n) for single GPU mode.
@@ -177,20 +160,19 @@ REM Support malloc device memory more than 4GB.
set "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1"
echo UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=%UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS%
echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
if not "%GGML_SYCL_DEVICE%"=="-1" (
echo Use %GGML_SYCL_DEVICE% as main GPU
REM Use single GPU only.
set "GPUS_SETTING=-mg %GGML_SYCL_DEVICE% -sm %SPLIT_MODE%"
) else (
echo Use Intel GPUs: %SYCL_DEVICES%
echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
) else (
echo Use all Intel GPUs, including iGPU ^& dGPU
set "GPUS_SETTING=-sm %SPLIT_MODE%"
)
echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device %SYCL_DEVICES% --mmap --host 0.0.0.0 --port 8000
echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap --host 0.0.0.0 --port 8000
set "ZES_ENABLE_SYSMAN=1"
%BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device "%SYCL_DEVICES%" --mmap --host 0.0.0.0 --port 8000
%BIN_FILE% -m "%MODEL_FILE%" -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap --host 0.0.0.0 --port 8000
endlocal
+5 -24
View File
@@ -19,7 +19,6 @@ set "MODEL_FILE=..\models\llama-2-7b.Q4_0.gguf"
set "NGL=99"
set "CONTEXT=4096"
set "GGML_SYCL_DEVICE=-1"
set "SYCL_DEVICES=SYCL0"
set "SPLIT_MODE=layer"
set "LOG_VERBOSE=3"
@@ -43,21 +42,6 @@ if /I "%~1"=="--context" (
goto parse_args
)
if /I "%~1"=="-d" (
if "%~2"=="" goto missing_value
set "SYCL_DEVICES=%~2"
shift
shift
goto parse_args
)
if /I "%~1"=="--device" (
if "%~2"=="" goto missing_value
set "SYCL_DEVICES=%~2"
shift
shift
goto parse_args
)
if /I "%~1"=="-p" (
if "%~2"=="" goto missing_value
set "INPUT_PROMPT=%~2"
@@ -167,7 +151,6 @@ echo This script processes files with specified options.
echo.
echo Options:
echo -h, --help Display this help message and exit.
echo -d, --device ^<value^> Set SYCL devices (default: SYCL0).
echo -c, --context ^<value^> Set context length. Bigger need more memory.
echo -p, --promote ^<value^> Prompt to start generation with.
echo -m, --model ^<value^> Full model file path.
@@ -199,21 +182,19 @@ REM Support malloc device memory more than 4GB.
set "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1"
echo UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=%UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS%
echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
if not "%GGML_SYCL_DEVICE%"=="-1" (
echo Use %GGML_SYCL_DEVICE% as main GPU
REM Use single GPU only.
set "GPUS_SETTING=-mg %GGML_SYCL_DEVICE% -sm %SPLIT_MODE%"
)
else (
echo Use Intel GPUs: %SYCL_DEVICES%
echo ONEAPI_DEVICE_SELECTOR=%ONEAPI_DEVICE_SELECTOR%
) else (
echo Use all Intel GPUs, including iGPU ^& dGPU
set "GPUS_SETTING=-sm %SPLIT_MODE%"
)
echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m %MODEL_FILE% -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device %SYCL_DEVICES% --mmap
echo run cmd: ZES_ENABLE_SYSMAN=1 %BIN_FILE% -m %MODEL_FILE% -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap
set "ZES_ENABLE_SYSMAN=1"
%BIN_FILE% -m "%MODEL_FILE%" -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --device "%SYCL_DEVICES%" --mmap
%BIN_FILE% -m "%MODEL_FILE%" -no-cnv -p "%INPUT_PROMPT%" -n 200 -e -ngl %NGL% -s %SEED% -c %CONTEXT% %GPUS_SETTING% -lv %LOG_VERBOSE% --mmap
endlocal
-3
View File
@@ -1,3 +0,0 @@
llama-build-install
install
build
-13
View File
@@ -1,13 +0,0 @@
cmake_minimum_required(VERSION 3.14)
project(llama-simple)
set(CMAKE_CXX_STANDARD 17)
find_package(llama 0.1.0 REQUIRED)
add_executable(test-cmake test-cmake.cpp)
target_link_libraries(test-cmake PRIVATE llama)
target_compile_definitions(test-cmake PRIVATE
LLAMA_BUILD_NUMBER=${LLAMA_BUILD_NUMBER}
LLAMA_BUILD_COMMIT="${LLAMA_BUILD_COMMIT}"
)
-36
View File
@@ -1,36 +0,0 @@
## cmake-test
This is just for manually testing/developing of a llama.cpp installation to
enable troubleshooting issues and exploration. The idea is that this can be used
after making changes to llama.cpp installation cmake configuration and then
verify it locally.
### Usage
The following will configure, build, and install llama.cpp
Configuring/build/install:
```console
./build-install.sh
```
The above command will create a directory named `install` in the current directory
which will have the follwing files in its lib directory:
```console
(venv) $ ls install/lib/
cmake libggml.so libllama-common.so.0 libllama.so.0.1.0 llama.cpp
libggml-base.so libggml.so.0 libllama-common.so.0.1.0 libmtmd.so pkgconfig
libggml-base.so.0 libggml.so.0.19.0 libllama.so libmtmd.so.0
libggml-base.so.0.19.0 libllama-common.so libllama.so.0 libmtmd.so.0.1.0
```
Build/run this project using the installation created above:
```console
(venv) $ ./build.sh
-- Configuring done (0.0s)
-- Generating done (0.0s)
-- Build files have been written to: /home/danbev/work/ai/llama.cpp/examples/test-cmake/build
[100%] Built target test-cmake
[test-cmake] Using llama.cpp version 0.1.0-dev-b10335
[test-cmake] Initializing backend...
load_backend: loaded CPU backend from /home/danbev/work/ai/llama.cpp/examples/test-cmake/install/lib/llama.cpp/libggml-cpu-alderlake.so
[test-cmake] Backend initialized.
```
-19
View File
@@ -1,19 +0,0 @@
#!/bin/bash
set -e
rm -rf llama-build-install install
cmake --fresh -S ../../. -B llama-build-install -DCMAKE_BUILD_TYPE=Release \
-DBUILD_SHARED_LIBS=ON \
-DGGML_BACKEND_DL=ON \
-DGGML_CPU_ALL_VARIANTS=ON \
-DLLAMA_TESTS_INSTALL=OFF \
-DCMAKE_INSTALL_PREFIX="${PWD}/install" \
-DGGML_BACKEND_DIR="${PWD}/install/lib/llama.cpp" \
-DGGML_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
-DLLAMA_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \
-DLLAMA_TOOLS_INSTALL=OFF
cmake --build llama-build-install --parallel 12
cmake --install llama-build-install
-7
View File
@@ -1,7 +0,0 @@
#!/bin/bash
set -e
cmake -S . -B build -DCMAKE_PREFIX_PATH="${PWD}/install"
cmake --build build
LD_LIBRARY_PATH="${PWD}/install/lib/llama.cpp:${PWD}/install/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" ./build/test-cmake
-12
View File
@@ -1,12 +0,0 @@
#include "llama.h"
#include <cstdio>
int main(void) {
printf("[test-cmake] version: %s, build: %d (%s)\n",
llama_version(), LLAMA_BUILD_NUMBER, LLAMA_BUILD_COMMIT);
printf("[test-cmake] Initializing backend...\n");
llama_backend_init();
printf("[test-cmake] Backend initialized.\n");
llama_backend_free();
return 0;
}
+4 -7
View File
@@ -4,8 +4,8 @@ project("ggml" C CXX ASM)
### GGML Version
set(GGML_VERSION_MAJOR 0)
set(GGML_VERSION_MINOR 19)
set(GGML_VERSION_PATCH 0)
set(GGML_VERSION_MINOR 18)
set(GGML_VERSION_PATCH 1)
set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
@@ -341,9 +341,6 @@ set(GGML_PUBLIC_HEADERS
include/gguf.h)
set_target_properties(ggml PROPERTIES PUBLIC_HEADER "${GGML_PUBLIC_HEADERS}")
#if (GGML_METAL)
# set_target_properties(ggml PROPERTIES RESOURCE "${CMAKE_CURRENT_SOURCE_DIR}/src/ggml-metal.metal")
#endif()
install(TARGETS ggml LIBRARY PUBLIC_HEADER)
install(TARGETS ggml-base LIBRARY)
@@ -402,7 +399,7 @@ configure_package_config_file(
GGML_BIN_INSTALL_DIR)
write_basic_package_version_file(
${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
VERSION ${GGML_INSTALL_VERSION}
COMPATIBILITY SameMajorVersion)
@@ -414,7 +411,7 @@ message(STATUS "ggml version: ${GGML_INSTALL_VERSION}")
message(STATUS "ggml commit: ${GGML_BUILD_COMMIT}")
install(FILES ${CMAKE_CURRENT_BINARY_DIR}/ggml-config.cmake
${CMAKE_CURRENT_BINARY_DIR}/ggml-config-version.cmake
${CMAKE_CURRENT_BINARY_DIR}/ggml-version.cmake
DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/ggml)
if (MSVC)
+1 -6
View File
@@ -113,7 +113,6 @@ set_and_check(GGML_LIB_DIR "@PACKAGE_GGML_LIB_INSTALL_DIR@")
if(NOT TARGET ggml::ggml)
find_package(Threads REQUIRED)
unset(GGML_LIBRARY CACHE)
find_library(GGML_LIBRARY ggml
REQUIRED
HINTS ${GGML_LIB_DIR}
@@ -122,10 +121,8 @@ if(NOT TARGET ggml::ggml)
add_library(ggml::ggml UNKNOWN IMPORTED)
set_target_properties(ggml::ggml
PROPERTIES
IMPORTED_LOCATION "${GGML_LIBRARY}"
INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}")
IMPORTED_LOCATION "${GGML_LIBRARY}")
unset(GGML_BASE_LIBRARY CACHE)
find_library(GGML_BASE_LIBRARY ggml-base
REQUIRED
HINTS ${GGML_LIB_DIR}
@@ -135,7 +132,6 @@ if(NOT TARGET ggml::ggml)
set_target_properties(ggml::ggml-base
PROPERTIES
IMPORTED_LOCATION "${GGML_BASE_LIBRARY}"
INTERFACE_INCLUDE_DIRECTORIES "${GGML_INCLUDE_DIR}"
INTERFACE_LINK_LIBRARIES "${GGML_BASE_INTERFACE_LINK_LIBRARIES}")
set(_ggml_all_targets "")
@@ -144,7 +140,6 @@ if(NOT TARGET ggml::ggml)
string(REPLACE "-" "_" _ggml_backend_pfx "${_ggml_backend}")
string(TOUPPER "${_ggml_backend_pfx}" _ggml_backend_pfx)
unset(${_ggml_backend_pfx}_LIBRARY CACHE)
find_library(${_ggml_backend_pfx}_LIBRARY ${_ggml_backend}
REQUIRED
HINTS ${GGML_LIB_DIR}
-2
View File
@@ -154,8 +154,6 @@ extern "C" {
bool buffer_from_host_ptr;
// event synchronization
bool events;
// mmap is supported for loading
bool mmap_support;
};
// all the device properties
-6
View File
@@ -2788,12 +2788,6 @@ extern "C" {
struct ggml_cgraph * cgraph,
struct ggml_tensor * tensor);
// add the tensor and its parents to the graph without marking them for compute
// the flag is set later, when the tensor is reached from a node that computes
GGML_API void ggml_build_forward_order(
struct ggml_cgraph * cgraph,
struct ggml_tensor * tensor);
GGML_API void ggml_build_backward_expand(
struct ggml_context * ctx, // context for gradient computation
struct ggml_cgraph * cgraph,
-2
View File
@@ -132,7 +132,6 @@ static void ggml_backend_meta_device_get_props(ggml_backend_dev_t dev, ggml_back
/* .host_buffer = */ false, // Not implemented.
/* .buffer_from_host_ptr = */ false, // Not implemented.
/* .events = */ false, // Not implemented.
/* .mmap_support = */ true,
};
for (ggml_backend_dev_t simple_dev : meta_dev_ctx->simple_devs) {
ggml_backend_dev_props tmp_props;
@@ -141,7 +140,6 @@ static void ggml_backend_meta_device_get_props(ggml_backend_dev_t dev, ggml_back
props->caps.host_buffer = props->caps.host_buffer && tmp_props.caps.host_buffer;
props->caps.buffer_from_host_ptr = props->caps.buffer_from_host_ptr && tmp_props.caps.buffer_from_host_ptr;
props->caps.events = props->caps.events && tmp_props.caps.events;
props->caps.mmap_support = props->caps.mmap_support && tmp_props.caps.mmap_support;
}
}
-1
View File
@@ -367,7 +367,6 @@ static void ggml_backend_blas_device_get_props(ggml_backend_dev_t dev, struct gg
/* .host_buffer = */ false,
/* .buffer_from_host_ptr = */ true,
/* .events = */ false,
/* .mmap_support = */ true,
};
}
-1
View File
@@ -2815,7 +2815,6 @@ static void ggml_backend_cann_device_get_props(ggml_backend_dev_t dev, ggml_back
/* .host_buffer = */ host_buffer,
/* .buffer_from_host_ptr = */ false,
/* .events = */ true,
/* .mmap_support = */ true,
};
}
+68 -6
View File
@@ -1,19 +1,81 @@
#include "ggml-backend-impl.h"
#include "ggml-feats.h"
#if defined(__aarch64__) || defined(_M_ARM64)
#if defined(__aarch64__)
#if defined(__linux__)
#include <sys/auxv.h>
#elif defined(__APPLE__)
#include <sys/sysctl.h>
#endif
#if !defined(HWCAP2_SVE2)
#define HWCAP2_SVE2 (1 << 1)
#endif
#if !defined(HWCAP2_I8MM)
#define HWCAP2_I8MM (1 << 13)
#endif
#if !defined(HWCAP2_SME)
#define HWCAP2_SME (1 << 23)
#endif
struct aarch64_features {
// has_neon not needed, aarch64 has NEON guaranteed
bool has_dotprod = false;
bool has_fp16_va = false;
bool has_sve = false;
bool has_sve2 = false;
bool has_i8mm = false;
bool has_sme = false;
bool has_sme2 = false;
aarch64_features() {
#if defined(__linux__)
uint32_t hwcap = getauxval(AT_HWCAP);
uint32_t hwcap2 = getauxval(AT_HWCAP2);
has_dotprod = !!(hwcap & HWCAP_ASIMDDP);
has_fp16_va = !!(hwcap & HWCAP_FPHP);
has_sve = !!(hwcap & HWCAP_SVE);
has_sve2 = !!(hwcap2 & HWCAP2_SVE2);
has_i8mm = !!(hwcap2 & HWCAP2_I8MM);
has_sme = !!(hwcap2 & HWCAP2_SME);
#elif defined(__APPLE__)
int oldp = 0;
size_t size = sizeof(oldp);
if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, NULL, 0) == 0) {
has_dotprod = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, NULL, 0) == 0) {
has_i8mm = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, NULL, 0) == 0) {
has_sme = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, NULL, 0) == 0) {
has_sme2 = static_cast<bool>(oldp);
}
// Apple apparently does not implement SVE yet
#endif
}
};
static int ggml_backend_cpu_aarch64_score() {
int score = 1;
const ggml_feats_arch64_runtime_t af = ggml_feats_get_arch64_runtime();
GGML_UNUSED(af);
aarch64_features af;
#ifdef GGML_USE_DOTPROD
if (!af.has_dotprod) { return 0; }
score += 1<<1;
#endif
#ifdef GGML_USE_FP16_VECTOR_ARITHMETIC
if (!af.has_fp16) { return 0; }
if (!af.has_fp16_va) { return 0; }
score += 1<<2;
#endif
#ifdef GGML_USE_SVE
@@ -38,4 +100,4 @@ static int ggml_backend_cpu_aarch64_score() {
GGML_BACKEND_DL_SCORE_IMPL(ggml_backend_cpu_aarch64_score)
# endif // defined(__aarch64__) || defined(_M_ARM64)
# endif // defined(__aarch64__)
+1 -1
View File
@@ -2608,7 +2608,7 @@ static bool ggml_thread_apply_priority(int32_t prio) {
return true;
}
#elif defined(__linux__)
#elif defined(__gnu_linux__)
// TODO: this may not work on BSD, to be verified
static bool ggml_thread_apply_affinity(const bool * mask) {
-1
View File
@@ -397,7 +397,6 @@ static void ggml_backend_cpu_device_get_props(ggml_backend_dev_t dev, struct ggm
/* .host_buffer = */ false,
/* .buffer_from_host_ptr = */ true,
/* .events = */ false,
/* .mmap_support = */ true,
};
}
+99 -213
View File
@@ -2,12 +2,10 @@
// SPDX-License-Identifier: MIT
//
#include <arm_neon.h>
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <assert.h>
#include <stdio.h>
#include <atomic>
#include <cfloat>
#include <cctype>
#include <algorithm>
#include <cmath>
#include <stdexcept>
@@ -19,21 +17,25 @@
#include <cstddef>
#include <cstdint>
#include <fstream>
#include <map>
#include <set>
#include <iostream>
#include <climits>
#include <charconv>
#include <system_error>
#if defined(__linux__)
#include <asm/hwcap.h>
#include <dirent.h>
#include <sys/auxv.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <unistd.h>
#ifndef HWCAP2_SME2
#define HWCAP2_SME2 (1UL << 37)
#endif
#elif defined(__APPLE__)
#include <string_view>
#include <sys/sysctl.h>
#include <sys/types.h>
#elif defined(_WIN32)
#include <windows.h>
#include <excpt.h>
#endif
#include "kleidiai.h"
@@ -41,7 +43,6 @@
#include "ggml-cpu.h"
#include "ggml-cpu-impl.h"
#include "ggml-impl.h"
#include "ggml-feats.h"
#include "ggml-backend-impl.h"
#include "ggml-threading.h"
#include "traits.h"
@@ -63,8 +64,8 @@ struct ggml_kleidiai_context {
ggml_kleidiai_kernels * kernels_q4;
ggml_kleidiai_kernels * kernels_q8;
ggml_kleidiai_kernels * kernels_f32;
int sme_thread_cap; // <= 0 means "SME disabled/unknown"
int thread_hint; // <= 0 means "no hint"
int sme_thread_cap; // <= 0 means SME disabled/unknown”;
int thread_hint; // <= 0 means no hint
int chunk_multiplier;
} static ctx = { CPU_FEATURE_NONE, nullptr, nullptr, nullptr, 0, -1, 4 };
@@ -92,117 +93,24 @@ static const char* cpu_feature_to_string(cpu_feature f) {
}
}
#if defined(__linux__) && defined(__aarch64__)
static bool parse_cpu_dir_name(const char* name, size_t* cpu) {
if (strncmp(name, "cpu", 3) != 0 ||
name[3] < '0' || name[3] > '9') {
return false;
}
const char* first = name + 3;
const char* last = name + strlen(name);
size_t value = 0;
const auto [end, ec] = std::from_chars(first, last, value, 10);
if (ec != std::errc{} || end != last) {
return false;
}
*cpu = value;
return true;
}
static std::vector<size_t> detect_cpu_ids() {
std::vector<size_t> cpus;
DIR * dir = opendir("/sys/devices/system/cpu");
if (dir == nullptr) {
return cpus;
}
while (dirent * entry = readdir(dir)) {
size_t cpu = 0;
if (parse_cpu_dir_name(entry->d_name, &cpu)) {
cpus.push_back(cpu);
}
}
closedir(dir);
std::sort(cpus.begin(), cpus.end());
cpus.erase(std::unique(cpus.begin(), cpus.end()), cpus.end());
return cpus;
}
#endif
#if defined(__APPLE__) && defined(__aarch64__)
static bool apple_sme_counted_perf_level(std::string name) {
for (std::string::size_type i = 0; i < name.size(); ++i) {
name[i] = (char) std::tolower((unsigned char) name[i]);
}
// Conservative ceiling: only count perf-level names observed to provide full SME throughput.
// Future names should be calibrated here before they raise the automatic SME thread cap.
return name.find("super") != std::string::npos ||
name.find("performance") != std::string::npos;
}
#endif
static void add_smcus_from_smidr(uint64_t smidr, size_t & num_private, std::map<uint32_t, size_t> & shared_counts) {
// Arm ARM: SMIDR_EL1. SH==0 is implementation-defined; keep the existing
// conservative policy and only treat zero affinity as private.
const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
const uint32_t nsmc = (uint32_t)((smidr >> 56) & 0xF);
const size_t shared_count = nsmc == 0xF ? 1 : (size_t)nsmc + 1;
const uint32_t affinity = (uint32_t)(smidr & 0xFFFu);
const uint32_t affinity2 = (uint32_t)((smidr >> 32) & 0xFFFFFu);
const uint32_t id = (affinity2 << 12) | affinity;
if (nsmc == 0xF) {
GGML_LOG_WARN("kleidiai: NSMC detected as 0xF indicating reseved value, setting min safe shared SMCU count to 1");
}
switch (sh) {
case 2: // private SMCU
++num_private;
break;
case 3: // shared SMCU
if (shared_counts[id] < shared_count) {
shared_counts[id] = shared_count;
}
break;
case 0:
if (id == 0) {
++num_private;
} else if (shared_counts[id] < shared_count) {
shared_counts[id] = shared_count;
}
break;
default:
break;
}
}
static size_t detect_num_smcus() {
const auto runtime_feat = ggml_feats_get_arch64_runtime();
if (!runtime_feat.has_sme) {
if (!ggml_cpu_has_sme()) {
return 0;
}
#if defined(__linux__) && defined(__aarch64__)
// Linux/aarch64: Best-effort count of Streaming Mode Compute Units (SMCUs) via SMIDR_EL1 sysfs.
size_t num_private = 0;
std::map<uint32_t, size_t> shared_counts;
std::set<uint32_t> shared_ids;
const std::vector<size_t> cpus = detect_cpu_ids();
for (const size_t cpu : cpus) {
for (size_t cpu = 0;; ++cpu) {
const std::string path =
"/sys/devices/system/cpu/cpu" + std::to_string(cpu) +
"/regs/identification/smidr_el1";
std::ifstream file(path);
if (!file.is_open()) {
continue;
break;
}
uint64_t smidr = 0;
@@ -210,69 +118,54 @@ static size_t detect_num_smcus() {
continue;
}
add_smcus_from_smidr(smidr, num_private, shared_counts);
// Arm ARM: SMIDR_EL1
const uint32_t sh = (uint32_t)((smidr >> 13) & 0x3);
// Build an "affinity-like" identifier for shared SMCUs.
// Keep the original packing logic, but isolate it here.
const uint32_t id = (uint32_t)((smidr & 0xFFFu) | ((smidr >> 20) & 0xFFFFF000u));
switch (sh) {
case 0b10: // private SMCU
++num_private;
break;
case 0b11: // shared SMCU
shared_ids.emplace(id);
break;
case 0b00:
// Ambiguous / implementation-defined. Be conservative:
// treat id==0 as private, otherwise as shared.
if (id == 0) ++num_private;
else shared_ids.emplace(id);
break;
default:
break;
}
}
size_t total = num_private;
for (const auto & entry : shared_counts) {
total += entry.second;
}
return total;
return num_private + shared_ids.size();
#elif defined(__APPLE__) && defined(__aarch64__)
int perf_levels = 0;
size_t size = sizeof(perf_levels);
if (sysctlbyname("hw.nperflevels", &perf_levels, &size, nullptr, 0) != 0 ||
size != sizeof(perf_levels) || perf_levels <= 0) {
return 0;
}
// table for known M4 variants. Users can override via GGML_KLEIDIAI_SME=<n>.
char chip_name[256] = {};
size_t size = sizeof(chip_name);
size_t units = 0;
for (int i = 0; i < perf_levels; ++i) {
char key[64] = {};
int physical_cpus = 0;
int cpus_per_l2 = 0;
if (sysctlbyname("machdep.cpu.brand_string", chip_name, &size, nullptr, 0) == 0) {
const std::string brand(chip_name);
snprintf(key, sizeof(key), "hw.perflevel%d.physicalcpu", i);
size = sizeof(physical_cpus);
if (sysctlbyname(key, &physical_cpus, &size, nullptr, 0) != 0 ||
size != sizeof(physical_cpus) || physical_cpus <= 0) {
continue;
}
struct ModelSMCU { const char *match; size_t smcus; };
static const ModelSMCU table[] = {
{ "M4 Ultra", 2 },
{ "M4 Max", 2 },
{ "M4 Pro", 2 },
{ "M4", 1 },
};
snprintf(key, sizeof(key), "hw.perflevel%d.cpusperl2", i);
size = sizeof(cpus_per_l2);
if (sysctlbyname(key, &cpus_per_l2, &size, nullptr, 0) != 0 ||
size != sizeof(cpus_per_l2) || cpus_per_l2 <= 0) {
continue;
}
snprintf(key, sizeof(key), "hw.perflevel%d.name", i);
size = 0;
if (sysctlbyname(key, nullptr, &size, nullptr, 0) != 0 || size == 0) {
continue;
}
std::string name(size, '\0');
if (sysctlbyname(key, &name[0], &size, nullptr, 0) != 0) {
continue;
}
name.resize(size);
while (!name.empty() && name.back() == '\0') {
name.pop_back();
}
if (apple_sme_counted_perf_level(name)) {
units += (size_t) ((physical_cpus + cpus_per_l2 - 1) / cpus_per_l2);
for (const auto &e : table) {
if (brand.find(e.match) != std::string::npos) {
return e.smcus;
}
}
}
return units;
#elif defined(_WIN32) && (defined(_M_ARM64) || defined(__aarch64__))
// No verified Windows arm64 SMCU detection path yet. Return unknown and use
// GGML_KLEIDIAI_SME=N as a diagnostics/debug override for SME thread cap
// calibration until a detection mechanism is verified on real hardware.
return 0;
#else
@@ -305,18 +198,15 @@ static void init_kleidiai_context(void) {
if (!initialized) {
initialized = true;
// Optional diagnostics/debug overrides; production defaults come from runtime detection.
const char *env_sme = getenv("GGML_KLEIDIAI_SME");
const char *env_threads = getenv("GGML_TOTAL_THREADS");
const char *env_chunk_mult = getenv("GGML_KLEIDIAI_CHUNK_MULTIPLIER");
const auto runtime_feat = ggml_feats_get_arch64_runtime();
size_t detected_smcus = 0;
ctx.features = (runtime_feat.has_dotprod ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
(runtime_feat.has_i8mm ? CPU_FEATURE_I8MM : CPU_FEATURE_NONE) |
(runtime_feat.sve_cnt == QK8_0 ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);
ctx.features = (ggml_cpu_has_dotprod() ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
(ggml_cpu_has_matmul_int8() ? CPU_FEATURE_I8MM : CPU_FEATURE_NONE) |
((ggml_cpu_has_sve() && ggml_cpu_get_sve_cnt() == QK8_0) ? CPU_FEATURE_SVE : CPU_FEATURE_NONE);
if (env_threads) {
bool ok = false;
@@ -334,54 +224,54 @@ static void init_kleidiai_context(void) {
}
}
// SME policy:
// - env unset => auto-detect SMCUs; enable SME only if detected > 0.
// - env=0 => force off.
// - env>0 => force N cores, if the binary was built with SME.
int sme_cores = 0;
bool sme_env_ok = false;
bool sme_env_set = (env_sme != nullptr);
const bool has_supported_sme_family = runtime_feat.has_sme;
bool sme_cap_detected = false;
if (has_supported_sme_family) {
detected_smcus = detect_num_smcus();
sme_cap_detected = detected_smcus > 0;
// Some platforms expose SME without exposing a calibrated SMCU count.
// Use one SME thread as the conservative default; add platform SMCU detection to raise it.
sme_cores = sme_cap_detected ? (int)detected_smcus : 1;
if (!sme_env_set && !sme_cap_detected) {
GGML_LOG_INFO("kleidiai: SME detected; SMCU count unavailable, using conservative SME thread cap=1\n");
}
}
// Runtime-detect SME support and available SMCUs first. The detected SMCU
// count is used as the SME thread cap, and GGML_KLEIDIAI_SME can debug-override that:
// - unset: use runtime detection.
// - 0: disable SME-family kernels.
// - N > 0: use N as the SME thread cap, if an SME-family kernel is selectable.
if (sme_env_set) {
bool ok = false;
int v = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
sme_env_ok = ok;
if (ok) {
if (has_supported_sme_family) {
sme_cores = v;
} else {
if (v > 0) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but SME is not supported on this CPU; disabling SME-family kernels\n", v);
}
sme_cores = 0;
}
if (!ok) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; falling back to runtime SME-core detection\n");
detected_smcus = detect_num_smcus();
sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
} else if (v == 0) {
sme_cores = 0;
} else if (!ggml_cpu_has_sme()) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but the binary was not built with SME; disabling SME\n", v);
sme_cores = 0;
} else {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; using automatic SME thread cap\n");
sme_cores = v;
}
} else {
detected_smcus = detect_num_smcus();
sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
}
if (sme_cores > 0 && has_supported_sme_family) {
if (!sme_env_set && ggml_cpu_has_sme() && sme_cores == 0) {
GGML_LOG_WARN("kleidiai: runtime SME-core detection returned 0; falling back to NEON\n");
}
if (sme_cores > 0) {
ctx.features |= CPU_FEATURE_SME;
if (runtime_feat.has_sme2) {
#if defined(__aarch64__) && defined(__linux__)
// ARM guarantees SME2 implies SME, so only check SME2 when SME is enabled.
if (getauxval(AT_HWCAP2) & HWCAP2_SME2) {
ctx.features |= CPU_FEATURE_SME2;
}
#elif defined(__aarch64__) && defined(__APPLE__)
int feat_sme2 = 0;
size_t size = sizeof(feat_sme2);
if (sysctlbyname("hw.optional.arm.FEAT_SME2", &feat_sme2, &size, NULL, 0) == 0 && feat_sme2) {
ctx.features |= CPU_FEATURE_SME2;
}
#endif
}
// Kernel selection
@@ -407,19 +297,16 @@ static void init_kleidiai_context(void) {
GGML_LOG_INFO("kleidiai: primary f32 kernel feature %s\n", cpu_feature_to_string(ctx.kernels_f32->required_cpu));
}
const bool has_selected_sme_family_kernel =
(ctx.kernels_q4 && is_sme_family(ctx.kernels_q4->required_cpu)) ||
(ctx.kernels_q8 && is_sme_family(ctx.kernels_q8->required_cpu)) ||
(ctx.kernels_f32 && is_sme_family(ctx.kernels_f32->required_cpu));
ctx.sme_thread_cap = has_selected_sme_family_kernel ? sme_cores : 0;
ctx.sme_thread_cap = (ctx.features & CPU_FEATURE_SME) ? sme_cores : 0;
if (has_selected_sme_family_kernel) {
if (ctx.features & CPU_FEATURE_SME) {
const bool has_sme2 = (ctx.features & CPU_FEATURE_SME2) != CPU_FEATURE_NONE;
if (sme_env_set && sme_env_ok && sme_cores > 0) {
GGML_LOG_INFO("kleidiai: SME enabled (GGML_KLEIDIAI_SME=%d debug override)\n", sme_cores);
} else if (sme_cap_detected) {
GGML_LOG_INFO("kleidiai: SME enabled (runtime-detected SME thread cap=%d)\n", sme_cores);
GGML_LOG_INFO("kleidiai: SME%s enabled (GGML_KLEIDIAI_SME=%d override)\n",
has_sme2 ? "2" : "", sme_cores);
} else {
GGML_LOG_INFO("kleidiai: SME enabled (runtime SME detected, conservative thread cap=%d)\n", sme_cores);
GGML_LOG_INFO("kleidiai: SME%s enabled (runtime-detected SME cores=%d)\n",
has_sme2 ? "2" : "", sme_cores);
}
} else {
GGML_LOG_INFO("kleidiai: SME disabled\n");
@@ -580,7 +467,7 @@ static int kleidiai_collect_kernel_chain_common(
}
if (is_sme_family(primary->required_cpu)) {
const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~(CPU_FEATURE_SME | CPU_FEATURE_SME2));
const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~CPU_FEATURE_SME & ~CPU_FEATURE_SME2);
if (fallback_mask != CPU_FEATURE_NONE) {
ggml_kleidiai_kernels * fallback = select_fallback(fallback_mask);
if (fallback && fallback != primary &&
@@ -1190,14 +1077,13 @@ class tensor_traits : public ggml::cpu::tensor_traits {
const int ith_total = params->ith;
int sme_slot = -1;
int non_sme_slot = -1;
for (int i = 0; i < runtime_count; ++i) {
if (is_sme_family(runtime[i].kernels->required_cpu)) {
sme_slot = i;
break;
}
}
int non_sme_slot = -1;
for (int i = 0; i < runtime_count; ++i) {
if (!is_sme_family(runtime[i].kernels->required_cpu)) {
non_sme_slot = i;
-2
View File
@@ -195,7 +195,6 @@ template <typename BLOC_TYPE, int64_t INTER_SIZE, int64_t NB_COLS> class tensor_
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q5_K:
//case GGML_TYPE_MXFP4:
@@ -215,7 +214,6 @@ template <typename BLOC_TYPE, int64_t INTER_SIZE, int64_t NB_COLS> class tensor_
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q5_K:
//case GGML_TYPE_MXFP4:
+22 -22
View File
@@ -253,9 +253,9 @@ static void ggml_cpy_f32_q8_0_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK8_0 == 0);
const int64_t num_blocks = (ne/QK8_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK8_0;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_q8_0, QK8_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_q8_0, QK8_0><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -264,9 +264,9 @@ static void ggml_cpy_q8_0_f32_cuda(
const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02,
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
const int64_t num_blocks = (ne/QK8_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_q_f32<cpy_blck_q8_0_f32, QK8_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_q_f32<cpy_blck_q8_0_f32, QK8_0><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -276,9 +276,9 @@ static void ggml_cpy_f32_q4_0_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK4_0 == 0);
const int64_t num_blocks = (ne/QK4_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK4_0;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_q4_0, QK4_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_q4_0, QK4_0><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -289,9 +289,9 @@ static void ggml_cpy_q4_0_f32_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
cudaStream_t stream) {
const int64_t num_blocks = (ne/QK4_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0><<<num_blocks, 1, 0, stream>>>(
cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -302,9 +302,9 @@ static void ggml_cpy_f32_q4_1_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK4_1 == 0);
const int64_t num_blocks = (ne/QK4_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK4_1;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_q4_1, QK4_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_q4_1, QK4_1><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -315,9 +315,9 @@ static void ggml_cpy_q4_1_f32_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
cudaStream_t stream) {
const int64_t num_blocks = (ne/QK4_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1><<<num_blocks, 1, 0, stream>>>(
cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -328,9 +328,9 @@ static void ggml_cpy_f32_q5_0_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK5_0 == 0);
const int64_t num_blocks = (ne/QK5_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK5_0;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_q5_0, QK5_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_q5_0, QK5_0><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -341,9 +341,9 @@ static void ggml_cpy_q5_0_f32_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
cudaStream_t stream) {
const int64_t num_blocks = (ne/QK5_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0><<<num_blocks, 1, 0, stream>>>(
cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -354,9 +354,9 @@ static void ggml_cpy_f32_q5_1_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK5_1 == 0);
const int64_t num_blocks = (ne/QK5_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK5_1;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_q5_1, QK5_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_q5_1, QK5_1><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -367,9 +367,9 @@ static void ggml_cpy_q5_1_f32_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
cudaStream_t stream) {
const int64_t num_blocks = (ne/QK5_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1><<<num_blocks, 1, 0, stream>>>(
cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
@@ -380,9 +380,9 @@ static void ggml_cpy_f32_iq4_nl_cuda(
const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
GGML_ASSERT(ne % QK4_NL == 0);
const int64_t num_blocks = (ne/QK4_NL + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
const int64_t num_blocks = ne / QK4_NL;
GGML_ASSERT(num_blocks <= INT_MAX);
cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL><<<num_blocks, 1, 0, stream>>>
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
+8 -130
View File
@@ -1865,37 +1865,6 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
}
// returns true when ggml_cuda_mul_mat_id takes the fallback path that requires stream synchronization
// [TAG_MUL_MAT_ID_CUDA_GRAPHS]
static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int cc) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
return true;
}
if (dst->ne[2] <= MMVQ_MAX_BATCH_SIZE) {
if (ggml_is_quantized(src0->type)) {
if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc)) {
return false;
}
} else if (GGML_CUDA_CC_IS_AMD(cc)) {
return false;
}
}
if (ggml_cuda_should_use_mmq(src0->type, cc, src1->ne[2], /*n_experts=*/src0->ne[2])) {
return false;
}
if (ggml_cuda_should_use_mmf(src0->type, cc, WARP_SIZE, src0->ne, src0->nb, src1->ne[2], /*mul_mat_id=*/true)) {
return false;
}
return true;
}
static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
@@ -1938,7 +1907,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
}
// note: this path should not be reached when recording CUDA graphs, because it requires stream synchronization
GGML_ASSERT(ggml_cuda_mul_mat_id_needs_sync(dst, cc));
// TODO: add asserts to verify this. should work with CUDA, HIP, etc.
cudaStream_t stream = ctx.stream();
GGML_ASSERT(nb12 % nb11 == 0);
@@ -2553,8 +2522,10 @@ static bool ggml_cuda_graph_check_compability(ggml_cgraph * cgraph) {
// [TAG_MUL_MAT_ID_CUDA_GRAPHS]
if (node->op == GGML_OP_MUL_MAT_ID) {
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
if (ggml_cuda_mul_mat_id_needs_sync(node, cc)) {
// the mul_mat_id fallback path synchronizes the stream, so we cannot use CUDA graphs
const int mmvq_mmid_max = get_mmvq_mmid_max_batch(node->src[0]->type, cc);
if (!ggml_is_quantized(node->src[0]->type) || node->ne[2] > mmvq_mmid_max) {
// under these conditions, the mul_mat_id operation will need to synchronize the stream, so we cannot use CUDA graphs
// TODO: figure out a way to enable for larger batch sizes, without hurting performance
// ref: https://github.com/ggml-org/llama.cpp/pull/18958
use_cuda_graph = false;
#ifndef NDEBUG
@@ -2680,52 +2651,6 @@ static bool ggml_cuda_should_fuse_rope_set_rows(const ggml_tensor * rope,
return true;
}
static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm,
const ggml_tensor * mul,
const ggml_tensor * rope) {
if (rms_norm->op != GGML_OP_RMS_NORM || mul->op != GGML_OP_MUL || rope->op != GGML_OP_ROPE) {
return false;
}
if (rms_norm->src[0]->type != GGML_TYPE_F32 || rms_norm->type != GGML_TYPE_F32 ||
mul->src[0]->type != GGML_TYPE_F32 || mul->src[1]->type != GGML_TYPE_F32 ||
mul->type != GGML_TYPE_F32 || rope->type != GGML_TYPE_F32) {
return false;
}
if (rope->src[0] != mul) {
return false;
}
//if rms norm is the B operand, then we don't handle broadcast
if (rms_norm == mul->src[1] && !ggml_are_same_shape(mul->src[0], rms_norm)) {
return false;
}
if (!ggml_are_same_shape(rms_norm, mul)) {
return false;
}
//rms_norm kernel assumes contiguous rows
if (!ggml_is_contiguous_rows(rms_norm->src[0]) ||
!ggml_is_contiguous_rows(mul->src[0]) || !ggml_is_contiguous_rows(mul->src[1])) {
return false;
}
// the fused kernel handles the norm/neox rope modes only
const int mode = ((const int32_t *) rope->op_params)[2];
if (mode != GGML_ROPE_TYPE_NORMAL && mode != GGML_ROPE_TYPE_NEOX) {
return false;
}
const int n_dims = ((const int32_t *) rope->op_params)[1];
if (n_dims % 2 != 0 || rope->src[0]->ne[0] % 2 != 0) {
return false;
}
return true;
}
// match gated_delta_net + the strided cpy that scatters its state snapshots into the cache
// (slot i -> rollback group i, slot 0 newest), so the kernel can write them and skip the cpy.
static int ggml_cuda_try_gdn_cache_fusion(
@@ -3055,36 +2980,6 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
}
}
std::initializer_list<enum ggml_op> rms_norm_mul_rope_ops = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE };
std::initializer_list<enum ggml_op> rms_norm_mul_rope_set_rows_ops = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS };
if (is_equal(rms_norm_mul_rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 4 })) {
const ggml_tensor * rms_norm = cgraph->nodes[node_idx];
const ggml_tensor * mul = cgraph->nodes[node_idx + 1];
const ggml_tensor * rope = cgraph->nodes[node_idx + 2];
const ggml_tensor * view = cgraph->nodes[node_idx + 3];
const ggml_tensor * set_rows = cgraph->nodes[node_idx + 4];
if (ggml_check_edges(cgraph, node_idx, {{1, 0, 0}, {2, 0, 1}, {3, 0, 2}, {4, 0, 3}}) &&
ggml_cuda_should_fuse_rms_norm_mul_rope(rms_norm, mul, rope) &&
ggml_cuda_should_fuse_rope_set_rows(rope, view, set_rows)) {
int out_nodes[] = { node_idx + 4 };
return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1);
}
}
if (is_equal(rms_norm_mul_rope_ops, ops) && ggml_can_fuse(cgraph, node_idx, ops)) {
const ggml_tensor * rms_norm = cgraph->nodes[node_idx];
const ggml_tensor * mul = cgraph->nodes[node_idx + 1];
const ggml_tensor * rope = cgraph->nodes[node_idx + 2];
if (ggml_cuda_should_fuse_rms_norm_mul_rope(rms_norm, mul, rope)) {
int out_nodes[] = { node_idx + 2 };
return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1);
}
return false;
}
std::initializer_list<enum ggml_op> rope_set_rows_ops = { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS };
if (is_equal(rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) {
@@ -3093,8 +2988,7 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
const ggml_tensor * set_rows = cgraph->nodes[node_idx + 2];
if (ggml_cuda_should_fuse_rope_set_rows(rope, view, set_rows)) {
int out_nodes[] = { node_idx + 2 };
return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1);
return true;
}
}
@@ -3946,16 +3840,6 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
return fused_node_count - 1;
}
if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }, {})) {
ggml_cuda_op_rms_norm_mul_rope_fused(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 4]);
return 4;
}
if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE }, {})) {
ggml_cuda_op_rms_norm_mul_rope_fused(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2], nullptr);
return 2;
}
if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD }, {})) {
ggml_cuda_op_rms_norm_fused_add(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2]);
return 2;
@@ -4149,11 +4033,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
continue;
}
#ifndef NDEBUG
// On integrated GPUs (APUs, e.g. RDNA3.5) the scheduler may place a
// node's output on the host-visible buffer, which the compute path
// handles. Allow that here, mirroring the src-tensor check below.
assert(node->buffer->buft == ggml_backend_cuda_buffer_type(cuda_ctx->device) ||
(integrated && ggml_backend_buft_is_cuda_host(node->buffer->buft)));
assert(node->buffer->buft == ggml_backend_cuda_buffer_type(cuda_ctx->device));
for (int j = 0; j < GGML_MAX_SRC; j++) {
if (node->src[j] != nullptr) {
assert(node->src[j]->buffer);
@@ -4830,7 +4710,6 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back
/* .host_buffer = */ host_buffer,
/* .buffer_from_host_ptr = */ false,
/* .events = */ events,
/* .mmap_support = */ props->type != GGML_BACKEND_DEVICE_TYPE_IGPU,
};
}
@@ -5215,7 +5094,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
return max_bias == 0.0f;
}
case GGML_OP_ROLL:
if(op->src[0]->type == GGML_TYPE_F32 && ggml_is_contiguous(op->src[0])) {
if(op->src[0]->type == GGML_TYPE_F32) {
return true;
}
return false;
@@ -5326,7 +5205,6 @@ static bool ggml_backend_cuda_device_offload_op(ggml_backend_dev_t dev, const gg
static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_t dev) {
#ifdef GGML_CUDA_NO_PEER_COPY
GGML_UNUSED(dev);
return nullptr;
#else
ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *)dev->context;
+1 -1
View File
@@ -8,6 +8,7 @@ struct __builtin_align__(32) float8 {
float x; float y; float z; float w;
float p; float q; float r; float s;
};
#endif
#if CUDART_VERSION >= 12080
static __device__ __forceinline__ float nvfp4_native_scale_error(
@@ -48,7 +49,6 @@ static __device__ __forceinline__ float nvfp4_native_scale_error(
return err;
}
#endif // CUDART_VERSION >= 12080
#endif // defined(BLACKWELL_MMA_AVAILABLE)
__launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1)
static __global__ void quantize_q8_1(
-235
View File
@@ -670,238 +670,3 @@ void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst)
void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * rope, ggml_tensor * set_rows) {
ggml_cuda_op_rope_impl<true>(ctx, rope, set_rows);
}
// fused RMS_NORM + MUL + ROPE (+ VIEW + SET_ROWS)
// one block per row: block_reduce gives the norm scale, then each thread applies mul and rope to the elements it owns
template <int block_size, bool has_ff, typename D>
static __global__ void rms_norm_mul_rope_f32(
const float * x, D * dst, const int ncols,
const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t s1, const int64_t s2, const int64_t s3,
const float eps,
const float * mul,
const int64_t mul_s01, const int64_t mul_s02, const int64_t mul_s03,
const uint3 mul_ncols_packed, const uint3 mul_nrows_packed,
const uint3 mul_nchannels_packed, const uint3 mul_nsamples_packed,
const int n_dims, const int32_t * pos,
const float freq_scale, const float ext_factor, const float attn_factor,
const rope_corr_dims corr_dims, const float theta_scale,
const float * freq_factors,
const int64_t * row_indices, const int set_rows_stride,
const bool is_neox) {
ggml_cuda_pdl_lc();
const int row = blockIdx.x;
const int channel = blockIdx.y;
const int sample = blockIdx.z;
const int tid = threadIdx.x;
x += sample*s03 + channel*s02 + row*s01;
const uint32_t mul_row = fastmodulo(row, mul_nrows_packed);
const uint32_t mul_channel = fastmodulo(channel, mul_nchannels_packed);
const uint32_t mul_sample = fastmodulo(sample, mul_nsamples_packed);
mul += mul_sample*mul_s03 + mul_channel*mul_s02 + mul_row*mul_s01;
float tmp = 0.0f;
ggml_cuda_pdl_sync();
for (int col = tid; col < ncols; col += block_size) {
const float xi = x[col];
tmp += xi * xi;
}
extern __shared__ float s_sum[];
tmp = block_reduce<block_reduce_method::SUM, block_size>(tmp, s_sum);
const float scale = rsqrtf(tmp/ncols + eps);
int64_t idst = sample*s3 + channel*s2 + row*s1;
if (set_rows_stride != 0) {
idst = row*s1 + row_indices[channel]*set_rows_stride;
}
dst += idst;
for (int i0 = 2*tid; i0 < ncols; i0 += 2*block_size) {
int ix0;
int ix1;
if (is_neox && i0 < n_dims) {
ix0 = i0/2;
ix1 = i0/2 + n_dims/2;
} else {
ix0 = i0 + 0;
ix1 = i0 + 1;
}
const float x0 = scale * x[ix0] * mul[fastmodulo(ix0, mul_ncols_packed)];
const float x1 = scale * x[ix1] * mul[fastmodulo(ix1, mul_ncols_packed)];
if (i0 >= n_dims) {
dst[ix0] = ggml_cuda_cast<D>(x0);
dst[ix1] = ggml_cuda_cast<D>(x1);
continue;
}
const float theta_base = pos[channel]*powf(theta_scale, i0/2.0f);
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn<true>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
dst[ix0] = ggml_cuda_cast<D>(x0*cos_theta - x1*sin_theta);
dst[ix1] = ggml_cuda_cast<D>(x0*sin_theta + x1*cos_theta);
}
}
template <typename D>
static void rms_norm_mul_rope_cuda(
const float * x, D * dst,
const int ncols, const int nrows, const int nchannels, const int nsamples,
const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t s1, const int64_t s2, const int64_t s3,
const float eps,
const float * mul,
const int64_t mul_s01, const int64_t mul_s02, const int64_t mul_s03,
const uint32_t mul_ncols, const uint32_t mul_nrows,
const uint32_t mul_nchannels, const uint32_t mul_nsamples,
const int n_dims, const int32_t * pos,
const float freq_scale, const float freq_base, const float ext_factor, const float attn_factor,
const rope_corr_dims corr_dims,
const float * freq_factors,
const int64_t * row_indices, const int set_rows_stride,
const bool is_neox, cudaStream_t stream) {
GGML_ASSERT(ncols % 2 == 0);
const dim3 blocks_num(nrows, nchannels, nsamples);
const float theta_scale = powf(freq_base, -2.0f/n_dims);
const uint3 mul_ncols_packed = init_fastdiv_values(mul_ncols);
const uint3 mul_nrows_packed = init_fastdiv_values(mul_nrows);
const uint3 mul_nchannels_packed = init_fastdiv_values(mul_nchannels);
const uint3 mul_nsamples_packed = init_fastdiv_values(mul_nsamples);
if (ncols < 1024) {
const dim3 block_dims(256, 1, 1);
const ggml_cuda_kernel_launch_params launch_params = {blocks_num, block_dims, 32*sizeof(float), stream};
if (freq_factors == nullptr) {
ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<256, false, D>, launch_params,
x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03,
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale,
freq_factors, row_indices, set_rows_stride, is_neox);
} else {
ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<256, true, D>, launch_params,
x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03,
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale,
freq_factors, row_indices, set_rows_stride, is_neox);
}
} else {
const dim3 block_dims(1024, 1, 1);
const ggml_cuda_kernel_launch_params launch_params = {blocks_num, block_dims, 32*sizeof(float), stream};
if (freq_factors == nullptr) {
ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<1024, false, D>, launch_params,
x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03,
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale,
freq_factors, row_indices, set_rows_stride, is_neox);
} else {
ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<1024, true, D>, launch_params,
x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03,
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale,
freq_factors, row_indices, set_rows_stride, is_neox);
}
}
}
void ggml_cuda_op_rms_norm_mul_rope_fused(ggml_backend_cuda_context & ctx,
ggml_tensor * rms_norm, ggml_tensor * mul, ggml_tensor * rope, ggml_tensor * set_rows) {
const ggml_tensor * x = rms_norm->src[0];
const ggml_tensor * mul_src = mul->src[0] == rms_norm ? mul->src[1] : mul->src[0];
float eps = 0.0f;
memcpy(&eps, rms_norm->op_params, sizeof(float));
GGML_ASSERT(eps >= 0.0f);
GGML_ASSERT(x->type == GGML_TYPE_F32);
GGML_ASSERT(mul_src->type == GGML_TYPE_F32);
GGML_ASSERT(rope->type == GGML_TYPE_F32);
void * dst_d = rope->data;
ggml_type dst_type = rope->type;
const int64_t * row_indices = nullptr;
int set_rows_stride = 0;
if (set_rows != nullptr) {
dst_d = set_rows->data;
dst_type = set_rows->type;
row_indices = (const int64_t *) set_rows->src[1]->data;
set_rows_stride = set_rows->nb[1] / ggml_type_size(set_rows->type);
}
const int n_dims = ((const int32_t *) rope->op_params)[1];
const int mode = ((const int32_t *) rope->op_params)[2];
const int n_ctx_orig = ((const int32_t *) rope->op_params)[4];
float freq_base;
float freq_scale;
float ext_factor;
float attn_factor;
float beta_fast;
float beta_slow;
memcpy(&freq_base, (const int32_t *) rope->op_params + 5, sizeof(float));
memcpy(&freq_scale, (const int32_t *) rope->op_params + 6, sizeof(float));
memcpy(&ext_factor, (const int32_t *) rope->op_params + 7, sizeof(float));
memcpy(&attn_factor, (const int32_t *) rope->op_params + 8, sizeof(float));
memcpy(&beta_fast, (const int32_t *) rope->op_params + 9, sizeof(float));
memcpy(&beta_slow, (const int32_t *) rope->op_params + 10, sizeof(float));
const bool is_neox = mode & GGML_ROPE_TYPE_NEOX;
const int32_t * pos = (const int32_t *) rope->src[1]->data;
const float * freq_factors = rope->src[2] != nullptr ? (const float *) rope->src[2]->data : nullptr;
rope_corr_dims corr_dims;
ggml_rope_yarn_corr_dims(n_dims, n_ctx_orig, freq_base, beta_fast, beta_slow, corr_dims.v);
const size_t ts0 = ggml_type_size(x->type);
GGML_ASSERT(x->nb[0] == ts0);
const int64_t s01 = x->nb[1] / ts0;
const int64_t s02 = x->nb[2] / ts0;
const int64_t s03 = x->nb[3] / ts0;
const size_t ts_mul = ggml_type_size(mul_src->type);
GGML_ASSERT(mul_src->nb[0] == ts_mul);
const int64_t mul_s01 = mul_src->nb[1] / ts_mul;
const int64_t mul_s02 = mul_src->nb[2] / ts_mul;
const int64_t mul_s03 = mul_src->nb[3] / ts_mul;
const size_t ts_dst = ggml_type_size(rope->type);
const int64_t s1 = rope->nb[1] / ts_dst;
const int64_t s2 = rope->nb[2] / ts_dst;
const int64_t s3 = rope->nb[3] / ts_dst;
cudaStream_t stream = ctx.stream();
if (dst_type == GGML_TYPE_F32) {
rms_norm_mul_rope_cuda((const float *) x->data, (float *) dst_d,
x->ne[0], x->ne[1], x->ne[2], x->ne[3], s01, s02, s03, s1, s2, s3, eps,
(const float *) mul_src->data, mul_s01, mul_s02, mul_s03,
mul_src->ne[0], mul_src->ne[1], mul_src->ne[2], mul_src->ne[3],
n_dims, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims,
freq_factors, row_indices, set_rows_stride, is_neox, stream);
} else if (dst_type == GGML_TYPE_F16) {
rms_norm_mul_rope_cuda((const float *) x->data, (half *) dst_d,
x->ne[0], x->ne[1], x->ne[2], x->ne[3], s01, s02, s03, s1, s2, s3, eps,
(const float *) mul_src->data, mul_s01, mul_s02, mul_s03,
mul_src->ne[0], mul_src->ne[1], mul_src->ne[2], mul_src->ne[3],
n_dims, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims,
freq_factors, row_indices, set_rows_stride, is_neox, stream);
} else {
GGML_ABORT("fatal error");
}
}
-2
View File
@@ -7,5 +7,3 @@ void ggml_cuda_op_rope(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * set_rows);
void ggml_cuda_op_rms_norm_mul_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * rms_norm, ggml_tensor * mul, ggml_tensor * rope, ggml_tensor * set_rows);
+1 -55
View File
@@ -141,57 +141,6 @@ static __global__ void rwkv_wkv7_f32(const int B, const int T, const int C, cons
}
}
template <int rows_per_block>
static __global__ void __launch_bounds__(WARP_SIZE * rows_per_block, 2)
rwkv_wkv7_f32_t1_warp_row(const int T, const int C, const int H, const float * r, const float * w, const float * k, const float * v, const float * a, const float * b, const float * s, float * dst) {
constexpr int head_size = CUDA_WKV_BLOCK_SIZE;
constexpr int half_head = head_size / 2;
const int lane = threadIdx.x;
const int row = blockIdx.y * rows_per_block + threadIdx.y;
const int bid = blockIdx.x;
const int batch_i = bid / H;
const int head_i = bid % H;
const int state_size = C * head_size;
const int head_off = head_i * head_size;
const int t = batch_i * C + head_off + row;
__shared__ float _r[head_size], _w[head_size], _k[head_size], _a[head_size], _b[head_size];
if (threadIdx.y == 0) {
_r[lane] = r[batch_i * C + head_off + lane];
_w[lane] = w[batch_i * C + head_off + lane];
_k[lane] = k[batch_i * C + head_off + lane];
_a[lane] = a[batch_i * C + head_off + lane];
_b[lane] = b[batch_i * C + head_off + lane];
_r[lane + half_head] = r[batch_i * C + head_off + lane + half_head];
_w[lane + half_head] = w[batch_i * C + head_off + lane + half_head];
_k[lane + half_head] = k[batch_i * C + head_off + lane + half_head];
_a[lane + half_head] = a[batch_i * C + head_off + lane + half_head];
_b[lane + half_head] = b[batch_i * C + head_off + lane + half_head];
}
__syncthreads();
const int64_t state_base = batch_i * state_size + head_i * head_size * head_size + row * head_size;
const float s0 = s[state_base + lane];
const float s1 = s[state_base + lane + half_head];
const float sa = warp_reduce_sum(_a[lane] * s0 + _a[lane + half_head] * s1);
const float vt = v[t];
const float st0 = s0 * _w[lane] + _k[lane] * vt + sa * _b[lane];
const float st1 = s1 * _w[lane + half_head] + _k[lane + half_head] * vt + sa * _b[lane + half_head];
const float y = warp_reduce_sum(st0 * _r[lane] + st1 * _r[lane + half_head]);
dst[T * C + state_base + lane] = st0;
dst[T * C + state_base + lane + half_head] = st1;
if (lane == 0) {
dst[t] = y;
}
}
void ggml_cuda_op_rwkv_wkv6(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const float * k_d = (const float *)dst->src[0]->data;
const float * v_d = (const float *)dst->src[1]->data;
@@ -242,10 +191,7 @@ void ggml_cuda_op_rwkv_wkv7(ggml_backend_cuda_context & ctx, ggml_tensor * dst)
GGML_ASSERT(C % H == 0);
GGML_ASSERT(C / H == CUDA_WKV_BLOCK_SIZE || C / H == CUDA_WKV_BLOCK_SIZE * 2);
if (T / B == 1 && C / H == CUDA_WKV_BLOCK_SIZE) {
constexpr int rows_per_block = 4;
rwkv_wkv7_f32_t1_warp_row<rows_per_block><<<dim3(B * H, CUDA_WKV_BLOCK_SIZE / rows_per_block), dim3(WARP_SIZE, rows_per_block), 0, stream>>>(T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
} else if (C / H == CUDA_WKV_BLOCK_SIZE) {
if (C / H == CUDA_WKV_BLOCK_SIZE) {
rwkv_wkv7_f32<CUDA_WKV_BLOCK_SIZE><<<B * H, C / H, 0, stream>>>(B, T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
} else {
rwkv_wkv7_f32<CUDA_WKV_BLOCK_SIZE * 2><<<B * H, C / H, 0, stream>>>(B, T, C, H, r_d, w_d, k_d, v_d, a_d, b_d, s_d, dst_d);
-1
View File
@@ -1646,7 +1646,6 @@ static void ggml_backend_et_device_get_props(ggml_backend_dev_t dev, struct ggml
/* .host_buffer = */ false,
/* .buffer_from_host_ptr = */ false,
/* .events = */ false,
/* .mmap_support = */ true,
};
}
-166
View File
@@ -1,166 +0,0 @@
#pragma once
#if defined(__aarch64__) || defined(_M_ARM64)
#if defined(__linux__)
#include <sys/auxv.h>
#include <sys/prctl.h>
#if !defined(HWCAP2_SVE2)
#define HWCAP2_SVE2 (1ULL << 1)
#endif
#if !defined(HWCAP_FPHP)
#define HWCAP_FPHP (1 << 9)
#endif
#if !defined(HWCAP_ASIMDHP)
#define HWCAP_ASIMDHP (1 << 10)
#endif
#if !defined(HWCAP2_I8MM)
#define HWCAP2_I8MM (1ULL << 13)
#endif
#if !defined(HWCAP_ASIMDDP)
#define HWCAP_ASIMDDP (1 << 20)
#endif
#if !defined(HWCAP_SVE)
#define HWCAP_SVE (1 << 22)
#endif
#if !defined(HWCAP2_SME)
#define HWCAP2_SME (1ULL << 23)
#endif
#if !defined(HWCAP2_SME2)
#define HWCAP2_SME2 (1ULL << 37)
#endif
#if !defined(PR_SVE_GET_VL)
#define PR_SVE_GET_VL 51
#endif
#if !defined(PR_SVE_VL_LEN_MASK)
#define PR_SVE_VL_LEN_MASK 0xffff
#endif
#elif defined(__APPLE__)
#include <sys/sysctl.h>
#elif defined(_WIN32)
#include <windows.h>
#if !defined(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE 43
#endif
#if !defined(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_SVE_INSTRUCTIONS_AVAILABLE 46
#endif
#if !defined(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE 47
#endif
#if !defined(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE 66
#endif
#if !defined(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE 67
#endif
#if !defined(PF_ARM_SME_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_SME_INSTRUCTIONS_AVAILABLE 70
#endif
#if !defined(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE)
#define PF_ARM_SME2_INSTRUCTIONS_AVAILABLE 71
#endif
#endif
typedef struct ggml_feats_arch64_runtime {
bool has_dotprod;
bool has_fp16;
bool has_sve;
bool has_sve2;
bool has_i8mm;
bool has_sme;
bool has_sme2;
int sve_cnt;
} ggml_feats_arch64_runtime_t;
static inline ggml_feats_arch64_runtime_t ggml_feats_get_arch64_runtime(void) {
ggml_feats_arch64_runtime_t runtime_feat = {};
#if defined(__linux__)
const unsigned long hwcap = getauxval(AT_HWCAP);
const unsigned long hwcap2 = getauxval(AT_HWCAP2);
runtime_feat.has_dotprod = !!(hwcap & HWCAP_ASIMDDP);
runtime_feat.has_fp16 = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP);;
runtime_feat.has_sve = !!(hwcap & HWCAP_SVE);
runtime_feat.has_sve2 = !!(hwcap2 & HWCAP2_SVE2);
runtime_feat.has_i8mm = !!(hwcap2 & HWCAP2_I8MM);
runtime_feat.has_sme = !!(hwcap2 & HWCAP2_SME);
runtime_feat.has_sme2 = !!(hwcap2 & HWCAP2_SME2);
if (runtime_feat.has_sve) {
const int vl = prctl(PR_SVE_GET_VL);
if (vl >= 0) {
runtime_feat.sve_cnt = vl & PR_SVE_VL_LEN_MASK;
}
}
#elif defined(__APPLE__)
int oldp = 0;
size_t size = sizeof(oldp);
if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_dotprod = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_FP16", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_fp16 = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SVE", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_sve = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SVE2", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_sve2 = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_i8mm = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_sme = static_cast<bool>(oldp);
}
if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, nullptr, 0) == 0) {
runtime_feat.has_sme2 = static_cast<bool>(oldp);
}
// Apple does not support userspace non-streaming SVE; keep SVE vector length unknown.
runtime_feat.sve_cnt = 0;
#elif defined (_WIN32)
runtime_feat.has_dotprod = IsProcessorFeaturePresent(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_fp16 = IsProcessorFeaturePresent(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_sve = IsProcessorFeaturePresent(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_sve2 = IsProcessorFeaturePresent(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_i8mm = IsProcessorFeaturePresent(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_sme = IsProcessorFeaturePresent(PF_ARM_SME_INSTRUCTIONS_AVAILABLE) != 0;
runtime_feat.has_sme2 = IsProcessorFeaturePresent(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE) != 0;
// Windows exposes SVE feature presence, but not the runtime SVE vector length here.
runtime_feat.sve_cnt = 0;
#endif
return runtime_feat;
}
#endif // defined(__aarch64__) || defined(_M_ARM64)
-1
View File
@@ -3930,7 +3930,6 @@ static void ggml_backend_hexagon_device_get_props(ggml_backend_dev_t dev, struct
/* .host_buffer = */ (bool) opt_hostbuf,
/* .buffer_from_host_ptr = */ false,
/* .events = */ false,
/* .mmap_support = */ false,
};
}
+3
View File
@@ -126,6 +126,9 @@ if (GGML_HIP_EXPORT_METRICS)
set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -Rpass-analysis=kernel-resource-usage --save-temps")
endif()
# Fast math for HIP, like CUDA's -use_fast_math. Not -ffast-math: that implies -ffinite-math-only, which breaks ggml's INFINITY masking and produces NaNs.
set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -funsafe-math-optimizations")
if (NOT GGML_CUDA_FA)
add_compile_definitions(GGML_CUDA_NO_FA)
endif()

Some files were not shown because too many files have changed in this diff Show More