Compare commits

..

6 Commits

Author SHA1 Message Date
Xuan Son Nguyen b4e5d8b616 rm pocket-tts 2026-08-17 10:14:14 +02:00
Xuan Son Nguyen 9edb2d2dba BailingMoeV3ForCausalLM 2026-08-17 10:11:26 +02:00
Xuan Son Nguyen 507ace2827 Merge branch 'master' into xsn/convert_add_example 2026-08-17 10:09:07 +02:00
Xuan Son Nguyen 6e97dc0518 add more variants 2026-08-17 09:12:19 +02:00
Xuan Son Nguyen 584b4f93c6 add docs 2026-08-17 02:36:33 +02:00
Xuan Son Nguyen 71c7866a16 convert: add @ModelBase.example 2026-08-17 02:31:10 +02:00
274 changed files with 2514 additions and 5976 deletions
+10 -10
View File
@@ -1,18 +1,18 @@
ARG OPENVINO_VERSION_MAJOR=2026.3
ARG OPENVINO_VERSION_FULL=2026.3.0.22451.bd8d6542e3c
ARG OPENVINO_VERSION_MAJOR=2026.2.1
ARG OPENVINO_VERSION_FULL=2026.2.1.21919.ede283a88e3
ARG UBUNTU_VERSION=24.04
# Intel GPU driver versions. https://github.com/intel/compute-runtime/releases
ARG IGC_VERSION=v2.38.2
ARG IGC_VERSION_FULL=2_2.38.2+22051
ARG COMPUTE_RUNTIME_VERSION=26.27.39122.11
ARG COMPUTE_RUNTIME_VERSION_FULL=26.27.39122.11-0
ARG IGC_VERSION=v2.36.3
ARG IGC_VERSION_FULL=2_2.36.3+21719
ARG COMPUTE_RUNTIME_VERSION=26.22.38646.4
ARG COMPUTE_RUNTIME_VERSION_FULL=26.22.38646.4-0
ARG IGDGMM_VERSION=22.10.0
# Intel NPU driver versions. https://github.com/intel/linux-npu-driver/releases
ARG NPU_DRIVER_VERSION=v1.35.0
ARG NPU_DRIVER_FULL=v1.35.0.20260722-29947505341
ARG LIBZE1_VERSION=1.28.2-1~24.04~ppa1
ARG NPU_DRIVER_VERSION=v1.33.0
ARG NPU_DRIVER_FULL=v1.33.0.20260529-26625960453
ARG LIBZE1_VERSION=1.27.0-1~24.04~ppa2
# Optional proxy build arguments
ARG http_proxy=
@@ -170,7 +170,7 @@ RUN --mount=type=cache,target=/var/cache/intel-npu,sharing=locked \
fi; \
DEB=/var/cache/intel-npu/libze1_${LIBZE1_VERSION}_amd64.deb; \
if [ ! -f "$DEB" ]; then \
wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260606T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
wget -q -O "$DEB" https://snapshot.ppa.launchpadcontent.net/kobuk-team/intel-graphics/ubuntu/20260324T100000Z/pool/main/l/level-zero-loader/libze1_${LIBZE1_VERSION}_amd64.deb; \
fi; \
mkdir /tmp/npu/ && cd /tmp/npu/ && tar -xf "$TGZ" && cp "$DEB" .; \
apt-get update; \
@@ -0,0 +1,20 @@
name: "Linux - Setup Vulkan SDK"
description: "Setup Vulkan SDK for Linux"
inputs:
path:
description: "Installation path"
required: true
version:
description: "Vulkan SDK version"
required: true
runs:
using: "composite"
steps:
- name: Setup Vulkan SDK
id: setup
uses: ./.github/actions/unarchive-tar
with:
url: https://sdk.lunarg.com/sdk/download/${{ inputs.version }}/linux/vulkan_sdk.tar.xz
path: ${{ inputs.path }}
strip: 1
@@ -6,7 +6,8 @@ inputs:
required: true
cuda_arch:
description: "CUDA target architecture"
required: true
required: false
default: "x64"
runs:
using: "composite"
+32 -5
View File
@@ -10,6 +10,33 @@ concurrency:
cancel-in-progress: true
jobs:
ubuntu-24-vulkan-cache:
runs-on: ubuntu-24.04
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
- name: Get latest Vulkan SDK version
id: vulkan_sdk_version
run: |
echo "VULKAN_SDK_VERSION=$(curl https://vulkan.lunarg.com/sdk/latest/linux.txt)" >> "$GITHUB_ENV"
- name: Setup Cache
uses: actions/cache@v5
id: cache-sdk
with:
path: ./vulkan_sdk
key: cache-gha-vulkan-sdk-${{ env.VULKAN_SDK_VERSION }}-${{ runner.os }}
- name: Setup Vulkan SDK
if: steps.cache-sdk.outputs.cache-hit != 'true'
uses: ./.github/actions/linux-setup-vulkan
with:
path: ./vulkan_sdk
version: ${{ env.VULKAN_SDK_VERSION }}
#ubuntu-24-spacemit-cache:
# runs-on: ubuntu-24.04
@@ -40,9 +67,9 @@ jobs:
runs-on: ubuntu-24.04
env:
# Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
# Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Clone
@@ -69,8 +96,8 @@ jobs:
env:
# Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Clone
+13 -1
View File
@@ -21,7 +21,6 @@ on:
paths: [
'.github/workflows/build-cpu.yml',
'.github/workflows/build-cmake-pkg.yml',
'ggml/src/ggml-rpc/**',
'**/CMakeLists.txt',
'**/.cmake',
'**/*.h',
@@ -124,6 +123,7 @@ jobs:
env:
OPENBLAS_VERSION: 0.3.23
SDE_VERSION: 9.33.0-2024-01-07
VULKAN_VERSION: 1.4.357.0
strategy:
matrix:
@@ -134,6 +134,9 @@ jobs:
- build: 'x64-openblas'
arch: 'x64'
defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_OPENMP=OFF -DGGML_BLAS=ON -DGGML_BLAS_VENDOR=OpenBLAS -DBLAS_INCLUDE_DIRS="$env:RUNNER_TEMP/openblas/include" -DBLAS_LIBRARIES="$env:RUNNER_TEMP/openblas/lib/openblas.lib"'
- build: 'x64-vulkan'
arch: 'x64'
defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DCMAKE_BUILD_TYPE=Release -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_VULKAN=ON'
- build: 'arm64'
arch: 'arm64'
defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON'
@@ -164,6 +167,15 @@ jobs:
$lib = $(join-path $msvc 'bin\Hostx64\x64\lib.exe')
& $lib /machine:x64 "/def:${env:RUNNER_TEMP}/openblas/lib/libopenblas.def" "/out:${env:RUNNER_TEMP}/openblas/lib/openblas.lib" /name:openblas.dll
- name: Install Vulkan SDK
id: get_vulkan
if: ${{ matrix.build == 'x64-vulkan' }}
run: |
curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe"
& "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install
Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}"
Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin"
- name: Install Ninja
id: install_ninja
run: |
+13 -19
View File
@@ -22,7 +22,6 @@ env:
jobs:
cuda:
name: windows-cuda (${{ matrix.cuda }}, ${{ matrix.arch }})
runs-on: windows-2022
permissions:
@@ -30,16 +29,7 @@ jobs:
strategy:
matrix:
include:
- cuda: '12.4'
arch: x64
defines: '-DGGML_CUDA_CUB_3DOT2=ON'
- cuda: '13.3'
arch: x64
defines: ''
- cuda: '13.4'
arch: arm64
defines: '-DCMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-msvc-cuda.cmake'
cuda: ['12.4', '13.3']
steps:
- name: Clone
@@ -49,13 +39,12 @@ jobs:
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
key: release-windows-2022-x64-cuda-${{ matrix.cuda }}
- name: Install Cuda Toolkit
uses: ./.github/actions/windows-setup-cuda
with:
cuda_version: ${{ matrix.cuda }}
cuda_arch: ${{ matrix.arch }}
- name: Install Ninja
id: install_ninja
@@ -65,21 +54,26 @@ jobs:
- name: Build
id: cmake_build
shell: cmd
# TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project
run: |
call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" ${{ matrix.arch == 'x64' && 'x64' || 'amd64_arm64' }}
call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64
cmake -S . -B build -G "Ninja Multi-Config" ^
-DGGML_BACKEND_DL=ON ^
-DLLAMA_BUILD_SERVER=ON ^
-DLLAMA_BUILD_BORINGSSL=ON ^
-DGGML_NATIVE=OFF ^
-DGGML_CPU=OFF ^
-DGGML_BACKEND_DL=ON ^
-DGGML_CPU_ALL_VARIANTS=ON ^
-DGGML_CUDA=ON ^
-DLLAMA_BUILD_BORINGSSL=ON ${{ matrix.defines }}
-DGGML_RPC=ON ^
-DGGML_CUDA_CUB_3DOT2=ON
set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1
cmake --build build --config Release -j %NINJA_JOBS% --target ggml-cuda
cmake --build build --config Release -j %NINJA_JOBS% -t ggml
cmake --build build --config Release
- name: ccache-clear
uses: ./.github/actions/ccache-clear
with:
key: release-windows-2022-${{ matrix.arch }}-cuda-${{ matrix.cuda }}
key: release-windows-2022-x64-cuda-${{ matrix.cuda }}
hip:
runs-on: windows-2022
+7 -7
View File
@@ -39,8 +39,8 @@ jobs:
env:
# Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Clone
@@ -81,7 +81,7 @@ jobs:
# TODO: fix and re-enable the `test-llama-archs` test below
run: |
cd ${{ github.workspace }}
ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" --verbose --timeout 2000
ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 2000
- name: Test (GPU)
id: cmake_test_gpu
@@ -89,15 +89,15 @@ jobs:
run: |
cd ${{ github.workspace }}
export GGML_OPENVINO_DEVICE=GPU
ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" --verbose --timeout 3000
ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 3000
openvino-windows-2022:
runs-on: windows-2022
env:
# Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Clone
@@ -166,4 +166,4 @@ jobs:
call "%OPENVINO_ROOT%\setupvars.bat"
cd build
ctest --test-dir ReleaseOV -L main -E "test-llama-archs|test-recurrent-state-rollback-nemotron-h" -C Release --verbose --timeout 3000
ctest --test-dir ReleaseOV -L main -E "test-llama-archs" -C Release --verbose --timeout 3000
+66
View File
@@ -0,0 +1,66 @@
name: CI (rpc)
on:
workflow_dispatch: # allows manual triggering
push:
branches:
- master
paths: [
'.github/workflows/build-rpc.yml',
'**/CMakeLists.txt',
'**/.cmake',
'**/*.h',
'**/*.hpp',
'**/*.c',
'**/*.cpp'
]
pull_request:
types: [opened, synchronize, reopened]
paths: [
'.github/workflows/build-rpc.yml',
'ggml/src/ggml-rpc/**'
]
concurrency:
group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
cancel-in-progress: true
env:
GGML_NLOOP: 3
GGML_N_THREADS: 1
LLAMA_ARG_LOG_COLORS: 1
LLAMA_ARG_LOG_PREFIX: 1
LLAMA_ARG_LOG_TIMESTAMPS: 1
jobs:
ubuntu-24-rpc:
runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }}
continue-on-error: true
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
- name: Dependencies
id: depends
run: |
sudo apt-get update
sudo apt-get install build-essential libssl-dev ninja-build
- name: Build
id: cmake_build
run: |
cmake -B build \
-G "Ninja" \
-DCMAKE_BUILD_TYPE=Release \
-DGGML_RPC=ON
time cmake --build build --config Release -j $(nproc)
- name: Test
id: cmake_test
run: |
cd build
ctest -L main --verbose
+2 -2
View File
@@ -288,8 +288,8 @@ jobs:
env:
# Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Clone
+11 -58
View File
@@ -93,13 +93,19 @@ jobs:
run: |
echo "VULKAN_SDK_VERSION=$(curl https://vulkan.lunarg.com/sdk/latest/linux.txt)" >> "$GITHUB_ENV"
- name: Setup Vulkan SDK
id: setup
uses: ./.github/actions/unarchive-tar
- name: Use Vulkan SDK Cache
uses: actions/cache@v5
id: cache-sdk
with:
url: https://sdk.lunarg.com/sdk/download/${{ env.VULKAN_SDK_VERSION }}/linux/vulkan_sdk.tar.xz
path: ./vulkan_sdk
strip: 1
key: cache-gha-vulkan-sdk-${{ env.VULKAN_SDK_VERSION }}-${{ runner.os }}
- name: Setup Vulkan SDK
if: steps.cache-sdk.outputs.cache-hit != 'true'
uses: ./.github/actions/linux-setup-vulkan
with:
path: ./vulkan_sdk
version: ${{ env.VULKAN_SDK_VERSION }}
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
@@ -127,56 +133,3 @@ jobs:
# This is using llvmpipe and runs slower than other backends
# test-backend-ops is too slow on llvmpipe, skip it
ctest -L main -E test-backend-ops --verbose --timeout 900
windows:
runs-on: windows-2025
env:
VULKAN_VERSION: 1.4.357.0
steps:
- name: Clone
id: checkout
uses: actions/checkout@v6
- name: ccache
uses: ggml-org/ccache-action@v1.2.21
with:
key: cpu-windows-2025-x64-vulkan
variant: ccache
evict-old-files: 1d
save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
- name: Install Vulkan SDK
id: get_vulkan
run: |
curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe"
& "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install
Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}"
Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin"
- name: Install Ninja
id: install_ninja
run: |
choco install ninja
- name: Build
id: cmake_build
run: |
cmake -S . -B build -G "Ninja Multi-Config" `
-D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake `
-DCMAKE_BUILD_TYPE=Release `
-DGGML_NATIVE=OFF `
-DLLAMA_BUILD_SERVER=ON `
-DGGML_RPC=ON `
-DGGML_BACKEND_DL=ON `
-DGGML_CPU_ALL_VARIANTS=ON `
-DGGML_VULKAN=ON `
-DLLAMA_BUILD_BORINGSSL=ON
cmake --build build --config Release -j ${env:NUMBER_OF_PROCESSORS}
- name: Test
id: cmake_test
run: |
cd build
ctest -L main -C Release --verbose --timeout 900
-39
View File
@@ -44,7 +44,6 @@ jobs:
uses: actions/checkout@v6
with:
fetch-depth: 0
ssh-key: ${{ secrets.DEPLOY_KEY_RELEASE }}
- name: Determine source tag name
id: srctag
@@ -394,11 +393,6 @@ jobs:
name: Create shared tags from digests
needs: [prepare_matrices, push_to_registry, create_tag]
runs-on: ubuntu-24.04
permissions:
contents: read
packages: write
id-token: write
attestations: write
strategy:
fail-fast: false
matrix:
@@ -433,7 +427,6 @@ jobs:
password: ${{ secrets.GITHUB_TOKEN }}
- name: Create tags from digests
id: create_tags
shell: bash
run: |
set -euo pipefail
@@ -445,7 +438,6 @@ jobs:
SRC_TAG="${{ needs.create_tag.outputs.source_tag }}"
BUILD_DATE="${{ steps.build_date.outputs.date }}"
COMMIT_SHA="${{ steps.checkout.outputs.commit }}"
echo "image_repo=${IMAGE_REPO}" >> "$GITHUB_OUTPUT"
TAGS="${{ matrix.config.tag }}"
ARCHES="${{ matrix.config.arches }}"
DIGEST_GLOB="/tmp/digests/*.tsv"
@@ -512,16 +504,6 @@ jobs:
echo "Creating ${merged_versioned_tag} from ${refs[*]}"
docker buildx imagetools create "${annotations[@]}" --tag "${merged_versioned_tag}" "${refs[@]}"
if [[ "$tag_name" == "${TAGS%% *}" ]]; then
local digest
digest="$(docker buildx imagetools inspect "${merged_versioned_tag}" --format '{{.Manifest.Digest}}')"
if [[ ! "$digest" =~ ^sha256:[0-9a-f]{64}$ ]]; then
echo "Invalid digest for ${merged_versioned_tag}: ${digest}" >&2
exit 1
fi
echo "${image_type}_digest=${digest}" >> "$GITHUB_OUTPUT"
fi
}
for tag in $TAGS; do
@@ -545,24 +527,3 @@ jobs:
done
env:
GITHUB_REPOSITORY_OWNER: '${{ github.repository_owner }}'
- name: Attest full image
if: ${{ matrix.config.full }}
uses: actions/attest@v4
with:
subject-name: ${{ steps.create_tags.outputs.image_repo }}
subject-digest: ${{ steps.create_tags.outputs.full_digest }}
- name: Attest light image
if: ${{ matrix.config.light }}
uses: actions/attest@v4
with:
subject-name: ${{ steps.create_tags.outputs.image_repo }}
subject-digest: ${{ steps.create_tags.outputs.light_digest }}
- name: Attest server image
if: ${{ matrix.config.server }}
uses: actions/attest@v4
with:
subject-name: ${{ steps.create_tags.outputs.image_repo }}
subject-digest: ${{ steps.create_tags.outputs.server_digest }}
-37
View File
@@ -3,11 +3,6 @@ name: Make Release
on:
workflow_dispatch:
inputs:
commit:
description: 'Commit SHA to release (empty = branch HEAD)'
required: false
default: ''
type: string
dry_run:
description: 'Dry run - validate without creating the tag'
required: true
@@ -27,17 +22,12 @@ jobs:
steps:
- name: Checkout
uses: actions/checkout@v6
with:
ssh-key: ${{ secrets.DEPLOY_KEY_RELEASE }}
ref: ${{ inputs.commit != '' && inputs.commit || github.ref_name }}
fetch-depth: 0
- name: Run release checks
id: checks
run: bash scripts/make-release-checks.sh ${{ github.event.inputs.dry_run == 'true' && '--dry-run' || '' }}
env:
GITHUB_REPOSITORY: ${{ github.repository }}
RELEASE_BRANCH: ${{ github.ref_name }}
- name: Create release tag
if: ${{ github.event.inputs.dry_run == 'false' }}
@@ -49,33 +39,6 @@ jobs:
git push origin "${VERSION}"
echo "Created and pushed tag ${VERSION}"
- name: Generate release description
id: desc
run: bash scripts/make-release-desc.sh "${{ steps.checks.outputs.version }}"
env:
GITHUB_REPOSITORY: ${{ github.repository }}
- name: Create release
if: ${{ github.event.inputs.dry_run == 'false' }}
uses: ggml-org/action-create-release@v1
env:
GITHUB_TOKEN: ${{ github.token }}
with:
tag_name: ${{ steps.checks.outputs.version }}
# TODO: remove the prerelease flag once the semantic versioning workflow is ready
# ref: https://github.com/ggml-org/ggml/discussions/1579
prerelease: true
body: |
> [!NOTE]
> Semantic versioning is still work in progress.
> More info can be found in https://github.com/ggml-org/ggml/discussions/1579
${{ steps.desc.outputs.nightly }}
## ${{ steps.desc.outputs.changelog_title }}
${{ steps.desc.outputs.changelog }}
- name: Dry run summary
if: ${{ github.event.inputs.dry_run == 'true' }}
run: |
+7 -31
View File
@@ -446,8 +446,8 @@ jobs:
env:
# Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Set OpenVINO version output
@@ -562,8 +562,8 @@ jobs:
env:
# Sync versions in build-openvino.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile
OPENVINO_VERSION_MAJOR: "2026.3"
OPENVINO_VERSION_FULL: "2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR: "2026.2.1"
OPENVINO_VERSION_FULL: "2026.2.1.21919.ede283a88e3"
steps:
- name: Set OpenVINO version output
@@ -1439,9 +1439,7 @@ jobs:
- name: xcodebuild for swift package
id: xcodebuild
run: |
# note: only macos and ios-device due to long build time
# ref: https://github.com/ggml-org/llama.cpp/pull/27252
./build-xcframework.sh macos ios-device
./build-xcframework.sh
- name: Build Xcode project
run: xcodebuild -project examples/llama.swiftui/llama.swiftui.xcodeproj -scheme llama.swiftui -sdk iphoneos CODE_SIGNING_REQUIRED=NO CODE_SIGN_IDENTITY= -destination 'generic/platform=iOS' FRAMEWORK_FOLDER_PATH=./build-ios build
@@ -1569,8 +1567,6 @@ jobs:
# https://docs.github.com/en/actions/security-for-github-actions/security-guides/automatic-token-authentication#modifying-the-permissions-for-the-github_token
permissions:
contents: write # for creating release
id-token: write
attestations: write
runs-on: ubuntu-slim
@@ -1579,14 +1575,14 @@ jobs:
- windows
- windows-cpu
- windows-cuda
- windows-sycl
#- windows-sycl
- windows-rocm
- windows-openvino
#- ubuntu-22-rocm
- ubuntu-cpu
- ubuntu-vulkan
- ubuntu-24-openvino
- ubuntu-24-sycl
#- ubuntu-24-sycl
- android-arm64
- macos-cpu
- ios-xcode
@@ -1602,7 +1598,6 @@ jobs:
uses: actions/checkout@v6
with:
fetch-depth: 0
ssh-key: ${{ secrets.DEPLOY_KEY_RELEASE }}
- name: Determine tag name
id: tag
@@ -1664,22 +1659,6 @@ jobs:
run: |
tar -czvf release/llama-${{ steps.tag.outputs.name }}-ui.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./ui-dist .
- name: Attest release artifacts
id: attest
uses: actions/attest@v4
with:
subject-path: 'release/*'
- name: Create and push git tag
run: |
TAG="${{ steps.tag.outputs.name }}"
if git rev-parse -q --verify "refs/tags/${TAG}" >/dev/null 2>&1; then
echo "Tag ${TAG} already exists, skipping creation"
else
git tag "${TAG}"
git push origin "${TAG}"
fi
- name: Create release
id: create_release
uses: ggml-org/action-create-release@v1
@@ -1697,9 +1676,6 @@ jobs:
**Website:**
- <https://llama.app>
**Attestations:**
- <${{ steps.attest.outputs.attestation-url }}>
**macOS/iOS:**
- [macOS Apple Silicon (arm64)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-macos-arm64.tar.gz)
- macOS Apple Silicon (arm64, KleidiAI enabled) [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23780)
-1
View File
@@ -2,7 +2,6 @@ You are a coding agent. Here are some very important rules that you must follow:
General:
- Be very precise and concise when writing code, comments, explanations, etc.
- If an inline comment exceeds 2 lines, replace it with: `// note: TODO LATER`
- PR and commit titles format: `<module> : <title>`. Lookup recents for examples
- Don't try to build or run the code unless you are explicitly asked to do so
- Use the `gh` CLI tool when querying PRs, issues, or other GitHub resources
+1 -462
View File
File diff suppressed because it is too large Load Diff
+2 -3
View File
@@ -5,7 +5,7 @@ include(CheckIncludeFileCXX)
### llama.cpp version
set(LLAMA_VERSION_MAJOR 0)
set(LLAMA_VERSION_MINOR 1)
set(LLAMA_VERSION_PATCH 2)
set(LLAMA_VERSION_PATCH 0)
set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
# whether this is a development/nightly build
@@ -224,10 +224,9 @@ add_subdirectory(src)
# utils, programs, examples and tests
#
add_subdirectory(vendor)
if (LLAMA_BUILD_COMMON)
add_subdirectory(common)
add_subdirectory(vendor/cpp-httplib)
endif()
if (LLAMA_BUILD_COMMON AND LLAMA_BUILD_TESTS AND NOT CMAKE_JS_VERSION)
+6 -7
View File
@@ -7,11 +7,10 @@
<b>LLM inference in C/C++</b>
[![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT)
[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp?filter=v*)](https://github.com/ggml-org/llama.cpp/releases?q=tag:v0)
[![Nightly](https://img.shields.io/github/v/release/ggml-org/llama.cpp?label=nightly)](https://github.com/ggml-org/llama.cpp/releases)
[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp)](https://github.com/ggml-org/llama.cpp/releases)
[![Server](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml)
[![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
[![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
[![Docker](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
[![Winget](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
[manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3A0cc4m%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [compile times](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-compile-times.md) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
@@ -120,7 +119,7 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
## Acknowledgements
- [yhirose/cpp-httplib](https://github.com/yhirose/cpp-httplib) - Single-header HTTP server, used by `llama-server` - MIT license
- [nothings/stb](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
- [stb-image](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
- [nlohmann/json](https://github.com/nlohmann/json) - Single-header JSON library, used by various tools/examples - MIT License
- [mackron/miniaudio](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
- [sheredom/subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
- [miniaudio.h](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
- [subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
+128 -210
View File
@@ -1,8 +1,5 @@
#!/usr/bin/env bash
#
# usage: ./build-xcframework.sh [BUILD ...] (default: all builds)
# builds: ios-sim ios-device macos visionos visionos-sim tvos-sim tvos-device
#
# Options
IOS_MIN_OS_VERSION=16.4
MACOS_MIN_OS_VERSION=13.3
@@ -22,43 +19,6 @@ GGML_METAL_EMBED_LIBRARY=ON
GGML_BLAS_DEFAULT=ON
GGML_OPENMP=OFF
# Max number of concurrent platform builds
MAX_PARALLEL_BUILDS=1
# Split the available cores between the concurrent builds (min 1)
JOBS_PER_BUILD=$(( $(sysctl -n hw.logicalcpu) / MAX_PARALLEL_BUILDS ))
if [[ "$JOBS_PER_BUILD" -lt 1 ]]; then
JOBS_PER_BUILD=1
fi
# echo "build_fn build_dir release_dir platform is_simulator min_os" for a build name
build_spec() {
case "$1" in
ios-sim) echo "build_ios_sim build-ios-sim Release-iphonesimulator ios true ${IOS_MIN_OS_VERSION}" ;;
ios-device) echo "build_ios_device build-ios-device Release-iphoneos ios false ${IOS_MIN_OS_VERSION}" ;;
macos) echo "build_macos build-macos Release macos false ${MACOS_MIN_OS_VERSION}" ;;
visionos) echo "build_visionos build-visionos Release-xros visionos false ${VISIONOS_MIN_OS_VERSION}" ;;
visionos-sim) echo "build_visionos_sim build-visionos-sim Release-xrsimulator visionos true ${VISIONOS_MIN_OS_VERSION}" ;;
tvos-sim) echo "build_tvos_sim build-tvos-sim Release-appletvsimulator tvos true ${TVOS_MIN_OS_VERSION}" ;;
tvos-device) echo "build_tvos_device build-tvos-device Release-appletvos tvos false ${TVOS_MIN_OS_VERSION}" ;;
*) return 1 ;;
esac
}
# Default: build everything
if [[ $# -eq 0 ]]; then
BUILDS=(ios-sim ios-device macos visionos visionos-sim tvos-sim tvos-device)
else
BUILDS=("$@")
fi
for b in "${BUILDS[@]}"; do
if ! build_spec "$b" >/dev/null; then
echo "Error: unknown build '$b'" >&2
echo "Valid builds: ios-sim ios-device macos visionos visionos-sim tvos-sim tvos-device" >&2
exit 1
fi
done
COMMON_C_FLAGS="-Wno-macro-redefined -Wno-shorten-64-to-32 -Wno-unused-command-line-argument -g"
COMMON_CXX_FLAGS="-Wno-macro-redefined -Wno-shorten-64-to-32 -Wno-unused-command-line-argument -g"
@@ -290,7 +250,6 @@ combine_static_libraries() {
"${base_dir}/${build_dir}/ggml/src/ggml-metal/${release_dir}/libggml-metal.a"
"${base_dir}/${build_dir}/ggml/src/ggml-blas/${release_dir}/libggml-blas.a"
"${base_dir}/${build_dir}/tools/mtmd/${release_dir}/libmtmd.a"
"${base_dir}/${build_dir}/vendor/hash/${release_dir}/libvendor-hash.a"
)
# Create temporary directory for processing
@@ -442,189 +401,148 @@ combine_static_libraries() {
rm -rf "${temp_dir}"
}
build_ios_sim() {
echo "Building for iOS simulator..."
cmake -B build-ios-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${IOS_MIN_OS_VERSION} \
-DIOS=ON \
-DCMAKE_SYSTEM_NAME=iOS \
-DCMAKE_OSX_SYSROOT=iphonesimulator \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=iphonesimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-ios-sim --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for iOS simulator..."
cmake -B build-ios-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${IOS_MIN_OS_VERSION} \
-DIOS=ON \
-DCMAKE_SYSTEM_NAME=iOS \
-DCMAKE_OSX_SYSROOT=iphonesimulator \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=iphonesimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-ios-sim --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
build_ios_device() {
echo "Building for iOS devices..."
cmake -B build-ios-device -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${IOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=iOS \
-DCMAKE_OSX_SYSROOT=iphoneos \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=iphoneos \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-ios-device --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for iOS devices..."
cmake -B build-ios-device -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${IOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=iOS \
-DCMAKE_OSX_SYSROOT=iphoneos \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=iphoneos \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-ios-device --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
build_macos() {
echo "Building for macOS..."
cmake -B build-macos -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${MACOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-S .
cmake --build build-macos --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for macOS..."
cmake -B build-macos -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${MACOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-S .
cmake --build build-macos --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
build_visionos() {
echo "Building for visionOS..."
cmake -B build-visionos -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${VISIONOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DCMAKE_SYSTEM_NAME=visionOS \
-DCMAKE_OSX_SYSROOT=xros \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=xros \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_SERVER=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-visionos --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for visionOS..."
cmake -B build-visionos -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${VISIONOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DCMAKE_SYSTEM_NAME=visionOS \
-DCMAKE_OSX_SYSROOT=xros \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=xros \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_SERVER=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-visionos --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
build_visionos_sim() {
echo "Building for visionOS simulator..."
cmake -B build-visionos-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${VISIONOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_SYSTEM_NAME=visionOS \
-DCMAKE_OSX_SYSROOT=xrsimulator \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=xrsimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_SERVER=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-visionos-sim --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for visionOS simulator..."
cmake -B build-visionos-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${VISIONOS_MIN_OS_VERSION} \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DCMAKE_SYSTEM_NAME=visionOS \
-DCMAKE_OSX_SYSROOT=xrsimulator \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=xrsimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DLLAMA_BUILD_SERVER=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-visionos-sim --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
# Add tvOS builds (might need the same u_int definitions as watchOS and visionOS)
build_tvos_sim() {
echo "Building for tvOS simulator..."
cmake -B build-tvos-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${TVOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=tvOS \
-DCMAKE_OSX_SYSROOT=appletvsimulator \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DGGML_METAL=ON \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=appletvsimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-tvos-sim --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
echo "Building for tvOS simulator..."
cmake -B build-tvos-sim -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${TVOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=tvOS \
-DCMAKE_OSX_SYSROOT=appletvsimulator \
-DCMAKE_OSX_ARCHITECTURES="arm64;x86_64" \
-DGGML_METAL=ON \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=appletvsimulator \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-tvos-sim --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
build_tvos_device() {
echo "Building for tvOS devices..."
cmake -B build-tvos-device -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${TVOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=tvOS \
-DCMAKE_OSX_SYSROOT=appletvos \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DGGML_METAL=ON \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=appletvos \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-tvos-device --config Release -j "${JOBS_PER_BUILD}" -- -quiet
}
run_builds_parallel() {
local -a pids=()
local -a names=()
local name i
for name in "$@"; do
# Wait for the oldest running build to free a slot
if [[ "${#pids[@]}" -ge "$MAX_PARALLEL_BUILDS" ]]; then
if ! wait "${pids[0]}"; then
echo "ERROR: build '${names[0]}' failed, log follows (${names[0]}.log):" >&2
kill "${pids[@]}" 2>/dev/null || true
cat "${names[0]}.log" >&2
exit 1
fi
pids=("${pids[@]:1}")
names=("${names[@]:1}")
fi
echo "Starting build: $name (log: ${name}.log, -j ${JOBS_PER_BUILD})"
"$name" > "${name}.log" 2>&1 &
pids+=("$!")
names+=("$name")
done
# Wait for the remaining builds
for i in "${!pids[@]}"; do
if ! wait "${pids[$i]}"; then
echo "ERROR: build '${names[$i]}' failed, log follows (${names[$i]}.log):" >&2
kill "${pids[@]}" 2>/dev/null || true
cat "${names[$i]}.log" >&2
exit 1
fi
done
}
BUILD_FNS=()
for b in "${BUILDS[@]}"; do
read -r fn _ < <(build_spec "$b")
BUILD_FNS+=("$fn")
done
echo "Building: ${BUILDS[*]} (max ${MAX_PARALLEL_BUILDS} at a time, -j ${JOBS_PER_BUILD} each)..."
run_builds_parallel "${BUILD_FNS[@]}"
echo "Building for tvOS devices..."
cmake -B build-tvos-device -G Xcode \
"${COMMON_CMAKE_ARGS[@]}" \
-DCMAKE_OSX_DEPLOYMENT_TARGET=${TVOS_MIN_OS_VERSION} \
-DCMAKE_SYSTEM_NAME=tvOS \
-DCMAKE_OSX_SYSROOT=appletvos \
-DCMAKE_OSX_ARCHITECTURES="arm64" \
-DGGML_METAL=ON \
-DCMAKE_XCODE_ATTRIBUTE_SUPPORTED_PLATFORMS=appletvos \
-DCMAKE_C_FLAGS="${COMMON_C_FLAGS}" \
-DCMAKE_CXX_FLAGS="${COMMON_CXX_FLAGS}" \
-DLLAMA_OPENSSL=OFF \
-DMTMD_VIDEO=OFF \
-S .
cmake --build build-tvos-device --config Release -j $(sysctl -n hw.logicalcpu) -- -quiet
# Setup frameworks and copy binaries and headers
echo "Setting up framework structures..."
for b in "${BUILDS[@]}"; do
read -r _ bdir _ platform _ min_os < <(build_spec "$b")
setup_framework_structure "$bdir" "$min_os" "$platform"
done
setup_framework_structure "build-ios-sim" ${IOS_MIN_OS_VERSION} "ios"
setup_framework_structure "build-ios-device" ${IOS_MIN_OS_VERSION} "ios"
setup_framework_structure "build-macos" ${MACOS_MIN_OS_VERSION} "macos"
setup_framework_structure "build-visionos" ${VISIONOS_MIN_OS_VERSION} "visionos"
setup_framework_structure "build-visionos-sim" ${VISIONOS_MIN_OS_VERSION} "visionos"
setup_framework_structure "build-tvos-sim" ${TVOS_MIN_OS_VERSION} "tvos"
setup_framework_structure "build-tvos-device" ${TVOS_MIN_OS_VERSION} "tvos"
# Create dynamic libraries from static libraries
echo "Creating dynamic libraries from static libraries..."
for b in "${BUILDS[@]}"; do
read -r _ bdir rdir platform is_sim _ < <(build_spec "$b")
combine_static_libraries "$bdir" "$rdir" "$platform" "$is_sim"
done
combine_static_libraries "build-ios-sim" "Release-iphonesimulator" "ios" "true"
combine_static_libraries "build-ios-device" "Release-iphoneos" "ios" "false"
combine_static_libraries "build-macos" "Release" "macos" "false"
combine_static_libraries "build-visionos" "Release-xros" "visionos" "false"
combine_static_libraries "build-visionos-sim" "Release-xrsimulator" "visionos" "true"
combine_static_libraries "build-tvos-sim" "Release-appletvsimulator" "tvos" "true"
combine_static_libraries "build-tvos-device" "Release-appletvos" "tvos" "false"
# Create XCFramework with correct debug symbols paths
echo "Creating XCFramework..."
XCFW_ARGS=()
for b in "${BUILDS[@]}"; do
read -r _ bdir _ _ _ _ < <(build_spec "$b")
XCFW_ARGS+=(-framework "$(pwd)/${bdir}/framework/llama.framework")
XCFW_ARGS+=(-debug-symbols "$(pwd)/${bdir}/dSYMs/llama.dSYM")
done
xcrun xcodebuild -create-xcframework \
"${XCFW_ARGS[@]}" \
-output "$(pwd)/build-apple/llama.xcframework"
-framework $(pwd)/build-ios-sim/framework/llama.framework \
-debug-symbols $(pwd)/build-ios-sim/dSYMs/llama.dSYM \
-framework $(pwd)/build-ios-device/framework/llama.framework \
-debug-symbols $(pwd)/build-ios-device/dSYMs/llama.dSYM \
-framework $(pwd)/build-macos/framework/llama.framework \
-debug-symbols $(pwd)/build-macos/dSYMs/llama.dSYM \
-framework $(pwd)/build-visionos/framework/llama.framework \
-debug-symbols $(pwd)/build-visionos/dSYMs/llama.dSYM \
-framework $(pwd)/build-visionos-sim/framework/llama.framework \
-debug-symbols $(pwd)/build-visionos-sim/dSYMs/llama.dSYM \
-framework $(pwd)/build-tvos-device/framework/llama.framework \
-debug-symbols $(pwd)/build-tvos-device/dSYMs/llama.dSYM \
-framework $(pwd)/build-tvos-sim/framework/llama.framework \
-debug-symbols $(pwd)/build-tvos-sim/dSYMs/llama.dSYM \
-output $(pwd)/build-apple/llama.xcframework
+1 -1
View File
@@ -190,7 +190,7 @@ if [ ! -z ${GG_BUILD_OPENVINO} ]; then
CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_OPENVINO=ON"
# TODO: fix and re-enable the `test-llama-archs` test below
CTEST_EXTRA="-E test-llama-archs|test-recurrent-state-rollback-nemotron-h"
CTEST_EXTRA="-E test-llama-archs"
fi
## helpers
+1 -2
View File
@@ -126,8 +126,7 @@ set_target_properties(${TARGET} PROPERTIES
MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
)
target_include_directories(${TARGET} PUBLIC .)
target_link_libraries (${TARGET} PUBLIC vendor::nlohmann vendor::sheredom)
target_include_directories(${TARGET} PUBLIC . ../vendor)
target_compile_features (${TARGET} PUBLIC cxx_std_17)
if (LLAMA_SUBPROCESS)
+1 -7
View File
@@ -3362,7 +3362,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
{"--tools"}, "TOOL1,TOOL2,...",
"experimental: whether to enable built-in tools for AI agents - do not enable in untrusted environments (default: no tools)\n"
"specify \"all\" to enable all tools\n"
"available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_info\n"
"available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime, get_info\n"
"note: for security reasons, this will limit --cors-origins to localhost by default",
[](common_params & params, const std::string & value) {
params.server_tools = parse_csv_row(value);
@@ -4658,12 +4658,6 @@ void common_params_add_preset_options(std::vector<common_arg> & args) {
[](common_params &, int) { /* unused */ }
).set_env(COMMON_ARG_PRESET_STOP_TIMEOUT).set_preset_only());
args.push_back(common_arg(
{"dedup-cache-models"}, "0|1",
"in server router mode, hide a cached model from the model list when this preset resolves to the same model file",
[](common_params &, const std::string &) { /* unused */ }
).set_env(COMMON_ARG_PRESET_DEDUP_CACHE_MODELS).set_preset_only());
// args.push_back(common_arg(
// {"pin"},
// "in server router mode, do not unload this model if models_max is exceeded",
+2 -3
View File
@@ -11,9 +11,8 @@
#include <memory>
// pseudo-env variable to identify preset-only arguments
#define COMMON_ARG_PRESET_LOAD_ON_STARTUP "__PRESET_LOAD_ON_STARTUP"
#define COMMON_ARG_PRESET_STOP_TIMEOUT "__PRESET_STOP_TIMEOUT"
#define COMMON_ARG_PRESET_DEDUP_CACHE_MODELS "__PRESET_DEDUP_CACHE_MODELS"
#define COMMON_ARG_PRESET_LOAD_ON_STARTUP "__PRESET_LOAD_ON_STARTUP"
#define COMMON_ARG_PRESET_STOP_TIMEOUT "__PRESET_STOP_TIMEOUT"
//
// CLI argument parsing
-2
View File
@@ -1778,8 +1778,6 @@ void common_threadpools::init(llama_context * ctx, const common_params & params)
struct ggml_threadpool_params tpp =
ggml_threadpool_params_from_cpu_params(params.cpuparams);
// each pool needs to match the respective n_threads exactly
// see: https://github.com/ggml-org/llama.cpp/pull/27138#issuecomment-5332307332
if (!ggml_threadpool_params_match(&tpp, &tpp_batch)) {
threadpool_batch = ggml_threadpool_new_fn(&tpp_batch);
if (!threadpool_batch) {
-20
View File
@@ -989,26 +989,6 @@ std::vector<common_cached_model_info> common_list_cached_models() {
return result;
}
std::string common_download_resolve_path(const std::string & hf_repo_with_tag, const std::string & hf_file) {
auto [repo, tag] = common_download_split_repo_tag(hf_repo_with_tag);
auto files = hf_cache::get_cached_files(repo);
if (files.empty()) {
return "";
}
if (!hf_file.empty()) {
for (const auto & f : files) {
if (f.path == hf_file) {
return f.local_path;
}
}
return "";
}
return find_best_model(files, tag).local_path;
}
bool common_download_remove(const std::string & hf_repo_with_tag) {
namespace fs = std::filesystem;
-4
View File
@@ -85,10 +85,6 @@ std::vector<std::string> common_download_get_all_parts(const std::string & url);
// returns list of cached models
std::vector<common_cached_model_info> common_list_cached_models();
// resolve the local cached file path for a HF repo without network access (hf_file, if given, must match exactly)
// returns an empty string if the model is not present in the cache
std::string common_download_resolve_path(const std::string & hf_repo_with_tag, const std::string & hf_file = "");
// download single file from url to local path
// returns status code or -1 on error
// skip_etag: if true, don't read/write .etag files (for HF cache where filename is the hash)
+8 -15
View File
@@ -926,9 +926,6 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
// draft-dspark: the draft carries a Markov head and uses an anchor-first block layout
const bool is_dspark;
// dspark speculators
bool sample_from_anchor = true;
const int32_t * target_layer_ids = nullptr; // model_dft's extract layer indices
uint32_t target_layer_ids_n = 0;
@@ -963,20 +960,16 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
if (llama_model_meta_val_str(model_dft, "dflash.block_size", buf, sizeof(buf)) >= 0) {
block_size = std::atoi(buf);
}
if (llama_model_meta_val_str(model_dft, "dflash.sample_from_anchor", buf, sizeof(buf)) >= 0) {
sample_from_anchor = std::strcmp(buf, "true") == 0;
}
}
mask_token_id = llama_vocab_mask(llama_model_get_vocab(model_dft));
LOG_INF("%s: adding speculative implementation '%s'\n", __func__, common_speculative_type_to_str(type).c_str());
LOG_INF("%s: - n_max=%d, n_min=%d, p_min=%.2f\n", __func__, this->params.n_max, this->params.n_min, this->params.p_min);
LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u, sample_from_anchor=%s\n", __func__,
block_size, mask_token_id, target_layer_ids_n, sample_from_anchor ? "true" : "false");
LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u\n", __func__, block_size, mask_token_id, target_layer_ids_n);
// DFlash input is [id_last, <mask> * (block_size-1)]: in-place denoising yields at most
// block_size-1 draft tokens, anchor-first DSpark yields a full block_size draft tokens
const int32_t n_draft_max = is_dspark && sample_from_anchor ? block_size : block_size - 1;
// block_size-1 draft tokens, DSpark yield a full block_size draft tokens
const int32_t n_draft_max = is_dspark ? block_size : block_size - 1;
if (this->params.n_max > n_draft_max || this->params.n_min > n_draft_max) {
LOG_WRN("%s: requested draft size (n_max=%d, n_min=%d) exceeds the trained block size %d -- clamping to %d\n",
__func__, this->params.n_max, this->params.n_min, block_size, n_draft_max);
@@ -1182,7 +1175,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
const int32_t n_draft = params.n_max;
const int32_t n_block_tokens = n_draft + (is_dspark && sample_from_anchor ? 0 : 1);
const int32_t n_block_tokens = n_draft + (is_dspark ? 0 : 1);
i_block_beg[seq_id] = batch.n_tokens;
n_block [seq_id] = n_block_tokens;
for (int32_t i = 0; i < n_block_tokens; ++i) {
@@ -1215,11 +1208,11 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
auto & result = *dp.result;
if (is_dspark) {
// DSpark: read from the first draft slot, truncate below the confidence threshold
// DSpark predicts the next token from position 0 and optionally truncates
// at the first position below the confidence threshold.
const float * conf = params.p_min > 0.0f ? llama_get_embeddings_nextn(ctx_dft) : nullptr;
// bonus-anchor drafts read the mask positions only, like DFlash
const int32_t i_draft_beg = sample_from_anchor ? 0 : 1;
for (int32_t i = i_draft_beg; i < n_block_tokens; ++i) {
for (int32_t i = 0; i < n_block_tokens; ++i) {
const int32_t idx = beg + i;
if (conf && conf[(size_t) idx * n_embd_dec] < params.p_min) {
-4
View File
@@ -55,8 +55,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"DeepseekV32ForCausalLM": "deepseek",
"DFlashDraftModel": "qwen",
"Qwen3DSparkModel": "qwen",
"DSparkDraftModel": "qwen",
"DSparkSpeculator": "qwen",
"DeepseekV4ForCausalLM": "deepseek",
"DeepseekV4DSparkModel": "deepseek",
"DistilBertForMaskedLM": "bert",
@@ -109,8 +107,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
"GraniteSwitchForCausalLM": "granite",
"GraniteSpeechForConditionalGeneration": "granite",
"GraniteSpeechPlusForConditionalGeneration": "granite",
"GraniteSWAForCausalLM": "granite",
"GraniteMoeSWAForCausalLM": "granite",
"Grok1ForCausalLM": "grok",
"GrokForCausalLM": "grok",
"GroveMoeForCausalLM": "grovemoe",
-102
View File
@@ -74,108 +74,6 @@ class GraniteModel(LlamaModel):
return super().filter_tensors(item)
@ModelBase.register("GraniteSWAForCausalLM")
class GraniteSWAModel(GraniteModel):
"""Conversion for IBM's GraniteSWAForCausalLM (interleaved sliding window attention)"""
model_arch = gguf.MODEL_ARCH.GRANITE_SWA
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if name.endswith("sinks"):
name += ".weight"
return super().filter_tensors((name, gen))
def set_gguf_parameters(self):
"""GraniteSWA uses Granite parameters plus sliding window configuration."""
super().set_gguf_parameters()
# Add sliding_window from config
sliding_window = self.hparams.get("sliding_window", 128)
self.gguf_writer.add_sliding_window(sliding_window)
logger.info("gguf: (granite_swa) sliding_window = %s", sliding_window)
# Derive sliding_window_pattern from layer_types
if layer_types := self.hparams.get("layer_types"):
is_swa = [t == "sliding_attention" for t in layer_types]
self.gguf_writer.add_sliding_window_pattern(is_swa)
logger.info("gguf: (granite_swa) sliding_window_pattern = %d SWA layers / %d total",
sum(is_swa), len(is_swa))
else:
# Fall back to period-based pattern: i % 4 != 0
# This matches the transformers default pattern
n_layers = self.block_count
is_swa = [i % 4 != 0 for i in range(n_layers)]
self.gguf_writer.add_sliding_window_pattern(is_swa)
logger.info("gguf: (granite_swa) sliding_window_pattern (inferred) = %d SWA layers / %d total",
sum(is_swa), n_layers)
# Add rope_pattern from no_rope_layers
if no_rope_layers := self.hparams.get("no_rope_layers"):
# Convert 1/0 to bool (1 = use RoPE, 0 = NoPE)
rope_pattern = [bool(x) for x in no_rope_layers]
self.gguf_writer.add_rope_pattern(rope_pattern)
logger.info("gguf: (granite_swa) rope_pattern = %d RoPE layers / %d total",
sum(rope_pattern), len(rope_pattern))
@ModelBase.register("GraniteMoeSWAForCausalLM")
class GraniteMoeSWAModel(GraniteSWAModel):
"""Conversion for IBM's GraniteMoeSWAForCausalLM (unified dense + MoE with iSWA)"""
model_arch = gguf.MODEL_ARCH.GRANITE_SWA
def set_gguf_parameters(self):
super().set_gguf_parameters()
if shared_intermediate_size := self.hparams.get("shared_intermediate_size"):
self.gguf_writer.add_expert_shared_feed_forward_length(shared_intermediate_size)
logger.info("gguf: (granitemoewa) shared_intermediate_size = %s", shared_intermediate_size)
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
"""Split merged MoE tensors (gate+up) following standard MoE pattern."""
# Handle expert FFN tensors (merged gate+up) - swash format: experts.gate_up_proj
# Kept fused since inference (build_moe_ffn) supports a single gate_up_exps
# tensor for the routed experts.
if name.endswith("block_sparse_moe.experts.gate_up_proj"):
ffn_dim = self.hparams["intermediate_size"]
assert data_torch.shape[-2] == 2 * ffn_dim, f"Merged FFN tensor size must be 2 * intermediate_size, got {data_torch.shape[-2]}"
yield from ModelBase.modify_tensors(self, data_torch, self.format_tensor_name(gguf.MODEL_TENSOR.FFN_GATE_UP_EXP, bid), bid)
return
# Handle expert FFN down projection - swash format: experts.down_proj
if name.endswith("block_sparse_moe.experts.down_proj"):
yield from ModelBase.modify_tensors(self, data_torch, self.format_tensor_name(gguf.MODEL_TENSOR.FFN_DOWN_EXP, bid), bid)
return
# Handle expert FFN tensors (merged gate+up) - standard granite format: input_linear.weight
# Kept fused since inference (build_moe_ffn) supports a single gate_up_exps
# tensor for the routed experts.
if name.endswith("block_sparse_moe.input_linear.weight"):
ffn_dim = self.hparams["intermediate_size"]
assert data_torch.shape[-2] == 2 * ffn_dim, "Merged FFN tensor size must be 2 * intermediate_size"
yield from ModelBase.modify_tensors(self, data_torch, self.format_tensor_name(gguf.MODEL_TENSOR.FFN_GATE_UP_EXP, bid), bid)
return
# Handle shared expert FFN tensors (if present) - kept fused since
# inference (build_ffn) supports a single ffn_up_shexp tensor with
# LLM_FFN_SWIGLU for the shared expert.
if name.endswith("shared_mlp.input_linear.weight"):
ffn_dim = self.hparams.get("shared_intermediate_size", self.hparams["intermediate_size"])
assert data_torch.shape[-2] == 2 * ffn_dim, "Merged FFN tensor size must be 2 * shared_intermediate_size"
yield from ModelBase.modify_tensors(self, data_torch, self.format_tensor_name(gguf.MODEL_TENSOR.FFN_UP_SHEXP, bid), bid)
return
# Handle shared expert output (if present)
if name.endswith("shared_mlp.output_linear.weight"):
yield from ModelBase.modify_tensors(self, data_torch, self.format_tensor_name(gguf.MODEL_TENSOR.FFN_DOWN_SHEXP, bid), bid)
return
# Pass through to parent for all other tensors (including sinks)
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("GraniteMoeForCausalLM", "GraniteMoeSharedForCausalLM")
@ModelBase.example("ibm-granite/granite-3.1-3b-a800m-instruct")
class GraniteMoeModel(GraniteModel):
+12 -73
View File
@@ -4,13 +4,12 @@ import json
from typing import Any, Callable, Iterable, TYPE_CHECKING
import numpy as np
import torch
if TYPE_CHECKING:
from torch import Tensor
from .base import LazyTorchTensor, ModelBase, TextModel, gguf, logger
from .base import ModelBase, TextModel, gguf, logger
@ModelBase.register("QWenLMHeadModel")
@@ -709,82 +708,22 @@ class DFlashModel(Qwen3Model):
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("Qwen3DSparkModel", "DSparkDraftModel", "DSparkSpeculator")
@ModelBase.register("Qwen3DSparkModel")
@ModelBase.example("satgeze/Qwen3.6-27B-DSpark")
class DSparkModel(DFlashModel):
# DSpark = DFlash + a semi-autoregressive Markov head.
# DSpark = DFlash + a semi-autoregressive Markov head
model_arch = gguf.MODEL_ARCH.DFLASH
def __init__(self, dir_model, *args, **kwargs):
hparams = kwargs.pop("hparams", None)
if hparams is None:
hparams = ModelBase.load_hparams(dir_model, False)
# EAGLE3-style exports use the 1+N bonus-anchor block, DFlash-lineage exports sample from the anchor
self._sample_from_anchor = hparams.get(
"sample_from_anchor",
"transformer_layer_config" not in hparams and "aux_hidden_state_layer_ids" not in hparams)
if "transformer_layer_config" in hparams:
hparams = {**hparams, **hparams["transformer_layer_config"]}
super().__init__(dir_model, *args, hparams=hparams, **kwargs)
# normalize both schemas to DFlash's nested dflash_config
if "aux_hidden_state_layer_ids" in self.hparams:
self.hparams.setdefault("dflash_config", {
"mask_token_id": self.hparams.get("mask_token_id"),
"target_layer_ids": [i - 1 for i in self.hparams["aux_hidden_state_layer_ids"]],
})
else:
self.hparams.setdefault("dflash_config", {
k: self.hparams[k] for k in ("target_layer_ids", "mask_token_id") if k in self.hparams
})
if (markov_head_type := self.hparams.get("markov_head_type", "vanilla")) != "vanilla":
raise ValueError(f"unsupported markov_head_type {markov_head_type!r} (only 'vanilla' is supported)")
n_vocab = self.hparams["vocab_size"]
self._n_vocab_draft = self.hparams.get("draft_vocab_size") or n_vocab
if self._n_vocab_draft > n_vocab:
raise ValueError(f"draft_vocab_size {self._n_vocab_draft} exceeds vocab_size {n_vocab}")
self._d2t: Tensor | None = None
def set_gguf_parameters(self):
super().set_gguf_parameters()
self.gguf_writer.add_sample_from_anchor(self._sample_from_anchor)
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
# normalize the flat DeepSpec schema to DFlash's nested dflash_config
self.hparams.setdefault("dflash_config", {
k: self.hparams[k] for k in ("target_layer_ids", "mask_token_id") if k in self.hparams
})
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
if item[0] == "t2d": # not used at runtime
name, gen = item
if name.endswith(("embed_tokens.weight", "lm_head.weight")):
return None
return super().filter_tensors(item)
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
if name == "model.d2t":
self._d2t = data_torch
return
if self._n_vocab_draft == self.hparams["vocab_size"] and name.endswith(("embed_tokens.weight", "lm_head.weight")):
return
yield from super().modify_tensors(data_torch, name, bid)
def prepare_tensors(self):
super().prepare_tensors()
n_vocab = self.hparams["vocab_size"]
if self._n_vocab_draft < n_vocab and self._d2t is None:
raise ValueError(f"draft_vocab_size {self._n_vocab_draft} < vocab_size {n_vocab} but no d2t table found")
# write d2t as absolute target token ids
if self._d2t is not None:
data = LazyTorchTensor.to_eager(self._d2t).to(torch.int64).cpu().numpy().reshape(-1)
if data.size != self._n_vocab_draft:
raise ValueError(f"d2t size {data.size} does not match draft_vocab_size {self._n_vocab_draft}")
data = data + np.arange(data.size, dtype=np.int64)
if np.any((data < 0) | (data >= n_vocab)):
raise ValueError(f"d2t target ids out of range for target vocab size {n_vocab}")
if np.unique(data).size != data.size:
raise ValueError("d2t contains duplicate target ids")
logger.info(f"{'d2t,':<30} --> I64, shape = {{{data.size}}}")
self.gguf_writer.add_tensor("d2t", data, raw_dtype=gguf.GGMLQuantizationType.I64)
return super().filter_tensors((name, gen))
+6 -6
View File
@@ -237,8 +237,8 @@ chmod +x ubuntu-llamacpp-ov-install.sh
# ============================================
set -euo pipefail
OPENVINO_VERSION_MAJOR="2026.3"
OPENVINO_VERSION_FULL="2026.3.0.22451.bd8d6542e3c"
OPENVINO_VERSION_MAJOR="2026.2.1"
OPENVINO_VERSION_FULL="2026.2.1.21919.ede283a88e3"
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
OPENVINO_INSTALL_DIR="/opt/intel/openvino_${OPENVINO_VERSION_MAJOR}"
@@ -334,7 +334,7 @@ echo " ./build/ReleaseOV/bin/llama-cli -m model.gguf"
```
> [!NOTE]
> The script pins OpenVINO `2026.3` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release.
> The script pins OpenVINO `2026.2.1` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release.
</details>
@@ -364,8 +364,8 @@ REM ============================================
REM llama.cpp OpenVINO Build Script (Ninja)
REM ============================================
set "OPENVINO_VERSION_MAJOR=2026.3"
set "OPENVINO_VERSION_FULL=2026.3.0.22451.bd8d6542e3c"
set "OPENVINO_VERSION_MAJOR=2026.2.1"
set "OPENVINO_VERSION_FULL=2026.2.1.21919.ede283a88e3"
set "SCRIPT_DIR=%~dp0"
set "VCPKG_DIR=C:\vcpkg"
@@ -547,7 +547,7 @@ endlocal
```
> [!NOTE]
> The script pins OpenVINO `2026.3` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release. From any new shell, source the matching `setupvars` script via the junction — `call "C:\Intel\openvino\setupvars.bat"` from `cmd`, or `& "C:\Intel\openvino\setupvars.ps1"` from PowerShell. If `winget` cannot register Visual Studio Build Tools on first run, install them once manually and re-run the script from an elevated **Developer Command Prompt for VS 2022**.
> The script pins OpenVINO `2026.2.1` via the `OPENVINO_VERSION_MAJOR` / `OPENVINO_VERSION_FULL` variables at the top — edit them to track a different release. From any new shell, source the matching `setupvars` script via the junction — `call "C:\Intel\openvino\setupvars.bat"` from `cmd`, or `& "C:\Intel\openvino\setupvars.ps1"` from PowerShell. If `winget` cannot register Visual Studio Build Tools on first run, install them once manually and re-run the script from an elevated **Developer Command Prompt for VS 2022**.
</details>
-6
View File
@@ -29,12 +29,6 @@ identify which PRs require a version bump before cutting a release._
Releases are created by running the [make-release](.github/workflows/make-release.yml)
which is a manual workflow.
The workflow runs against the branch selected in the "Run workflow" dialog
(default `master`) and takes an optional `commit` SHA. When a commit is given,
the workflow validates that the commit belongs to the branch and is not older
than 3 days from the branch HEAD, then releases that commit instead of the
branch HEAD.
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.
-4
View File
@@ -106,10 +106,6 @@ acceptance (from the draft's confidence head, if present) falls below `P` (defau
Currently only drafts with a Qwen3 backbone are supported; support for other backbones
(e.g. Gemma4) is planned.
DSpark drafts exported in the [speculators](https://github.com/vllm-project/speculators) format
(for example [`RedHatAI/gemma-4-31B-it-speculator.dspark`](https://huggingface.co/RedHatAI/gemma-4-31B-it-speculator.dspark))
convert the same way.
See:
- #25173
+17 -1
View File
@@ -2,5 +2,21 @@ set(TARGET llama-gguf-hash)
add_executable(${TARGET} gguf-hash.cpp)
install(TARGETS ${TARGET} RUNTIME)
target_link_libraries(${TARGET} PRIVATE vendor::hash ggml ${CMAKE_THREAD_LIBS_INIT})
# clibs dependencies
include_directories(deps/)
add_library(xxhash OBJECT deps/xxhash/xxhash.c deps/xxhash/xxhash.h)
target_link_libraries(${TARGET} PRIVATE xxhash)
add_library(sha1 OBJECT deps/sha1/sha1.c deps/sha1/sha1.h)
target_link_libraries(${TARGET} PRIVATE sha1)
if (NOT MSVC)
# disable warnings in 3rd party code
target_compile_options(sha1 PRIVATE -w)
endif()
add_library(sha256 OBJECT deps/sha256/sha256.c deps/sha256/sha256.h)
target_link_libraries(${TARGET} PRIVATE sha256)
target_link_libraries(${TARGET} PRIVATE ggml ${CMAKE_THREAD_LIBS_INIT})
target_compile_features(${TARGET} PRIVATE cxx_std_17)
@@ -0,0 +1,13 @@
{
"name": "rotate-bits",
"version": "0.1.1",
"repo": "jb55/rotate-bits.h",
"description": "rotate bits",
"keywords": ["rotl", "rotr"],
"src": ["rotate-bits.h"],
"license": "Public Domain",
"development": {
"thlorenz/tap.c": "*"
}
}
@@ -0,0 +1,9 @@
{
"name": "sha1",
"version": "0.0.1",
"repo": "clibs/sha1",
"description": "sha1 hash algorithm",
"keywords": ["sha1", "hash"],
"license": "public domain",
"src": ["sha1.c", "sha1.h"]
}
@@ -25,8 +25,6 @@ A million repetitions of "a"
#include "sha1.h"
namespace vendor_hash {
#define rol(value, bits) (((value) << (bits)) | ((value) >> (32 - (bits))))
@@ -295,5 +293,3 @@ void SHA1(
SHA1Final((unsigned char *)hash_out, &ctx);
}
} // namespace vendor_hash
@@ -9,7 +9,9 @@
#include "stdint.h"
namespace vendor_hash {
#if defined(__cplusplus)
extern "C" {
#endif
typedef struct
{
@@ -43,6 +45,8 @@ void SHA1(
const char *str,
uint32_t len);
} // namespace vendor_hash
#if defined(__cplusplus)
}
#endif
#endif /* SHA1_H */
@@ -0,0 +1,15 @@
{
"name": "sha256",
"version": "0.0.2",
"repo": "jb55/sha256.c",
"description": "sha256 in c",
"keywords": ["sha256", "sha2"],
"src": ["sha256.c", "sha256.h"],
"dependencies": {
"jb55/rotate-bits.h": "0.1.1"
},
"development": {
"thlorenz/tap.c": "*"
}
}
+12
View File
@@ -0,0 +1,12 @@
{
"name": "xxhash",
"version": "0.8.2",
"repo": "Cyan4973/xxhash",
"description": "Extremely fast non-cryptographic hash algorithm",
"keywords": ["xxhash", "hashing"],
"license": "BSD-2-Clause",
"src": [
"xxhash.c",
"xxhash.h"
]
}
+3 -6
View File
@@ -17,17 +17,14 @@
extern "C" {
#endif
#include "hash/xxhash/xxhash.h"
#include "hash/sha256/sha256.h"
#include "xxhash/xxhash.h"
#include "sha1/sha1.h"
#include "sha256/sha256.h"
#ifdef __cplusplus
}
#endif
// sha1 is compiled as C++ and lives in a namespace, see scripts/sync_vendor.py
#include "hash/sha1/sha1.h"
using namespace vendor_hash;
// uuid.uuid5(uuid.NAMESPACE_URL, 'en.wikipedia.org/wiki/Llama.cpp')
#define UUID_NAMESPACE_LLAMA_CPP "ef001206-dadc-5f6d-a15f-3359e577d4e5"
+1 -1
View File
@@ -5,7 +5,7 @@ project("ggml" C CXX ASM)
### GGML Version
set(GGML_VERSION_MAJOR 0)
set(GGML_VERSION_MINOR 20)
set(GGML_VERSION_PATCH 2)
set(GGML_VERSION_PATCH 1)
set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
+1 -1
View File
@@ -7,7 +7,7 @@ extern "C" {
#endif
#define RPC_PROTO_MAJOR_VERSION 5
#define RPC_PROTO_MINOR_VERSION 1
#define RPC_PROTO_MINOR_VERSION 0
#define RPC_PROTO_PATCH_VERSION 0
#ifdef __cplusplus
-8
View File
@@ -1981,14 +1981,6 @@ extern "C" {
float beta_fast,
float beta_slow);
// set the offset dims for RoPE
// a must be GGML_OP_ROPE or GGML_OP_ROPE_BACK
// vision RoPE is not supported
// example: (marking: x = rotated, 0 = unrotated)
// n_embd = 10, n_dims = 4, offset = 2 --> [00xxxx0000]
GGML_API struct ggml_tensor * ggml_rope_set_offset(
struct ggml_tensor * a,
int n_offs);
// clamp
// in-place, returns view(a)
-1
View File
@@ -83,7 +83,6 @@ extern "C" {
GGML_API ggml_backend_buffer_t ggml_backend_multi_buffer_alloc_buffer(ggml_backend_buffer_t * buffers, size_t n_buffers);
GGML_API bool ggml_backend_buffer_is_multi_buffer(ggml_backend_buffer_t buffer);
GGML_API void ggml_backend_multi_buffer_set_usage(ggml_backend_buffer_t buffer, enum ggml_backend_buffer_usage usage);
GGML_API void ggml_backend_meta_buffer_set_usage (ggml_backend_buffer_t buffer, enum ggml_backend_buffer_usage usage);
//
// Backend (meta)
+2 -18
View File
@@ -1118,6 +1118,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
}
static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(const struct ggml_tensor * tensor, bool assume_sync) {
GGML_ASSERT(ggml_backend_buffer_is_meta(tensor->buffer));
ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) tensor->buffer->context;
return ggml_backend_meta_get_split_state(buf_ctx->get_simple_tensor_container(tensor), tensor, assume_sync);
}
@@ -1177,15 +1178,7 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_m
t_ij->flags = tensor->flags;
memcpy(t_ij->op_params, tensor->op_params, sizeof(tensor->op_params));
ggml_set_name(t_ij, tensor->name);
t_ij->buffer = simple_buf;
if (simple_buf) {
// the backend that owns the buffer will set .extra
ggml_backend_buffer_init_tensor(simple_buf, t_ij);
} else {
t_ij->extra = tensor->extra;
}
t_ij->view_src = tensor->view_src;
t_ij->view_offs = tensor->view_offs;
if (t_ij->view_src != nullptr && ggml_backend_buffer_is_meta(t_ij->view_src->buffer)) {
@@ -1216,6 +1209,7 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_m
t_ij->data = (char *) ggml_backend_buffer_get_base(simple_buf)
+ size_t(tensor->data) - size_t(ggml_backend_buffer_get_base(tensor->buffer));
}
t_ij->extra = tensor->extra;
for (int i = 0; i < GGML_MAX_SRC; i++) {
t_ij->src[i] = tensor->src[i];
if (tensor->src[i] == tensor) {
@@ -1508,16 +1502,6 @@ bool ggml_backend_buffer_is_meta(ggml_backend_buffer_t buf) {
return buf != nullptr && buf->iface.free_buffer == ggml_backend_meta_buffer_iface.free_buffer;
}
void ggml_backend_meta_buffer_set_usage(ggml_backend_buffer_t buffer, enum ggml_backend_buffer_usage usage) {
GGML_ASSERT(ggml_backend_buffer_is_meta(buffer));
ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buffer->context;
for (size_t i = 0; i < buf_ctx->bufs.size(); i++) {
if (buf_ctx->bufs[i]) {
ggml_backend_buffer_set_usage(buf_ctx->bufs[i].get(), usage);
}
}
}
static ggml_backend_buffer_t ggml_backend_meta_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) {
const size_t n_simple_bufts = ggml_backend_meta_buft_n_bufts(buft);
-2
View File
@@ -182,8 +182,6 @@ void ggml_backend_buffer_set_usage(ggml_backend_buffer_t buffer, enum ggml_backe
// FIXME: add a generic callback to the buffer interface
if (ggml_backend_buffer_is_multi_buffer(buffer)) {
ggml_backend_multi_buffer_set_usage(buffer, usage);
} else if (ggml_backend_buffer_is_meta(buffer)) {
ggml_backend_meta_buffer_set_usage(buffer, usage);
}
}
-3
View File
@@ -2534,9 +2534,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
}
case GGML_OP_ROPE:
{
if (((const int32_t *) op->op_params)[15] != 0) {
return false; // FIXME: support ggml_rope_set_offset
}
if (op->src[0]->ne[0] > 896) {
return false;
}
+3 -14
View File
@@ -5979,8 +5979,6 @@ static void ggml_compute_forward_rope_flt(
memcpy(&beta_slow, (int32_t *) dst->op_params + 10, sizeof(float));
memcpy(&sections, (int32_t *) dst->op_params + 11, sizeof(int)*4);
const int n_offs = ((int32_t *) dst->op_params)[15];
GGML_TENSOR_UNARY_OP_LOCALS
//printf("ne0: %d, ne1: %d, ne2: %d, ne3: %d\n", ne0, ne1, ne2, ne3);
@@ -5997,10 +5995,6 @@ static void ggml_compute_forward_rope_flt(
GGML_ASSERT(n_dims <= ne0);
GGML_ASSERT(n_dims % 2 == 0);
GGML_ASSERT(n_offs >= 0);
GGML_ASSERT(n_offs % 2 == 0);
GGML_ASSERT(n_offs + n_dims <= ne0);
// rows per thread
const int dr = (nr + nth - 1)/nth;
@@ -6026,7 +6020,6 @@ static void ggml_compute_forward_rope_flt(
if (is_vision) {
GGML_ASSERT(n_dims == ne0/2);
GGML_ASSERT(n_offs == 0);
}
const float * freq_factors = NULL;
@@ -6075,12 +6068,12 @@ static void ggml_compute_forward_rope_flt(
switch (mode) {
case GGML_ROPE_TYPE_NORMAL:
rotate_pairs<T>(n_dims, 1, cache, src + n_offs, dst_data + n_offs, 1);
rotate_pairs<T>(n_dims, 1, cache, src, dst_data, 1);
break;
case GGML_ROPE_TYPE_NEOX:
case GGML_ROPE_TYPE_MROPE:
case GGML_ROPE_TYPE_IMROPE:
rotate_pairs<T>(n_dims, n_dims/2, cache, src + n_offs, dst_data + n_offs);
rotate_pairs<T>(n_dims, n_dims/2, cache, src, dst_data);
break;
case GGML_ROPE_TYPE_VISION:
rotate_pairs<T>(ne0, n_dims, cache, src, dst_data);
@@ -6091,11 +6084,7 @@ static void ggml_compute_forward_rope_flt(
if (!is_vision) {
// fill the remain channels with data from src tensor
for (int64_t i0 = 0; i0 < ne0; i0 += 2) {
if (i0 == n_offs) {
i0 += n_dims - 2; // skip the rotated channels
continue;
}
for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) {
const T * const src = (T *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
T * dst_data = (T *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
+3 -5
View File
@@ -29,15 +29,13 @@ extern "C" {
// FP16 to FP32 conversion
// 16-bit float
// on Arm, we use __fp16, which requires the IEEE fp16 format: implied on
// AArch64, selected by -mfp16-format=ieee on 32 bit Arm, where the compiler
// may otherwise reject the type
// on Arm, we use __fp16
// on x86, we use uint16_t
//
// for old CUDA compilers (<= 11), we use uint16_t: ref https://github.com/ggml-org/llama.cpp/pull/10616
// for MUSA compilers , we use uint16_t: ref https://github.com/ggml-org/llama.cpp/pull/11843
//
#if defined(__ARM_NEON) && defined(__ARM_FP16_FORMAT_IEEE) && !(defined(__CUDACC__) && __CUDACC_VER_MAJOR__ <= 11) && !defined(__MUSACC__)
#if defined(__ARM_NEON) && !(defined(__CUDACC__) && __CUDACC_VER_MAJOR__ <= 11) && !defined(__MUSACC__)
#define GGML_CPU_COMPUTE_FP16_TO_FP32(x) neon_compute_fp16_to_fp32(x)
#define GGML_CPU_COMPUTE_FP32_TO_FP16(x) neon_compute_fp32_to_fp16(x)
@@ -328,7 +326,7 @@ inline static float ggml_lookup_fp16_to_fp32(ggml_fp16_t f) {
#define GGML_F16_VEC_REDUCE GGML_F32Cx4_REDUCE
#endif
#elif defined(__ARM_NEON) && defined(__ARM_FEATURE_FMA) && defined(__ARM_FP16_FORMAT_IEEE)
#elif defined(__ARM_NEON) && defined(__ARM_FEATURE_FMA)
#define GGML_SIMD
+2 -8
View File
@@ -2723,12 +2723,6 @@ static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm
return false;
}
// ggml_rope_set_offset is not yet supported in the fused kernel
const int n_offs = ((const int32_t *) rope->op_params)[15];
if (n_offs != 0) {
return false;
}
return true;
}
@@ -4776,7 +4770,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t *
}
// ref: https://github.com/ggml-org/llama.cpp/pull/17368
#if defined(__linux__) && !defined(GGML_USE_HIP)
#if defined(__linux__)
// Check if this is a UMA (Unified Memory Architecture) system
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device)));
@@ -4796,7 +4790,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t *
GGML_LOG_ERROR("%s: /proc/meminfo reading failed, using cudaMemGetInfo\n", __func__);
}
}
#endif // defined(__linux__) && !defined(GGML_USE_HIP)
#endif // defined(__linux__)
// virtual devices sharing one physical GPU share its memory pool; split it between them
const int share_count = ggml_cuda_physical_device_share_count(ctx->device);
+31 -91
View File
@@ -4,7 +4,6 @@
#include "vecdotq.cuh"
#include <cstdint>
#include <type_traits>
typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
@@ -70,8 +69,7 @@ enum mmvq_parameter_table_id {
MMVQ_PARAMETERS_GCN,
MMVQ_PARAMETERS_RDNA2,
MMVQ_PARAMETERS_RDNA3_0,
MMVQ_PARAMETERS_RDNA4,
MMVQ_PARAMETERS_GB10
MMVQ_PARAMETERS_RDNA4
};
static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
@@ -85,8 +83,6 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
return MMVQ_PARAMETERS_GCN;
#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE
return MMVQ_PARAMETERS_TURING;
#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK
return MMVQ_PARAMETERS_GB10;
#else
return MMVQ_PARAMETERS_GENERIC;
#endif
@@ -108,9 +104,6 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) {
return MMVQ_PARAMETERS_TURING;
}
if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) {
return MMVQ_PARAMETERS_GB10;
}
return MMVQ_PARAMETERS_GENERIC;
}
@@ -358,7 +351,7 @@ static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() {
#endif
}
static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id, bool small_k = false, bool halve_iters = false) {
static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id) {
if (table_id == MMVQ_PARAMETERS_GENERIC) {
switch (ncols_dst) {
case 1:
@@ -461,32 +454,11 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
return 1;
}
}
if (table_id == MMVQ_PARAMETERS_GB10) {
const int generic = calc_nwarps(type, ncols_dst, MMVQ_PARAMETERS_GENERIC);
// Only worth the wider block when it actually retires the K loop in half the trips (Observation)
if (ncols_dst == 1 && !small_k && halve_iters) {
switch (type) {
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_IQ4_NL:
return 2 * generic;
default:
break;
}
}
return generic;
}
return 1;
}
static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int table_id, bool small_k = false, int nwarps = 1) {
if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING || table_id == MMVQ_PARAMETERS_GB10) {
if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING) {
switch (ncols_dst) {
case 1:
return small_k ? nwarps : 1;
@@ -505,8 +477,8 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
return 1;
}
template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false, bool halve_iters = false>
__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id(), small_k, halve_iters)*ggml_cuda_get_physical_warp_size(), 1)
template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
static __global__ void mul_mat_vec_q(
const void * vx_ptr, const void * vy_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -523,7 +495,7 @@ static __global__ void mul_mat_vec_q(
constexpr int qi = ggml_cuda_type_traits<type>::qi;
constexpr int vdr = get_vdr_mmvq(type);
constexpr mmvq_parameter_table_id table_id = get_device_table_id();
constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id);
constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
@@ -801,8 +773,8 @@ static __global__ void mul_mat_vec_q_moe(
template<ggml_type type>
static std::pair<dim3, dim3> calc_launch_params(
const int ncols_dst, const int nrows_x, const int nchannels_dst, const int nsamples_or_ntokens,
const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false, const bool halve_iters = false) {
const int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false) {
const int nwarps = calc_nwarps(type, ncols_dst, table_id);
const int rpb = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
const int64_t nblocks = (nrows_x + rpb - 1) / rpb;
const dim3 block_nums(nblocks, nchannels_dst, nsamples_or_ntokens);
@@ -810,7 +782,7 @@ static std::pair<dim3, dim3> calc_launch_params(
return {block_nums, block_dims};
}
template<ggml_type type, int c_ncols_dst, bool small_k = false, bool halve_iters = false>
template<ggml_type type, int c_ncols_dst, bool small_k = false>
static void mul_mat_vec_q_switch_fusion(
const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -825,7 +797,7 @@ static void mul_mat_vec_q_switch_fusion(
if constexpr (c_ncols_dst == 1) {
if (has_fusion) {
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k, halve_iters>, launch_params,
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k>, launch_params,
vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -836,7 +808,7 @@ static void mul_mat_vec_q_switch_fusion(
GGML_ASSERT(!has_fusion && "fusion only supported for ncols_dst=1");
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k, halve_iters>, launch_params,
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k>, launch_params,
vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -888,18 +860,16 @@ static void mul_mat_vec_q_switch_ncols_dst(
const bool has_ids = ids != nullptr;
// How the K loop divides up at the baseline block width, both decisions below use these.
constexpr int qk = ggml_cuda_type_traits<type>::qk;
constexpr int qi = ggml_cuda_type_traits<type>::qi;
constexpr int vdr = get_vdr_mmvq(type);
const int blocks_per_row_x = ncols_x / qk;
const int blocks_per_iter_1warp = vdr * warp_size / qi;
const auto should_use_small_k = [&](int c_ncols_dst) {
// When K is small, increase rows_per_block to match nwarps so each warp has more work to do
// Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle.
const int nwarps = calc_nwarps(type, c_ncols_dst, table_id);
bool use = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;
constexpr int qk = ggml_cuda_type_traits<type>::qk;
constexpr int qi = ggml_cuda_type_traits<type>::qi;
constexpr int vdr = get_vdr_mmvq(type);
const int blocks_per_row_x = ncols_x / qk;
const int blocks_per_iter_1warp = vdr * warp_size / qi;
const int nwarps = calc_nwarps(type, c_ncols_dst, table_id);
bool use = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;
constexpr std::array<ggml_type, 2> iq_slow_turing = {
GGML_TYPE_IQ3_XXS,
@@ -932,28 +902,6 @@ static void mul_mat_vec_q_switch_ncols_dst(
return use;
};
// Whether doubling nwarps pays off on the ncols_dst == 1 path, where K sets the K loop trip count.
const auto should_halve_iters = [&] {
if (table_id != MMVQ_PARAMETERS_GB10) {
return false;
}
// Expert rows are gathered per token, so a wider block adds reduction work without reuse.
if (has_ids) {
return false;
}
const int blocks_per_iter = calc_nwarps(type, 1, table_id) * blocks_per_iter_1warp;
const int iters = (blocks_per_row_x + blocks_per_iter - 1) / blocks_per_iter;
const int iters_wide = (blocks_per_row_x + blocks_per_iter * 2 - 1) / (blocks_per_iter * 2);
// An odd trip count leaves half the wider block idle for its last iteration, that tail is
// only affordable once the loop is long enough to dilute it to an eighth of the work (observation).
const int idle = iters_wide * 2 - iters;
return idle * 8 <= iters_wide * 2;
};
if (has_ids && ncols_dst > 1) {
// Multi-token MUL_MAT_ID path - dedicated MoE kernel
mul_mat_vec_q_moe_launch<type>(
@@ -966,34 +914,26 @@ static void mul_mat_vec_q_switch_ncols_dst(
switch (ncols_dst) {
case 1: {
// static, else MSVC lambda capture breaks the constexpr uses below
static constexpr int c_ncols_dst = 1;
constexpr int c_ncols_dst = 1;
// Tag types keep the flags compile-time, so __launch_bounds__ matches what is launched.
const auto launch = [&](auto small_k_tag, auto halve_iters_tag) {
constexpr bool c_small_k = decltype(small_k_tag)::value;
// Types the table does not promote would compile a second, identical kernel.
constexpr bool c_promoted =
calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, true) !=
calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, false);
bool use_small_k = should_use_small_k(c_ncols_dst);
constexpr bool c_halve_iters = decltype(halve_iters_tag)::value && c_promoted;
const std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
nsamples_dst, warp_size, table_id, c_small_k, c_halve_iters);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, c_small_k, c_halve_iters>(
if (use_small_k) {
std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
nsamples_dst, warp_size, table_id, true);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, true>(
vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
stream);
};
if (should_use_small_k(c_ncols_dst)) {
launch(std::true_type{}, std::false_type{});
} else if (should_halve_iters()) {
launch(std::false_type{}, std::true_type{});
} else {
launch(std::false_type{}, std::false_type{});
std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
nsamples_dst, warp_size, table_id);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst>(
vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
stream);
}
} break;
case 2: {
+59 -93
View File
@@ -53,7 +53,6 @@ static __global__ void rope_norm(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int32_t * pos,
const float freq_scale,
const float ext_factor,
@@ -62,8 +61,7 @@ static __global__ void rope_norm(const T * x,
const float theta_scale,
const float * freq_factors,
const int64_t * row_indices,
const int set_rows_stride,
const bool inplace) {
const int set_rows_stride) {
const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y);
if (i0 >= ne00) {
@@ -94,24 +92,19 @@ static __global__ void rope_norm(const T * x,
ggml_cuda_memcpy_1<4>(dst + idst, &v);
}
};
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
if (i0 >= n_dims) {
store_coaelsced(x[ix + 0], x[ix + 1]);
return;
}
const int iw = i0 - n_offs; // relative idx
const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
const float theta_base = pos[i2]*powf(theta_scale, iw/2.0f);
const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
const float x0 = x[ix + 0];
const float x1 = x[ix + 1];
@@ -132,7 +125,6 @@ static __global__ void rope_neox(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int32_t * pos,
const float freq_scale,
const float ext_factor,
@@ -141,8 +133,7 @@ static __global__ void rope_neox(const T * x,
const float theta_scale,
const float * freq_factors,
const int64_t * row_indices,
const int set_rows_stride,
const bool inplace) {
const int set_rows_stride) {
ggml_cuda_pdl_lc();
const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y);
@@ -167,33 +158,27 @@ static __global__ void rope_neox(const T * x,
idst += row_indices[i2] * set_rows_stride;
}
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
if (i0 >= n_dims) {
dst[idst + i0 / 2 + 0] = ggml_cuda_cast<D>(x[ix + i0 / 2 + 0]);
dst[idst + i0 / 2 + 1] = ggml_cuda_cast<D>(x[ix + i0 / 2 + 1]);
return;
}
const int iw = i0 - n_offs; // relative idx
const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
const float theta_base = pos[i2]*powf(theta_scale, iw/2.0f);
const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
// idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
const float x0 = x[ix + n_offs/2 + 0];
const float x1 = x[ix + n_offs/2 + n_dims/2];
const float x0 = x[ix + 0];
const float x1 = x[ix + n_dims/2];
dst[idst + n_offs/2 + 0] = ggml_cuda_cast<D>(x0 * cos_theta - x1 * sin_theta);
dst[idst + n_offs/2 + n_dims / 2] = ggml_cuda_cast<D>(x0 * sin_theta + x1 * cos_theta);
dst[idst + 0] = ggml_cuda_cast<D>(x0 * cos_theta - x1 * sin_theta);
dst[idst + n_dims / 2] = ggml_cuda_cast<D>(x0 * sin_theta + x1 * cos_theta);
}
template <bool forward, bool has_ff, typename T>
@@ -209,7 +194,6 @@ static __global__ void rope_multi(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int32_t * pos,
const float freq_scale,
const float ext_factor,
@@ -218,8 +202,7 @@ static __global__ void rope_multi(const T * x,
const float theta_scale,
const float * freq_factors,
const mrope_sections sections,
const bool is_imrope,
const bool inplace) {
const bool is_imrope) {
const int i0 = 2 * (blockDim.y * blockIdx.y + threadIdx.y);
if (i0 >= ne00) {
@@ -236,58 +219,52 @@ static __global__ void rope_multi(const T * x,
const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03;
ggml_cuda_pdl_sync();
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
if (i0 >= n_dims) {
dst[idst + i0/2 + 0] = x[ix + i0/2 + 0];
dst[idst + i0/2 + 1] = x[ix + i0/2 + 1];
return;
}
const int iw = i0 - n_offs; // relative idx
const int sect_dims = sections.v[0] + sections.v[1] + sections.v[2] + sections.v[3];
const int sec_w = sections.v[1] + sections.v[0];
const int sector = (iw / 2) % sect_dims;
const int sector = (i0 / 2) % sect_dims;
float theta_base = 0.0;
if (is_imrope) {
if (sector % 3 == 1 && sector < 3 * sections.v[1]) { // h
theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, i0 / 2.0f);
} else if (sector % 3 == 2 && sector < 3 * sections.v[2]) { // w
theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, i0 / 2.0f);
} else if (sector % 3 == 0 && sector < 3 * sections.v[0]) { // t
theta_base = pos[i2] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2] * powf(theta_scale, i0 / 2.0f);
} else {
theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, i0 / 2.0f);
}
} else {
if (sector < sections.v[0]) {
theta_base = pos[i2] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2] * powf(theta_scale, i0 / 2.0f);
} else if (sector >= sections.v[0] && sector < sec_w) {
theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 1] * powf(theta_scale, i0 / 2.0f);
} else if (sector >= sec_w && sector < sec_w + sections.v[2]) {
theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 2] * powf(theta_scale, i0 / 2.0f);
} else if (sector >= sec_w + sections.v[2]) {
theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, iw / 2.0f);
theta_base = pos[i2 + ne02 * 3] * powf(theta_scale, i0 / 2.0f);
}
}
const float freq_factor = has_ff ? freq_factors[iw/2] : 1.0f;
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta);
rope_yarn<forward>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
// idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
const float x0 = x[ix + n_offs/2 + 0];
const float x1 = x[ix + n_offs/2 + n_dims/2];
const float x0 = x[ix + 0];
const float x1 = x[ix + n_dims/2];
dst[idst + n_offs/2 + 0] = x0*cos_theta - x1*sin_theta;
dst[idst + n_offs/2 + n_dims/2] = x0*sin_theta + x1*cos_theta;
dst[idst + 0] = x0*cos_theta - x1*sin_theta;
dst[idst + n_dims/2] = x0*sin_theta + x1*cos_theta;
}
template <bool forward, bool has_ff, typename T>
@@ -367,7 +344,6 @@ static void rope_norm_cuda(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int nr,
const int32_t * pos,
const float freq_scale,
@@ -378,7 +354,6 @@ static void rope_norm_cuda(const T * x,
const float * freq_factors,
const int64_t * row_indices,
const int set_rows_stride,
const bool inplace,
cudaStream_t stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -389,12 +364,12 @@ static void rope_norm_cuda(const T * x,
if (freq_factors == nullptr) {
rope_norm<forward, false><<<block_nums, block_dims, 0, stream>>>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
} else {
rope_norm<forward, true><<<block_nums, block_dims, 0, stream>>>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
}
}
@@ -411,7 +386,6 @@ static void rope_neox_cuda(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int nr,
const int32_t * pos,
const float freq_scale,
@@ -422,7 +396,6 @@ static void rope_neox_cuda(const T * x,
const float * freq_factors,
const int64_t * row_indices,
const int set_rows_stride,
const bool inplace,
cudaStream_t stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -434,12 +407,12 @@ static void rope_neox_cuda(const T * x,
if (freq_factors == nullptr) {
ggml_cuda_kernel_launch(rope_neox<forward, false, T, D>, launch_params,
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
} else {
ggml_cuda_kernel_launch(rope_neox<forward, true, T, D>, launch_params,
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride);
}
}
@@ -456,7 +429,6 @@ static void rope_multi_cuda(const T * x,
const int s2,
const int s3,
const int n_dims,
const int n_offs,
const int nr,
const int32_t * pos,
const float freq_scale,
@@ -467,7 +439,6 @@ static void rope_multi_cuda(const T * x,
const float * freq_factors,
const mrope_sections sections,
const bool is_imrope,
const bool inplace,
cudaStream_t stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1);
@@ -479,13 +450,13 @@ static void rope_multi_cuda(const T * x,
if (freq_factors == nullptr) {
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
ggml_cuda_kernel_launch(rope_multi<forward, false, T>, launch_params,
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope);
} else {
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
ggml_cuda_kernel_launch(rope_multi<forward, true, T>, launch_params,
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope, inplace);
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, pos, freq_scale, ext_factor,
attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope);
}
}
@@ -581,12 +552,8 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
const int mode = ((int32_t *) dst->op_params)[2];
//const int n_ctx = ((int32_t *) dst->op_params)[3];
const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
const int n_offs = ((int32_t *) dst->op_params)[15];
mrope_sections sections;
// when dst aliases src0, the channels outside the rotated window already hold the correct data
const bool inplace = dst_d == src0->data;
// RoPE alteration for extended context
float freq_base;
float freq_scale;
@@ -614,7 +581,6 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
if (is_vision) {
GGML_ASSERT(n_dims == ne00/2);
GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row
}
const int32_t * pos = (const int32_t *) src1_d;
@@ -631,31 +597,31 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
if (is_neox) {
if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
rope_neox_cuda<forward, float, float>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_neox_cuda<forward, float, half>((const float *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_neox_cuda<forward, half, half>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else {
GGML_ABORT("fatal error");
}
} else if (is_mrope && !is_vision) {
if (src0->type == GGML_TYPE_F32) {
rope_multi_cuda<forward>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02, s03, s1,
s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
corr_dims, freq_factors, sections, is_imrope, inplace, stream);
s2, s3, n_dims, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
corr_dims, freq_factors, sections, is_imrope, stream);
} else if (src0->type == GGML_TYPE_F16) {
rope_multi_cuda<forward>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02, s03, s1,
s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
corr_dims, freq_factors, sections, is_imrope, inplace, stream);
s2, s3, n_dims, nr, pos, freq_scale, freq_base, ext_factor, attn_factor,
corr_dims, freq_factors, sections, is_imrope, stream);
} else {
GGML_ABORT("fatal error");
}
@@ -674,19 +640,19 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx,
} else {
if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
rope_norm_cuda<forward, float, float>((const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_norm_cuda<forward, float, half>((const float *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_norm_cuda<forward, half, half>((const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02,
s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else {
GGML_ABORT("fatal error");
}
+1 -3
View File
@@ -1061,11 +1061,9 @@ static bool ggml_backend_et_device_supports_op(ggml_backend_dev_t dev, const ggm
const bool zero_view_offset = op->src[0]->view_src == nullptr || op->src[0]->view_offs == 0;
const bool has_sections = ggml_get_op_params_i32(op, 11) > 0 || ggml_get_op_params_i32(op, 12) > 0 ||
ggml_get_op_params_i32(op, 13) > 0;
// FIXME: support ggml_rope_set_offset
const bool zero_rot_offset = ggml_get_op_params_i32(op, 15) == 0;
supported =
zero_view_offset && zero_rot_offset && ndims <= 512 &&
zero_view_offset && ndims <= 512 &&
(is_normal || (is_neox && ndims % 16 == 0) || (is_imrope && ndims % 16 == 0 && has_sections));
} else {
supported = false;
-4
View File
@@ -3180,10 +3180,6 @@ static bool ggml_hexagon_supported_argsort(const struct ggml_hexagon_session * s
static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
const int32_t * op_params = &op->op_params[0];
if (op_params[15] != 0) {
return false; // FIXME: support ggml_rope_set_offset
}
int mode = op_params[2];
// n_dims == ne0/2, so the rotation spans the full row
+33 -43
View File
@@ -132,8 +132,8 @@ struct hmx_fa_context {
__fp16 * vtcm_v_tiles[2]; // V tiles (column-major, double-buffered)
__fp16 * vtcm_s_tiles[2]; // S = QK^T [g_br, Bc] (double-buffered)
__fp16 * vtcm_p_tiles[2]; // P = softmax(S) [g_br, Bc]
__fp16 * vtcm_d_tiles[2]; // Diagonal rescale, g_br/32 packed diagonal tiles (double-buffered)
__fp16 * vtcm_d_inv_l; // Diagonal rescale (1/l), same packed layout
__fp16 * vtcm_d_tiles; // Diagonal rescale [g_br, g_br]
__fp16 * vtcm_d_inv_l; // Diagonal rescale (1/l) [g_br, g_br]
HVX_Vector * vtcm_m_vec; // Row max [g_br]
HVX_Vector * vtcm_l_vec; // Row sum [g_br]
HVX_Vector * vtcm_s_rowmax; // Softmax intermediate [g_br]
@@ -782,14 +782,13 @@ static void fa_q_load_thread(unsigned int n, unsigned int i, void * data) {
}
}
// Zero the whole rescale region: vtcm_d_tiles[0], the optional vtcm_d_tiles[1]
// and vtcm_d_inv_l are equal-sized and allocated back to back, so one run covers
// them all. The scatter only ever writes the diagonal, ignore the rest.
// Initialize vtcm_d_tiles and vtcm_d_inv_l to 0
const size_t d_bytes_per_t = hex_align_up(d_tile_bytes / n, 128);
const size_t d_start = i * d_bytes_per_t;
const size_t d_end = hex_smin(d_start + d_bytes_per_t, d_tile_bytes);
if (d_start < d_tile_bytes) {
hvx_splat_u8_a((char *) factx->vtcm_d_tiles[0] + d_start, 0, d_end - d_start);
hvx_splat_u8_a((char *) factx->vtcm_d_tiles + d_start, 0, d_end - d_start);
hvx_splat_u8_a((char *) factx->vtcm_d_inv_l + d_start, 0, d_end - d_start);
}
}
@@ -1433,19 +1432,17 @@ static inline void fa_softmax_impl(
const HVX_VectorPred q_32_mask = Q6_Q_vsetq_R(32 * sizeof(__fp16));
HVX_Vector v_exp_m_diff = exp_m_diff_f16;
__fp16 * const d_tiles_out = factx->vtcm_d_tiles[args->buf_idx];
size_t t0 = r_vec_idx * 2;
if (t0 < args->n_row_tiles) {
const HVX_Vector v_content = v_exp_m_diff;
__fp16 * out_base = d_tiles_out + t0 * HMX_FP16_TILE_N_ELMS;
__fp16 * out_base = factx->vtcm_d_tiles + t0 * (args->n_row_tiles_g_br + 1) * HMX_FP16_TILE_N_ELMS;
Q6_vscatter_QRMVhV(q_32_mask, (size_t) out_base, HMX_FP16_TILE_SIZE - 1, v_offsets, v_content);
}
size_t t1 = r_vec_idx * 2 + 1;
if (t1 < args->n_row_tiles) {
const HVX_Vector v_content = Q6_V_vror_VR(v_exp_m_diff, 64);
__fp16 * out_base = d_tiles_out + t1 * HMX_FP16_TILE_N_ELMS;
__fp16 * out_base = factx->vtcm_d_tiles + t1 * (args->n_row_tiles_g_br + 1) * HMX_FP16_TILE_N_ELMS;
Q6_vscatter_QRMVhV(q_32_mask, (size_t) out_base, HMX_FP16_TILE_SIZE - 1, v_offsets, v_content);
}
}
@@ -1509,7 +1506,7 @@ static __attribute__((noinline)) void fa_build_d_diag_inv_l(struct hmx_fa_contex
v_content = Q6_V_vror_VR(v_content, 64);
}
__fp16 * out_base = factx->vtcm_d_inv_l + i * HMX_FP16_TILE_N_ELMS;
__fp16 * out_base = factx->vtcm_d_inv_l + i * (n_row_tiles_g_br + 1) * HMX_FP16_TILE_N_ELMS;
Q6_vscatter_QRMVhV(q_32_mask, (size_t) out_base, HMX_FP16_TILE_SIZE - 1, v_offsets, v_content);
}
}
@@ -1618,7 +1615,7 @@ static void hmx_fa_o_update_worker(void * data) {
const size_t o_stride = n_row_tiles_g_br * HMX_FP16_TILE_N_ELMS;
const size_t v_stride = n_tiles_per_bc * HMX_FP16_TILE_N_ELMS;
for (size_t r = 0; r < n_row_tiles; ++r) {
const __fp16 * d_diag = d_tiles + r * HMX_FP16_TILE_N_ELMS;
const __fp16 * d_diag = d_tiles + r * (n_row_tiles_g_br + 1) * HMX_FP16_TILE_N_ELMS;
const __fp16 * p_tile_in = p_tiles + (r * n_tiles_per_bc) * HMX_FP16_TILE_N_ELMS;
const __fp16 * o_rc = o_prev + r * HMX_FP16_TILE_N_ELMS;
const __fp16 * v_tile_in = v_tiles;
@@ -1657,7 +1654,7 @@ static void hmx_fa_o_norm_worker(void * data) {
asm volatile(HMX_SET_BIAS("%0") :: "r"((unsigned int)job->hmx_scales));
const size_t o_stride = n_row_tiles_g_br * HMX_FP16_TILE_N_ELMS;
for (size_t r = 0; r < n_row_tiles; ++r) {
const __fp16 * d_diag = d_tiles + r * HMX_FP16_TILE_N_ELMS;
const __fp16 * d_diag = d_tiles + r * (n_row_tiles_g_br + 1) * HMX_FP16_TILE_N_ELMS;
const __fp16 * o_rc = o_prev + r * HMX_FP16_TILE_N_ELMS;
__fp16 * o_out = o_curr + r * DV_tiles * HMX_FP16_TILE_N_ELMS;
@@ -1885,8 +1882,7 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
factx.vtcm_s_tiles[1] = VTCM_LAYOUT_PTR_OPTIONAL(__fp16, base, L.off_s_tiles[1], pipeline);
factx.vtcm_p_tiles[0] = VTCM_LAYOUT_PTR(__fp16, base, L.off_p_tiles[0]);
factx.vtcm_p_tiles[1] = VTCM_LAYOUT_PTR_OPTIONAL(__fp16, base, L.off_p_tiles[1], pipeline);
factx.vtcm_d_tiles[0] = VTCM_LAYOUT_PTR(__fp16, base, L.off_d_tiles[0]);
factx.vtcm_d_tiles[1] = VTCM_LAYOUT_PTR_OPTIONAL(__fp16, base, L.off_d_tiles[1], pipeline);
factx.vtcm_d_tiles = VTCM_LAYOUT_PTR(__fp16, base, L.off_d_tiles);
factx.vtcm_d_inv_l = VTCM_LAYOUT_PTR(__fp16, base, L.off_d_inv_l);
factx.vtcm_m_vec = VTCM_LAYOUT_PTR(HVX_Vector, base, L.off_m_vec);
factx.vtcm_l_vec = VTCM_LAYOUT_PTR(HVX_Vector, base, L.off_l_vec);
@@ -2043,30 +2039,7 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
}
}
// ---- 3. Start HMX O update for block kv_blk - 1 (reads P[1 - buf_idx], V[1 - buf_idx], D) ----
// O update relys on the previous block's P and V tiles.
// O update MUST be pushed before the next block's QK-dot: hmx_queue_pop() retires the
// oldest descriptor, so push order alone decides which pop waits for which job.
// If OU went in after QK(i+1), the pop below would retire QK(i+1) and leave
// OU(i-1) in flight into the next iteration, where V-prep overwrites V[prev_buf].
if (kv_blk > 0) {
const size_t prev_buf = 1 - buf_idx;
ou_job[prev_buf].o_curr = o_tile_curr;
ou_job[prev_buf].o_prev = o_tile_prev;
ou_job[prev_buf].p_tiles = factx.vtcm_p_tiles[prev_buf];
ou_job[prev_buf].v_tiles = factx.vtcm_v_tiles[prev_buf];
ou_job[prev_buf].d_tiles = factx.vtcm_d_tiles[prev_buf];
ou_job[prev_buf].hmx_scales = factx.vtcm_hmx_scales_id;
ou_job[prev_buf].n_row_tiles = n_row_tiles;
ou_job[prev_buf].n_col_tiles =
hmx_ceil_div(hex_smin(Bc, nek1 - (kv_blk - 1) * Bc), HMX_FP16_TILE_N_COLS);
ou_job[prev_buf].n_row_tiles_g_br = n_row_tiles_g_br;
ou_job[prev_buf].n_tiles_per_bc = n_tiles_per_bc;
ou_job[prev_buf].DV = DV;
hmx_queue_push(hmx_q, hmx_queue_make_desc(hmx_fa_o_update_worker, &ou_job[prev_buf]));
}
// ---- 4. Pop and run K-prep for next block & push next QK-dot ----
// ---- 3. Pop and run K-prep for next block & push next QK-dot ----
if (kv_blk + 1 < factx.n_kv_blocks) {
const uint32_t next_start = (kv_blk + 1) * Bc;
const uint32_t next_rows = hex_smin(Bc, nek1 - next_start);
@@ -2086,10 +2059,10 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
hmx_queue_push(hmx_q, hmx_queue_make_desc(hmx_fa_qk_dot_worker, &qk_job[next_buf]));
}
// ---- 5. Wait for current block's QK-dot to finish ----
// ---- 4. Wait for current block's QK-dot to finish ----
hmx_queue_pop(hmx_q);
// ---- 6. Phase 2: softmax + build_D ----
// ---- 5. Phase 2: softmax + build_D ----
fa_softmax_args_t sargs;
memset(&sargs, 0, sizeof(sargs));
sargs.factx = &factx;
@@ -2112,6 +2085,23 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
sargs.mask_vtcm_row_stride = factx.mask_buf_row_stride;
sargs.slopes = factx.vtcm_slopes;
// Start HMX O update for block kv_blk - 1 (reads P[1 - buf_idx], V[1 - buf_idx])
if (kv_blk > 0) {
const size_t prev_buf = 1 - buf_idx;
ou_job[prev_buf].o_curr = o_tile_curr;
ou_job[prev_buf].o_prev = o_tile_prev;
ou_job[prev_buf].p_tiles = factx.vtcm_p_tiles[prev_buf];
ou_job[prev_buf].v_tiles = factx.vtcm_v_tiles[prev_buf];
ou_job[prev_buf].d_tiles = factx.vtcm_d_tiles;
ou_job[prev_buf].hmx_scales = factx.vtcm_hmx_scales_id;
ou_job[prev_buf].n_row_tiles = n_row_tiles;
ou_job[prev_buf].n_col_tiles = hmx_ceil_div(hex_smin(Bc, nek1 - (kv_blk - 1) * Bc), HMX_FP16_TILE_N_COLS);
ou_job[prev_buf].n_row_tiles_g_br = n_row_tiles_g_br;
ou_job[prev_buf].n_tiles_per_bc = n_tiles_per_bc;
ou_job[prev_buf].DV = DV;
hmx_queue_push(hmx_q, hmx_queue_make_desc(hmx_fa_o_update_worker, &ou_job[prev_buf]));
}
// Run Softmax on HVX (blocking call)
fa_phase_softmax_and_build_d(&factx, &sargs, n_row_tiles, n_row_tiles_g_br);
@@ -2138,7 +2128,7 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
ou_job[0].o_prev = o_tile_prev;
ou_job[0].p_tiles = factx.vtcm_p_tiles[1 - buf_idx];
ou_job[0].v_tiles = factx.vtcm_v_tiles[1 - buf_idx];
ou_job[0].d_tiles = factx.vtcm_d_tiles[1 - buf_idx];
ou_job[0].d_tiles = factx.vtcm_d_tiles;
ou_job[0].hmx_scales = factx.vtcm_hmx_scales_id;
ou_job[0].n_row_tiles = n_row_tiles;
ou_job[0].n_col_tiles = last_cols;
@@ -2242,7 +2232,7 @@ int hmx_flash_attn_ext(struct htp_ops_context * octx) {
ou_job.o_prev = o_tile_prev;
ou_job.p_tiles = factx.vtcm_p_tiles[0];
ou_job.v_tiles = factx.vtcm_v_tiles[0];
ou_job.d_tiles = factx.vtcm_d_tiles[0];
ou_job.d_tiles = factx.vtcm_d_tiles;
ou_job.hmx_scales = factx.vtcm_hmx_scales_id;
ou_job.n_row_tiles = n_row_tiles;
ou_job.n_col_tiles = n_col_tiles;
+5 -14
View File
@@ -109,7 +109,7 @@ struct hmx_fa_vtcm_layout {
size_t off_v_tiles[2];
size_t off_s_tiles[2];
size_t off_p_tiles[2];
size_t off_d_tiles[2];
size_t off_d_tiles;
size_t off_d_inv_l;
size_t off_m_vec;
size_t off_l_vec;
@@ -125,7 +125,7 @@ struct hmx_fa_vtcm_layout {
size_t q_tile_bytes;
size_t o_tile_bytes;
size_t s_tile_bytes; // S and P tiles (same size)
size_t d_tile_bytes; // d_tiles[0..1] + d_inv_l, allocated back to back
size_t d_tile_bytes;
size_t m_line_bytes; // one mask row
size_t m_buf_slot_bytes; // one dma_cache slot = align_up(Br * m_line_bytes, 4096)
size_t col_vec_bytes;
@@ -149,12 +149,7 @@ static inline void hmx_fa_vtcm_layout_build(struct hmx_fa_vtcm_layout * L,
const size_t k_tile_size = hex_align_up(Bc * DK * sizeof(__fp16), HTP_FA_HMX_TILE_SIZE);
const size_t v_tile_size = hex_align_up(Bc * DV * sizeof(__fp16), HTP_FA_HMX_TILE_SIZE);
const size_t s_tile_size = hex_align_up(g_br * Bc * sizeof(__fp16), HTP_FA_HMX_TILE_SIZE);
// The rescale matrices are diagonal: the HMX kernels only ever load the g_br/32
// tiles that sit on the diagonal, so store just those, packed back to back with
// a stride of one tile. The old [g_br, g_br] square layout allocated g_br/32
// times more than it used, which is also why a second D buffer was unaffordable.
const size_t d_tile_size = (g_br / HMX_FP16_TILE_N_ROWS) * HTP_FA_HMX_TILE_SIZE;
const size_t d_tile_size = hex_align_up(g_br * g_br * sizeof(__fp16), HTP_FA_HMX_TILE_SIZE);
const size_t q_dma_size = hex_align_up(g_br * DK * (is_q_fp32 ? sizeof(float) : sizeof(__fp16)), 128);
const size_t k_dma_size = hex_align_up(Bc * hex_round_up(DK * sizeof(__fp16), 128), 128);
@@ -172,8 +167,7 @@ static inline void hmx_fa_vtcm_layout_build(struct hmx_fa_vtcm_layout * L,
VTCM_LAYOUT_ALLOC(off, off_q_tiles, q_tile_size);
VTCM_LAYOUT_ALLOC(off, off_o_tiles[0], o_tile_size);
VTCM_LAYOUT_ALLOC(off, off_o_tiles[1], o_tile_size);
VTCM_LAYOUT_ALLOC(off, off_d_tiles[0], d_tile_size);
VTCM_LAYOUT_ALLOC_OPTIONAL(off, off_d_tiles[1], d_tile_size, pipeline);
VTCM_LAYOUT_ALLOC(off, off_d_tiles, d_tile_size);
VTCM_LAYOUT_ALLOC(off, off_d_inv_l, d_tile_size);
// Group B & C share start offset (Group B tiles must be 2KB aligned)
@@ -219,10 +213,7 @@ static inline void hmx_fa_vtcm_layout_build(struct hmx_fa_vtcm_layout * L,
L->o_tile_bytes = o_tile_size;
L->col_vec_bytes = col_vec_size;
L->s_tile_bytes = s_tile_size;
// Measured from the actual offsets rather than assumed to be N * d_tile_size, so
// that inserting a region between them (or adding padding to VTCM_LAYOUT_ALLOC)
// cannot silently leave the tail of the run unzeroed.
L->d_tile_bytes = (L->off_d_inv_l + d_tile_size) - L->off_d_tiles[0];
L->d_tile_bytes = d_tile_size;
L->m_line_bytes = m_line_size;
L->m_buf_slot_bytes = m_buf_slot;
L->row_buf_stride = row_vec_size / 128;
-2
View File
@@ -329,7 +329,6 @@ typedef struct {
uint64_t nb3;
int32_t n_past;
int32_t n_dims;
int32_t n_offs;
int32_t n_ctx_orig;
float freq_base;
float freq_scale;
@@ -342,7 +341,6 @@ typedef struct {
int32_t sect_2;
int32_t sect_3;
bool src2;
bool inplace;
} ggml_metal_kargs_rope;
typedef struct {
-7
View File
@@ -3884,11 +3884,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
const int sect_2 = ((const int32_t *) op->op_params)[13];
const int sect_3 = ((const int32_t *) op->op_params)[14];
const int n_offs = ((const int32_t *) op->op_params)[15];
// when dst aliases src0, the channels outside the rotated window already hold the correct data
const bool inplace = op->data == op->src[0]->data;
ggml_metal_kargs_rope args = {
/*.ne00 =*/ ne00,
/*.ne01 =*/ ne01,
@@ -3908,7 +3903,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
/*.nb3 =*/ nb3,
/*.n_past =*/ n_past,
/*.n_dims =*/ n_dims,
/*.n_offs =*/ n_offs,
/*.n_ctx_orig =*/ n_ctx_orig,
/*.freq_base =*/ freq_base,
/*.freq_scale =*/ freq_scale,
@@ -3921,7 +3915,6 @@ int ggml_metal_op_rope(ggml_metal_op_t ctx, int idx) {
/* sect_2 =*/ sect_2,
/* sect_3 =*/ sect_3,
/* src2 =*/ op->src[2] != nullptr,
/* inplace =*/ inplace,
};
auto pipeline = ggml_metal_library_get_pipeline_rope(lib, op);
+23 -36
View File
@@ -656,13 +656,13 @@ void dequantize_q5_1_t4(device const block_q5_1 * xb, short il, thread type4 & r
template <typename type4x4>
void dequantize_q8_0(device const block_q8_0 *xb, short il, thread type4x4 & reg) {
device const packed_char4 * qs = (device const packed_char4 *) xb->qs;
device const int8_t * qs = ((device const int8_t *)xb->qs);
const float d = xb->d;
float4x4 reg_f;
for (int i = 0; i < 4; ++i) {
reg_f[i] = float4(qs[4*il + i]) * d;
for (int i = 0; i < 16; i++) {
reg_f[i/4][i%4] = (qs[i + 16*il] * d);
}
reg = (type4x4) reg_f;
@@ -670,10 +670,12 @@ void dequantize_q8_0(device const block_q8_0 *xb, short il, thread type4x4 & reg
template <typename type4>
void dequantize_q8_0_t4(device const block_q8_0 *xb, short il, thread type4 & reg) {
device const packed_char4 * qs = (device const packed_char4 *) xb->qs;
device const int8_t * qs = ((device const int8_t *)xb->qs);
const float d = xb->d;
reg = (type4) (float4(qs[il]) * d);
for (int i = 0; i < 4; i++) {
reg[i] = (qs[4*(il%4) + i + 16*(il/4)] * d);
}
}
template <typename type4x4>
@@ -4686,15 +4688,14 @@ kernel void kernel_rope_norm(
float sin_theta;
for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
const int iw = i0 - args.n_offs; // relative idx
const int ic = iw/2;
if (i0 < args.n_dims) {
const int ic = i0/2;
const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
@@ -4705,10 +4706,6 @@ kernel void kernel_rope_norm(
dst_data[0] = x0*cos_theta - x1*sin_theta;
dst_data[1] = x0*sin_theta + x1*cos_theta;
} else {
if (args.inplace) {
continue;
}
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
@@ -4744,18 +4741,17 @@ kernel void kernel_rope_neox(
float sin_theta;
for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
const int iw = i0 - args.n_offs; // relative idx
const int ic = iw/2;
if (i0 < args.n_dims) {
const int ic = i0/2;
const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + (args.n_offs + ic)*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + (args.n_offs + ic)*args.nb0);
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + ic*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + ic*args.nb0);
const float x0 = src[0];
const float x1 = src[args.n_dims/2];
@@ -4763,10 +4759,6 @@ kernel void kernel_rope_neox(
dst_data[0] = x0*cos_theta - x1*sin_theta;
dst_data[args.n_dims/2] = x0*sin_theta + x1*cos_theta;
} else {
if (args.inplace) {
continue;
}
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
@@ -4801,9 +4793,8 @@ kernel void kernel_rope_multi(
float sin_theta;
for (int i0 = 2*tiitg; i0 < args.ne0; i0 += 2*tptg.x) {
if (i0 >= args.n_offs && i0 < args.n_offs + args.n_dims) {
const int iw = i0 - args.n_offs; // relative idx
const int ic = iw/2;
if (i0 < args.n_dims) {
const int ic = i0/2;
// mrope theta calculations
// note: the rest is the same as kernel_rope_neox
@@ -4836,14 +4827,14 @@ kernel void kernel_rope_multi(
}
// end of mrope
const float theta = theta_base * pow(args.freq_base, inv_ndims*iw);
const float theta = theta_base * pow(args.freq_base, inv_ndims*i0);
const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, iw, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + (args.n_offs + ic)*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + (args.n_offs + ic)*args.nb0);
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + ic*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + ic*args.nb0);
const float x0 = src[0];
const float x1 = src[args.n_dims/2];
@@ -4851,10 +4842,6 @@ kernel void kernel_rope_multi(
dst_data[0] = x0*cos_theta - x1*sin_theta;
dst_data[args.n_dims/2] = x0*sin_theta + x1*cos_theta;
} else {
if (args.inplace) {
continue;
}
device const T * const src = (device T *)(src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01 + i0*args.nb00);
device T * dst_data = (device T *)( dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
-1
View File
@@ -202,7 +202,6 @@ set(GGML_OPENCL_KERNELS
sqr
sqrt
ssm_conv
ssm_scan
gated_delta_net
sub
sum_rows
-146
View File
@@ -866,9 +866,6 @@ struct ggml_backend_opencl_context {
// [size_idx][kda][tgpp] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1.
// tgpp 0 = TG variant (COLS_PER_LANE_GROUP=1), tgpp 1 = prefill variant (COLS_PER_LANE_GROUP=4).
cl_kernel kernel_gated_delta_net_f32[4][2][2] = {};
cl_kernel kernel_ssm_scan_f32_mamba2_d128 = nullptr;
cl_kernel kernel_ssm_scan_f32_mamba2_d256 = nullptr;
cl_kernel kernel_timestep_embedding;
cl_kernel kernel_gemv_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns_bin;
cl_kernel kernel_gemm_moe_q8_0_f32_ns;
@@ -3157,24 +3154,6 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
// ssm_scan (Mamba-2 fused per-token recurrent step; d_state in {128, 256})
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "ssm_scan.cl.h"
};
#else
const std::string kernel_src = read_file("ssm_scan.cl");
#endif
cl_program prog =
build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_ssm_scan_f32_mamba2_d128 = clCreateKernel(prog, "kernel_ssm_scan_f32_mamba2_d128", &err), err));
CL_CHECK((backend_ctx->kernel_ssm_scan_f32_mamba2_d256 = clCreateKernel(prog, "kernel_ssm_scan_f32_mamba2_d256", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// gated_delta_net: one kernel per (S_V, KDA, tgpp) triple.
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -7322,23 +7301,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
(op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
case GGML_OP_SSM_CONV:
return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
case GGML_OP_SSM_SCAN: {
// Mamba-2 fused per-token scan. Requires src3->ne[0] == 1 (scalar
// A per head); d_state in {128, 256}; all sources f32. Falls back
// to CPU otherwise (incl. Mamba-1 element-wise A).
for (int i = 0; i < 6; ++i) {
if (op->src[i]->type != GGML_TYPE_F32) {
return false;
}
}
if (op->type != GGML_TYPE_F32) {
return false;
}
const int K = ggml_get_op_params_i32(op, 0);
const int d_state = (int) op->src[0]->ne[0];
const bool is_mamba2 = (op->src[3]->ne[0] == 1);
return is_mamba2 && (d_state == 128 || d_state == 256) && (K == 1);
}
case GGML_OP_GATED_DELTA_NET:
{
// Match the Vulkan backend: only F32 -> F32, S_v in {16, 32, 64, 128}.
@@ -7414,9 +7376,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
case GGML_OP_DIAG_MASK_INF:
return op->ne[3] == 1;
case GGML_OP_ROPE: {
if (((const int32_t *) op->op_params)[15] != 0) {
return false; // FIXME: support ggml_rope_set_offset
}
const int mode = ((const int32_t *) op->op_params)[2];
const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE;
const bool is_vision = mode == GGML_ROPE_TYPE_VISION;
@@ -12298,103 +12257,6 @@ static void ggml_cl_mean(ggml_backend_t backend, const ggml_tensor * src0, const
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
}
static void ggml_cl_ssm_scan(ggml_backend_t backend, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0]; // s
const ggml_tensor * src1 = dst->src[1]; // x
const ggml_tensor * src2 = dst->src[2]; // dt
const ggml_tensor * src3 = dst->src[3]; // A
const ggml_tensor * src4 = dst->src[4]; // B
const ggml_tensor * src5 = dst->src[5]; // C
const ggml_tensor * src6 = dst->src[6]; // ids
GGML_ASSERT(src0 && src1 && src2 && src3 && src4 && src5 && src6 && dst);
ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
ggml_tensor_extra_cl * e0 = (ggml_tensor_extra_cl *) src0->extra;
ggml_tensor_extra_cl * e1 = (ggml_tensor_extra_cl *) src1->extra;
ggml_tensor_extra_cl * e2 = (ggml_tensor_extra_cl *) src2->extra;
ggml_tensor_extra_cl * e3 = (ggml_tensor_extra_cl *) src3->extra;
ggml_tensor_extra_cl * e4 = (ggml_tensor_extra_cl *) src4->extra;
ggml_tensor_extra_cl * e5 = (ggml_tensor_extra_cl *) src5->extra;
ggml_tensor_extra_cl * e6 = (ggml_tensor_extra_cl *) src6->extra;
ggml_tensor_extra_cl * ed = (ggml_tensor_extra_cl *) dst->extra;
cl_ulong o0 = e0->offset + src0->view_offs;
cl_ulong o1 = e1->offset + src1->view_offs;
cl_ulong o2 = e2->offset + src2->view_offs;
cl_ulong o3 = e3->offset + src3->view_offs;
cl_ulong o4 = e4->offset + src4->view_offs;
cl_ulong o5 = e5->offset + src5->view_offs;
cl_ulong o6 = e6->offset + src6->view_offs;
cl_ulong od = ed->offset + dst->view_offs;
const int d_state = (int) src0->ne[0];
const int head_dim = (int) src0->ne[1];
const int n_head = (int) src1->ne[1];
const int n_group = (int) src4->ne[1];
const int n_tokens = (int) src1->ne[2];
const int n_seqs = (int) src1->ne[3];
// Mirror CPU ref: s_off = ggml_nelements(src1) * sizeof(float)
const cl_ulong s_off_bytes = (cl_ulong) ggml_nelements(src1) * sizeof(float);
cl_kernel kernel = (d_state == 128)
? backend_ctx->kernel_ssm_scan_f32_mamba2_d128
: backend_ctx->kernel_ssm_scan_f32_mamba2_d256;
GGML_ASSERT(kernel != nullptr);
cl_ulong s0_nb2 = src0->nb[2];
cl_ulong s0_nb3 = src0->nb[3];
cl_ulong x_nb2 = src1->nb[2];
cl_ulong x_nb3 = src1->nb[3];
cl_ulong dt_nb1 = src2->nb[1];
cl_ulong dt_nb2 = src2->nb[2];
cl_ulong A_nb1 = src3->nb[1];
cl_ulong B_nb2 = src4->nb[2];
cl_ulong B_nb3 = src4->nb[3];
cl_ulong C_nb2 = src5->nb[2];
cl_ulong C_nb3 = src5->nb[3];
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &e0->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &o0));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &e1->data_device));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &o1));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &e2->data_device));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &o2));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &e3->data_device));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &o3));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), &e4->data_device));
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &o4));
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), &e5->data_device));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &o5));
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), &e6->data_device));
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &o6));
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), &ed->data_device));
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &od));
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &s0_nb2));
CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_ulong), &s0_nb3));
CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &x_nb2));
CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &x_nb3));
CL_CHECK(clSetKernelArg(kernel, 20, sizeof(cl_ulong), &dt_nb1));
CL_CHECK(clSetKernelArg(kernel, 21, sizeof(cl_ulong), &dt_nb2));
CL_CHECK(clSetKernelArg(kernel, 22, sizeof(cl_ulong), &A_nb1));
CL_CHECK(clSetKernelArg(kernel, 23, sizeof(cl_ulong), &B_nb2));
CL_CHECK(clSetKernelArg(kernel, 24, sizeof(cl_ulong), &B_nb3));
CL_CHECK(clSetKernelArg(kernel, 25, sizeof(cl_ulong), &C_nb2));
CL_CHECK(clSetKernelArg(kernel, 26, sizeof(cl_ulong), &C_nb3));
CL_CHECK(clSetKernelArg(kernel, 27, sizeof(cl_ulong), &s_off_bytes));
CL_CHECK(clSetKernelArg(kernel, 28, sizeof(int), &head_dim));
CL_CHECK(clSetKernelArg(kernel, 29, sizeof(int), &n_head));
CL_CHECK(clSetKernelArg(kernel, 30, sizeof(int), &n_group));
CL_CHECK(clSetKernelArg(kernel, 31, sizeof(int), &n_tokens));
size_t global_work_size[] = { (size_t)n_head * head_dim * 64, (size_t)n_seqs, 1 };
size_t local_work_size[] = { 64, 1, 1 };
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
}
static void ggml_cl_ssm_conv(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
GGML_ASSERT(src0->extra);
@@ -24881,14 +24743,6 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
}
func = ggml_cl_ssm_conv;
break;
case GGML_OP_SSM_SCAN:
if (!any_on_device) {
return false;
}
// SSM_SCAN has 7 source tensors, so it cannot use the standard
// (src0, src1, dst) func signature. Dispatch directly and return.
ggml_cl_ssm_scan(backend, tensor);
return true;
case GGML_OP_GATED_DELTA_NET:
if (!any_on_device) {
return false;
@@ -118,17 +118,6 @@ __kernel void flash_attn_f16(
__local DATA_TYPE4 l_v[BLOCK_N][DV_VEC];
for (int k_start = 0; k_start < n_kv; k_start += BLOCK_N) {
#if WG_SIZE > FA_SG
// WAR on l_k/l_v: a thread that finishes the compute below early either
// it skipped it (my_query_row >= n_q, the continue) or its subgroup simply
// ran ahead wraps around and reloads the tiles while another subgroup is
// still reading them. Any WG that is exactly one lockstep subgroup
// (WG_SIZE == FA_SG) cannot diverge and hides this; a WG spanning multiple
// subgroups (Intel sg=32, or BLOCK_M > 64 on Adreno) corrupts the result.
// All threads reach this each iteration (no-op on the first), so it does
// not diverge with the continue. Compiled out when WG == one subgroup.
barrier(CLK_LOCAL_MEM_FENCE);
#endif
for (int i = tid; i < BLOCK_N * DK_VEC; i += WG_SIZE) {
const int row = i / DK_VEC;
const int col = i % DK_VEC;
@@ -119,15 +119,13 @@ __kernel void flash_attn_f32(
__local DATA_TYPE4 l_v[BLOCK_N][DV_VEC];
for (int k_start = 0; k_start < n_kv; k_start += BLOCK_N) {
#if WG_SIZE > FA_SG
// WAR on l_k/l_v: a thread that finishes the compute below early either
// it skipped it (my_query_row >= n_q, the continue) or its subgroup simply
// ran ahead wraps around and reloads the tiles while another subgroup is
// still reading them. Any WG that is exactly one lockstep subgroup
// (WG_SIZE == FA_SG) cannot diverge and hides this; a WG spanning multiple
// subgroups (Intel sg=32, or BLOCK_M > 64 on Adreno) corrupts the result.
// All threads reach this each iteration (no-op on the first), so it does
// not diverge with the continue. Compiled out when WG == one subgroup.
#if FA_SG < 64
// WAR on l_k/l_v: threads with my_query_row >= n_q skip the compute below
// (continue) and would race ahead to reload the tiles while active threads
// still read them. A single 64-wide Adreno subgroup (WG == sg) runs lockstep
// and hides this; a WG that spans multiple narrower subgroups (Intel sg=32)
// corrupts the result. All threads reach this each iteration (no-op on the
// first), so it does not diverge with the continue. Compiled out at sg=64.
barrier(CLK_LOCAL_MEM_FENCE);
#endif
for (int i = tid; i < BLOCK_N * DK_VEC; i += WG_SIZE) {
-216
View File
@@ -1,216 +0,0 @@
// Mamba2 fused SSM scan kernel. One workgroup per (head, dim, seq); WG size =
// 64 threads. Each thread owns c_factor = d_state/64 state elements in
// private registers; the state stays resident across the n_tokens t-loop
//
// References:
// ggml/src/ggml-cuda/ssm-scan.cu:117 ssm_scan_f32_group
// ggml/src/ggml-cpu/ops.cpp:9368 ggml_compute_forward_ssm_scan_f32
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifdef cl_khr_subgroups
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#endif
#if defined(cl_qcom_reqd_sub_group_size)
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
#else
#define REQD_SUBGROUP_SIZE_64
#endif
inline float softplus_f32(float x) {
return (x <= 20.0f) ? log(1.0f + exp(x)) : x;
}
// d_state = 128 (most Mamba-2 models, e.g. mamba2-2.7B, Codestral-Mamba).
// WG = 64 threads, each holds 2 state elements (tid and tid+64).
REQD_SUBGROUP_SIZE_64
kernel void kernel_ssm_scan_f32_mamba2_d128(
global const char * src0_base, ulong src0_off,
global const char * src1_base, ulong src1_off,
global const char * src2_base, ulong src2_off,
global const char * src3_base, ulong src3_off,
global const char * src4_base, ulong src4_off,
global const char * src5_base, ulong src5_off,
global const char * src6_base, ulong src6_off,
global char * dst_base, ulong dst_off,
ulong s0_nb2, ulong s0_nb3,
ulong x_nb2, ulong x_nb3,
ulong dt_nb1, ulong dt_nb2,
ulong A_nb1,
ulong B_nb2, ulong B_nb3,
ulong C_nb2, ulong C_nb3,
ulong s_off_bytes,
int head_dim, int n_head, int n_group, int n_tokens
) {
const int d_state = 128;
const int tid = (int) get_local_id(0);
const int wg_x = (int) get_group_id(0);
const int seq_id = (int) get_group_id(1);
const int head_id = wg_x / head_dim;
const int dim_id = wg_x - head_id * head_dim;
const int g = head_id / (n_head / n_group);
src0_base += src0_off;
src1_base += src1_off;
src2_base += src2_off;
src3_base += src3_off;
src4_base += src4_off;
src5_base += src5_off;
src6_base += src6_off;
dst_base += dst_off;
const int seq_slot = ((global const int *) src6_base)[seq_id];
const ulong state_base_off = (ulong)seq_slot * s0_nb3 + (ulong)head_id * s0_nb2
+ (ulong)dim_id * d_state * sizeof(float);
global const float * s0_warp = (global const float *)(src0_base + state_base_off);
const ulong state_out_off = (ulong)seq_id * s0_nb3 + (ulong)head_id * s0_nb2
+ (ulong)dim_id * d_state * sizeof(float);
global float * s_warp = (global float *)(dst_base + s_off_bytes + state_out_off);
global const char * x_seq = src1_base + (ulong)seq_id * x_nb3;
global const char * dt_seq = src2_base + (ulong)seq_id * dt_nb2;
global const char * B_seq = src4_base + (ulong)seq_id * B_nb3 + (ulong)g * d_state * sizeof(float);
global const char * C_seq = src5_base + (ulong)seq_id * C_nb3 + (ulong)g * d_state * sizeof(float);
const ulong y_dim_total = (ulong)n_head * head_dim;
global float * y_seq = (global float *)dst_base
+ (ulong)seq_id * (ulong)n_tokens * y_dim_total;
const float A_val = ((global const float *)src3_base)[(ulong)head_id * A_nb1 / sizeof(float)];
// c_factor = 2: each thread owns 2 state elements (tid and tid+64).
float state0 = s0_warp[tid];
float state1 = s0_warp[tid + 64];
for (int t = 0; t < n_tokens; ++t) {
const float dt_h = ((global const float *)(dt_seq + (ulong)t * dt_nb1))[head_id];
const float dt_softplus = softplus_f32(dt_h);
const float dA = exp(dt_softplus * A_val);
const float x_val = ((global const float *)(x_seq + (ulong)t * x_nb2))[(ulong)head_id * head_dim + dim_id];
const float x_dt = x_val * dt_softplus;
const float B0 = ((global const float *)(B_seq + (ulong)t * B_nb2))[tid];
const float B1 = ((global const float *)(B_seq + (ulong)t * B_nb2))[tid + 64];
const float C0 = ((global const float *)(C_seq + (ulong)t * C_nb2))[tid];
const float C1 = ((global const float *)(C_seq + (ulong)t * C_nb2))[tid + 64];
state0 = state0 * dA + B0 * x_dt;
state1 = state1 * dA + B1 * x_dt;
const float partial = state0 * C0 + state1 * C1;
const float sum = sub_group_reduce_add(partial);
if (tid == 0) {
y_seq[(ulong)t * y_dim_total + (ulong)head_id * head_dim + dim_id] = sum;
}
}
s_warp[tid] = state0;
s_warp[tid + 64] = state1;
}
// d_state = 256 (Falcon-H1). WG = 64 threads, each holds 4 state elements.
REQD_SUBGROUP_SIZE_64
kernel void kernel_ssm_scan_f32_mamba2_d256(
global const char * src0_base, ulong src0_off,
global const char * src1_base, ulong src1_off,
global const char * src2_base, ulong src2_off,
global const char * src3_base, ulong src3_off,
global const char * src4_base, ulong src4_off,
global const char * src5_base, ulong src5_off,
global const char * src6_base, ulong src6_off,
global char * dst_base, ulong dst_off,
ulong s0_nb2, ulong s0_nb3,
ulong x_nb2, ulong x_nb3,
ulong dt_nb1, ulong dt_nb2,
ulong A_nb1,
ulong B_nb2, ulong B_nb3,
ulong C_nb2, ulong C_nb3,
ulong s_off_bytes,
int head_dim, int n_head, int n_group, int n_tokens
) {
const int d_state = 256;
const int tid = (int) get_local_id(0);
const int wg_x = (int) get_group_id(0);
const int seq_id = (int) get_group_id(1);
const int head_id = wg_x / head_dim;
const int dim_id = wg_x - head_id * head_dim;
const int g = head_id / (n_head / n_group);
src0_base += src0_off;
src1_base += src1_off;
src2_base += src2_off;
src3_base += src3_off;
src4_base += src4_off;
src5_base += src5_off;
src6_base += src6_off;
dst_base += dst_off;
const int seq_slot = ((global const int *) src6_base)[seq_id];
const ulong state_base_off = (ulong)seq_slot * s0_nb3 + (ulong)head_id * s0_nb2
+ (ulong)dim_id * d_state * sizeof(float);
global const float * s0_warp = (global const float *)(src0_base + state_base_off);
const ulong state_out_off = (ulong)seq_id * s0_nb3 + (ulong)head_id * s0_nb2
+ (ulong)dim_id * d_state * sizeof(float);
global float * s_warp = (global float *)(dst_base + s_off_bytes + state_out_off);
global const char * x_seq = src1_base + (ulong)seq_id * x_nb3;
global const char * dt_seq = src2_base + (ulong)seq_id * dt_nb2;
global const char * B_seq = src4_base + (ulong)seq_id * B_nb3 + (ulong)g * d_state * sizeof(float);
global const char * C_seq = src5_base + (ulong)seq_id * C_nb3 + (ulong)g * d_state * sizeof(float);
const ulong y_dim_total = (ulong)n_head * head_dim;
global float * y_seq = (global float *)dst_base
+ (ulong)seq_id * (ulong)n_tokens * y_dim_total;
const float A_val = ((global const float *)src3_base)[(ulong)head_id * A_nb1 / sizeof(float)];
// c_factor = 4: each thread owns 4 state elements.
float state0 = s0_warp[tid];
float state1 = s0_warp[tid + 64];
float state2 = s0_warp[tid + 128];
float state3 = s0_warp[tid + 192];
for (int t = 0; t < n_tokens; ++t) {
const float dt_h = ((global const float *)(dt_seq + (ulong)t * dt_nb1))[head_id];
const float dt_softplus = softplus_f32(dt_h);
const float dA = exp(dt_softplus * A_val);
const float x_val = ((global const float *)(x_seq + (ulong)t * x_nb2))[(ulong)head_id * head_dim + dim_id];
const float x_dt = x_val * dt_softplus;
global const float * B_t = (global const float *)(B_seq + (ulong)t * B_nb2);
global const float * C_t = (global const float *)(C_seq + (ulong)t * C_nb2);
const float B0 = B_t[tid];
const float B1 = B_t[tid + 64];
const float B2 = B_t[tid + 128];
const float B3 = B_t[tid + 192];
const float C0 = C_t[tid];
const float C1 = C_t[tid + 64];
const float C2 = C_t[tid + 128];
const float C3 = C_t[tid + 192];
state0 = state0 * dA + B0 * x_dt;
state1 = state1 * dA + B1 * x_dt;
state2 = state2 * dA + B2 * x_dt;
state3 = state3 * dA + B3 * x_dt;
const float partial = state0 * C0 + state1 * C1 + state2 * C2 + state3 * C3;
const float sum = sub_group_reduce_add(partial);
if (tid == 0) {
y_seq[(ulong)t * y_dim_total + (ulong)head_id * head_dim + dim_id] = sum;
}
}
s_warp[tid] = state0;
s_warp[tid + 64] = state1;
s_warp[tid + 128] = state2;
s_warp[tid + 192] = state3;
}
-4
View File
@@ -1227,10 +1227,6 @@ static bool is_op_unsupported_case(const ggml_tensor * op) {
const int32_t * op_params = op->op_params;
const int n_dims = op_params[1];
const int mode = op_params[2];
if (op_params[15] != 0) {
// FIXME: support ggml_rope_set_offset
return true;
}
if (mode != GGML_ROPE_TYPE_NORMAL && mode != GGML_ROPE_TYPE_NEOX && mode != GGML_ROPE_TYPE_IMROPE) {
// GGML_LOG_WARN("OpenVINO backend does not support ROPE with mode %d\n", mode);
return true;
+7 -16
View File
@@ -47,7 +47,7 @@ struct rpc_tensor {
uint64_t data;
char name[GGML_MAX_NAME];
int32_t use_count;
char padding[4];
};
static_assert(sizeof(rpc_tensor) % 8 == 0, "rpc_tensor size must be multiple of 8");
@@ -447,7 +447,7 @@ static rpc_tensor serialize_tensor(const ggml_tensor * tensor) {
// Avoid sending uninitialized data over the wire
memset(result.name, 0, sizeof(result.name));
result.use_count = 0;
memset(result.padding, 0, sizeof(result.padding));
snprintf(result.name, GGML_MAX_NAME, "%s", tensor->name);
return result;
@@ -675,7 +675,7 @@ static void ggml_backend_rpc_synchronize(ggml_backend_t backend) {
// this is no-op because we don't have any async operations
}
static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
static void add_tensor(ggml_tensor * tensor, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
if (tensor == nullptr) {
return;
}
@@ -684,15 +684,10 @@ static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::ve
}
visited.insert(tensor);
for (int i = 0; i < GGML_MAX_SRC; i++) {
add_tensor(tensor->src[i], cgraph, tensors, visited);
add_tensor(tensor->src[i], tensors, visited);
}
add_tensor(tensor->view_src, cgraph, tensors, visited);
rpc_tensor result = serialize_tensor(tensor);
const size_t hash_pos = ggml_hash_find(&cgraph->visited_hash_set, tensor);
if (hash_pos != GGML_HASHSET_FULL && ggml_bitset_get(cgraph->visited_hash_set.used, hash_pos)) {
result.use_count = cgraph->use_counts[hash_pos];
}
tensors.push_back(result);
add_tensor(tensor->view_src, tensors, visited);
tensors.push_back(serialize_tensor(tensor));
}
static void serialize_graph(uint32_t device, const ggml_cgraph * cgraph, std::vector<uint8_t> & output) {
@@ -700,7 +695,7 @@ static void serialize_graph(uint32_t device, const ggml_cgraph * cgraph, std::ve
std::vector<rpc_tensor> tensors;
std::unordered_set<ggml_tensor*> visited;
for (uint32_t i = 0; i < n_nodes; i++) {
add_tensor(cgraph->nodes[i], cgraph, tensors, visited);
add_tensor(cgraph->nodes[i], tensors, visited);
}
// serialization format:
// | device (4 bytes) | n_nodes (4 bytes) | nodes (n_nodes * sizeof(uint64_t) | n_tensors (4 bytes) | tensors (n_tensors * sizeof(rpc_tensor)) |
@@ -1456,10 +1451,6 @@ bool rpc_server::graph_compute(const std::vector<uint8_t> & input) {
GGML_LOG_ERROR("[%s] failed to create graph node %d (id=%" PRId64 ")\n", __func__, i, id);
return false;
}
if (graph->nodes[i] != nullptr) {
const size_t hash_pos = ggml_hash_insert(&graph->visited_hash_set, graph->nodes[i]);
graph->use_counts[hash_pos] = tensor_ptrs.at(id)->use_count;
}
}
ggml_status status = ggml_backend_graph_compute(backends[device], graph);
GGML_ASSERT(status == GGML_STATUS_SUCCESS && "Unsuccessful graph computations are not supported with RPC");
-119
View File
@@ -1,119 +0,0 @@
#include "fwht.hpp"
#include <cmath>
template <int N>
static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows,
const float scale, const sycl::nd_item<2> & item) {
const sycl::sub_group sg = item.get_sub_group();
const int64_t r = item.get_global_id(0);
if (r >= n_rows) {
return;
}
src += r * N;
dst += r * N;
constexpr int el_w = N / WARP_SIZE;
static_assert(el_w >= 1 && N % WARP_SIZE == 0, "row must be a whole number of sub-group widths");
float reg[el_w];
const int lane = sg.get_local_linear_id();
#pragma unroll
for (int i = 0; i < el_w; ++i) {
reg[i] = src[i * WARP_SIZE + lane] * scale;
}
// Butterflies inside the sub-group. The partner of a lane with bit h clear is the
// lower index of the pair, so it takes the sum and the upper takes lower - upper.
#pragma unroll
for (int h = 1; h < WARP_SIZE; h *= 2) {
#pragma unroll
for (int j = 0; j < el_w; ++j) {
const float val = reg[j];
const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE);
reg[j] = (lane & h) == 0 ? val + val2 : val2 - val;
}
}
// Butterflies across registers: h is a multiple of WARP_SIZE, so the partner of
// element i*WARP_SIZE + lane lives in reg[i + h/WARP_SIZE] on the same lane.
#pragma unroll
for (int h = WARP_SIZE; h < N; h *= 2) {
const int step = h / WARP_SIZE;
#pragma unroll
for (int j = 0; j < el_w; j += 2 * step) {
#pragma unroll
for (int k = 0; k < step; ++k) {
const float x = reg[j + k];
const float y = reg[j + k + step];
reg[j + k] = x + y;
reg[j + k + step] = x - y;
}
}
}
#pragma unroll
for (int i = 0; i < el_w; ++i) {
dst[i * WARP_SIZE + lane] = reg[i];
}
}
template <int N>
static void launch_fwht(const float * src, float * dst, const int64_t n_rows, const float scale,
dpct::queue_ptr stream) {
constexpr int rows_per_block = 4;
const int64_t num_blocks = (n_rows + rows_per_block - 1) / rows_per_block;
// dim 1 is the fastest-varying, so a sub-group is exactly one row's WARP_SIZE lanes.
const sycl::range<2> global(num_blocks * rows_per_block, WARP_SIZE);
const sycl::range<2> local(rows_per_block, WARP_SIZE);
stream->parallel_for(sycl::nd_range<2>(global, local),
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
fwht_kernel<N>(src, dst, n_rows, scale, item);
});
}
bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
return false;
}
if (!ggml_are_same_shape(src, dst)) {
return false;
}
if (!ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) {
return false;
}
const int n = (int) src->ne[0];
const int64_t rows = ggml_nrows(src);
const float * src_d = (const float *) src->data;
float * dst_d = (float *) dst->data;
dpct::queue_ptr stream = ctx.stream();
const float scale = 1.0f / std::sqrt((float) n);
switch (n) {
case 64:
launch_fwht<64>(src_d, dst_d, rows, scale, stream);
return true;
case 128:
launch_fwht<128>(src_d, dst_d, rows, scale, stream);
return true;
case 256:
launch_fwht<256>(src_d, dst_d, rows, scale, stream);
return true;
case 512:
launch_fwht<512>(src_d, dst_d, rows, scale, stream);
return true;
default:
return false;
}
}
-12
View File
@@ -1,12 +0,0 @@
#ifndef GGML_SYCL_FWHT_HPP
#define GGML_SYCL_FWHT_HPP
#include "common.hpp"
// Fast Walsh-Hadamard transform, the fast path for a MUL_MAT whose src0 ggml has
// tagged GGML_HINT_SRC0_IS_HADAMARD. src0 is not read at all. Returns false if the
// shape is not one this can serve, in which case the caller must fall through to the
// ordinary mat-mul dispatch.
bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst);
#endif // GGML_SYCL_FWHT_HPP
+1 -23
View File
@@ -58,7 +58,6 @@
#include "ggml-sycl/backend.hpp"
#include "ggml-sycl/common.hpp"
#include "ggml-sycl/element_wise.hpp"
#include "ggml-sycl/fwht.hpp"
#include "ggml-sycl/gemm.hpp"
#include "ggml-sycl/getrows.hpp"
#include "ggml-sycl/norm.hpp"
@@ -109,14 +108,7 @@ int g_ggml_sycl_enable_host_pinned_mem = 1;
static ggml_sycl_device_info ggml_sycl_init() {
ggml_sycl_device_info info = {};
// Do not hard crash when there exists no SYCL devices.
// We want to allow the user to use non-SYCL tools when SYCL is compiled (such as llama-quantize)
try {
info.device_count = dpct::dev_mgr::instance().device_count();
} catch (sycl::exception const &exc) {
GGML_LOG_INFO("%s: no SYCL device available: %s\n", __func__, exc.what());
info.device_count = 0;
}
info.device_count = dpct::dev_mgr::instance().device_count();
if (info.device_count == 0) {
GGML_LOG_ERROR("%s: failed to initialize: %s\n", GGML_SYCL_NAME, __func__);
return info;
@@ -4481,18 +4473,6 @@ static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor *
static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
// Handle HADAMARAD hint given from further up the pipeline and pass it to the correct
// kernel.
//
// The op check is not redundant: this backend also routes MUL_MAT_ID through here with a
// stack copy of dst, which carries MUL_MAT_ID's own op_params. ggml_mul_mat_set_hint()
// asserts GGML_OP_MUL_MAT for the same reason.
if (dst->op == GGML_OP_MUL_MAT && ggml_get_op_params_i32(dst, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
ggml_sycl_op_fwht(ctx, src1, dst)) {
return;
}
const bool split = ggml_backend_buffer_is_sycl_split(src0->buffer);
int64_t min_compute_capability = INT_MAX;
@@ -6242,8 +6222,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
}
case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK:
// FIXME: support ggml_rope_set_offset
return ((const int32_t *) op->op_params)[15] == 0;
case GGML_OP_IM2COL:
case GGML_OP_IM2COL_3D:
case GGML_OP_UPSCALE:
+9 -107
View File
@@ -913,7 +913,6 @@ struct vk_device_struct {
vk_pipeline pipeline_quantize_q8_1_x4;
vk_pipeline pipeline_dequant[GGML_TYPE_COUNT];
vk_pipeline pipeline_dequant_transpose[GGML_TYPE_COUNT]; // fused dequant+transpose for FA quant-KV
vk_pipeline pipeline_dequant_mul_mat_vec_f32_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT][mul_mat_vec_max_cols];
vk_pipeline pipeline_dequant_mul_mat_vec_f16_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT][mul_mat_vec_max_cols];
vk_pipeline pipeline_dequant_mul_mat_vec_id_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT];
@@ -963,7 +962,6 @@ struct vk_device_struct {
vk_pipeline pipeline_cpy_f32_quant[GGML_TYPE_COUNT];
vk_pipeline pipeline_cpy_quant_f32[GGML_TYPE_COUNT];
vk_pipeline pipeline_cpy_transpose_16, pipeline_cpy_transpose_32;
vk_pipeline pipeline_cpy_transpose_02_16, pipeline_cpy_transpose_02_32;
// [src0 0=fp32,1=fp16][dst]
vk_pipeline pipeline_set_rows_i32[2][GGML_TYPE_COUNT];
vk_pipeline pipeline_set_rows_i64[2][GGML_TYPE_COUNT];
@@ -1646,7 +1644,6 @@ struct vk_op_rope_push_constants {
uint32_t rope_mode;
uint32_t nrows;
uint32_t n_dims;
uint32_t n_offs;
float freq_scale;
float freq_base;
float ext_factor;
@@ -3385,10 +3382,10 @@ static void ggml_vk_queue_command_pools_cleanup(vk_device& device) {
// Arbitrary frequency to cleanup/reuse command buffers
static constexpr uint32_t cleanup_frequency = 10;
if (device->compute_queue && device->compute_queue->cmd_pool.buffers_in_use() >= cleanup_frequency) {
if (device->compute_queue->cmd_pool.buffers_in_use() >= cleanup_frequency) {
ggml_vk_command_pool_cleanup(device, device->compute_queue->cmd_pool);
}
if (device->transfer_queue && device->transfer_queue->cmd_pool.buffers_in_use() >= cleanup_frequency) {
if (device->transfer_queue->cmd_pool.buffers_in_use() >= cleanup_frequency) {
ggml_vk_command_pool_cleanup(device, device->transfer_queue->cmd_pool);
}
}
@@ -5392,7 +5389,6 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q5_0], "dequant_q5_0", dequant_q5_0_len, dequant_q5_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q5_1], "dequant_q5_1", dequant_q5_1_len, dequant_q5_1_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q8_0], "dequant_q8_0", dequant_q8_0_len, dequant_q8_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant_transpose[GGML_TYPE_Q8_0], "dequant_q8_0_transpose", dequant_q8_0_transpose_len, dequant_q8_0_transpose_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q2_K], "dequant_q2_k", dequant_q2_k_len, dequant_q2_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_TQ2_0], "dequant_tq2_0", dequant_tq2_0_len, dequant_tq2_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q3_K], "dequant_q3_k", dequant_q3_k_len, dequant_q3_k_data, "main", 2, 5 * sizeof(uint32_t), {256 * 64, 1, 1}, {}, 1);
@@ -5529,8 +5525,6 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_32, "cpy_transpose_32", cpy_transpose_32_len, cpy_transpose_32_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_16, "cpy_transpose_16", cpy_transpose_16_len, cpy_transpose_16_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_02_32, "cpy_transpose_02_32", cpy_transpose_02_32_len, cpy_transpose_02_32_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_transpose_02_16, "cpy_transpose_02_16", cpy_transpose_02_16_len, cpy_transpose_02_16_data, "main", 2, sizeof(vk_op_unary_push_constants), {1, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_Q1_0], "cpy_f32_q1_0", cpy_f32_q1_0_len, cpy_f32_q1_0_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_Q2_0], "cpy_f32_q2_0", cpy_f32_q2_0_len, cpy_f32_q2_0_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
@@ -8937,18 +8931,6 @@ static vk_pipeline ggml_vk_get_cpy_pipeline(ggml_backend_vk_context * ctx, const
}
}
// Same, for a 0<->2 swap: src dim2 is the innermost dimension.
bool transpose02 = dst && !contig && src->nb[2] == ggml_type_size(to) &&
ggml_is_contiguous(dst) && ggml_are_same_shape(dst, src);
if (transpose02 && src->type == to) {
if (ggml_type_size(to) == 4) {
return ctx->device->pipeline_cpy_transpose_02_32;
} else if (ggml_type_size(to) == 2) {
return ctx->device->pipeline_cpy_transpose_02_16;
}
}
if (src->type == GGML_TYPE_F32 && to == GGML_TYPE_F32) {
if (contig) {
return ctx->device->pipeline_contig_cpy_f32_f32;
@@ -10825,32 +10807,9 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
const bool f32acc = !ctx->device->fp16 || dst->op_params[3] == GGML_PREC_F32 || k->type == GGML_TYPE_BF16;
// dequant K/V once into an f16 scratch, reordered KV layout so FA can read without a stride
auto is_dense_kv_cache = [](const ggml_tensor * t) {
return t->nb[0] == ggml_type_size(t->type) &&
t->nb[2] == ggml_row_size(t->type, t->ne[0]) &&
t->nb[1] == t->nb[2] * t->ne[2] &&
t->nb[3] == t->nb[1] * t->ne[1];
};
const bool k_quant = k->type != GGML_TYPE_F16 && k->type != GGML_TYPE_BF16 && k->type != GGML_TYPE_F32;
const bool v_quant = v->type != GGML_TYPE_F16 && v->type != GGML_TYPE_BF16 && v->type != GGML_TYPE_F32;
const bool use_dequant_kv = k_quant && v_quant && neq1 >= 64 &&
is_dense_kv_cache(k) && is_dense_kv_cache(v) &&
(uint64_t)ggml_nelements(k) * sizeof(ggml_fp16_t) <= ctx->device->properties.limits.maxStorageBufferRange &&
(uint64_t)ggml_nelements(v) * sizeof(ggml_fp16_t) <= ctx->device->properties.limits.maxStorageBufferRange &&
ctx->device->pipeline_dequant_transpose[k->type] != nullptr &&
ctx->device->pipeline_dequant_transpose[v->type] != nullptr &&
// coopmat2 path does not benefit from the f16 scratch
!ctx->device->coopmat2 &&
// Intel Xe1 regresses, see PR 25494
(ctx->device->vendor_id != VK_VENDOR_ID_INTEL ||
(ctx->device->coopmat_support && ctx->device->architecture != vk_device_architecture::INTEL_XE1));
const ggml_type k_type_eff = use_dequant_kv ? GGML_TYPE_F16 : k->type;
const ggml_type v_type_eff = use_dequant_kv ? GGML_TYPE_F16 : v->type;
// For scalar/coopmat1 FA, we can use the "large" size to accommodate qga.
// For coopmat2 FA, we always use the small size (which is still pretty large for gqa).
vk_fa_tuning_params tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, 512, KV, k_type_eff, v_type_eff, f32acc);
vk_fa_tuning_params tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, 512, KV, k->type, v->type, f32acc);
const uint32_t max_gqa = std::min(tuning_params.block_rows, 32u);
if (N <= 8 && qk_ratio > 1 && qk_ratio <= max_gqa &&
@@ -10863,7 +10822,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
workgroups_y /= gqa_ratio;
}
tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, N, KV, k_type_eff, v_type_eff, f32acc);
tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, N, KV, k->type, v->type, f32acc);
const uint32_t q_stride = (uint32_t)(nbq1 / ggml_type_size(q->type));
uint32_t k_stride = (uint32_t)(nbk1 / ggml_type_size(k->type));
@@ -10877,17 +10836,6 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
v_stride /= 4;
}
uint32_t nbk2_eff = (uint32_t)nbk2, nbk3_eff = (uint32_t)nbk3;
uint32_t nbv2_eff = (uint32_t)nbv2, nbv3_eff = (uint32_t)nbv3;
if (use_dequant_kv) {
k_stride = HSK;
v_stride = HSV;
nbk2_eff = (uint32_t)((uint64_t)HSK * KV * sizeof(ggml_fp16_t));
nbk3_eff = (uint32_t)((uint64_t)HSK * KV * nek2 * sizeof(ggml_fp16_t));
nbv2_eff = (uint32_t)((uint64_t)HSV * KV * sizeof(ggml_fp16_t));
nbv3_eff = (uint32_t)((uint64_t)HSV * KV * nev2 * sizeof(ggml_fp16_t));
}
const uint32_t alignment = tuning_params.block_cols;
bool aligned = (KV % alignment) == 0 &&
// the "aligned" shader variant will forcibly align strides, for performance
@@ -10914,7 +10862,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768 && nem0 >= tuning_params.block_cols * 16
&& (ctx->device->architecture != vk_device_architecture::AMD_GCN || HSK > 256 || HSV > 256);
vk_fa_pipeline_state fa_pipeline_state = get_fa_pipeline_state(ctx->device, tuning_params, HSK, HSV, aligned, f32acc,
mask != nullptr, use_mask_opt, logit_softcap != 0, k_type_eff, v_type_eff);
mask != nullptr, use_mask_opt, logit_softcap != 0, k->type, v->type);
vk_pipeline pipeline = nullptr;
@@ -11018,34 +10966,6 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
vk_subbuffer sinks_buf = sinks ? ggml_vk_tensor_subbuffer(ctx, sinks) : q_buf;
vk_subbuffer mask_opt_buf = use_mask_opt ? ggml_vk_subbuffer(ctx, ctx->prealloc_y, 0) : q_buf;
if (use_dequant_kv) {
const uint64_t fp = sizeof(ggml_fp16_t);
const uint64_t k_f16_sz = (uint64_t)ggml_nelements(k) * fp;
const uint64_t v_f16_sz = (uint64_t)ggml_nelements(v) * fp;
if (ctx->prealloc_size_x < k_f16_sz + v_f16_sz) {
ctx->prealloc_size_x = k_f16_sz + v_f16_sz;
ggml_vk_preallocate_buffers(ctx, subctx);
}
vk_pipeline tr_k = ctx->device->pipeline_dequant_transpose[k->type];
vk_pipeline tr_v = ctx->device->pipeline_dequant_transpose[v->type];
ggml_pipeline_request_descriptor_sets(ctx, tr_k, 1);
ggml_pipeline_request_descriptor_sets(ctx, tr_v, 1);
if (ctx->prealloc_x_need_sync) {
ggml_vk_sync_buffers(ctx, subctx);
}
vk_subbuffer k_dst = vk_subbuffer{ ctx->prealloc_x, 0, k_f16_sz };
vk_subbuffer v_dst = vk_subbuffer{ ctx->prealloc_x, k_f16_sz, v_f16_sz };
const uint32_t k_nel = (uint32_t)ggml_nelements(k);
const uint32_t v_nel = (uint32_t)ggml_nelements(v);
{ const std::vector<uint32_t> pc = { (uint32_t)HSK, (uint32_t)nek2, (uint32_t)KV, 0, k_nel };
ggml_vk_dispatch_pipeline(ctx, subctx, tr_k, { k_buf, k_dst }, pc, { k_nel, 1, 1 }); }
{ const std::vector<uint32_t> pc = { (uint32_t)HSV, (uint32_t)nev2, (uint32_t)KV, 0, v_nel };
ggml_vk_dispatch_pipeline(ctx, subctx, tr_v, { v_buf, v_dst }, pc, { v_nel, 1, 1 }); }
ggml_vk_sync_buffers(ctx, subctx);
k_buf = k_dst;
v_buf = v_dst;
}
uint32_t mask_n_head_log2 = ((sinks != nullptr) << 24) | n_head_log2;
if (use_mask_opt)
@@ -11075,8 +10995,8 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
(uint32_t)nev2, (uint32_t)nev3,
nem1, nem2, nem3,
q_stride, (uint32_t)nbq2, (uint32_t)nbq3,
k_stride, nbk2_eff, nbk3_eff,
v_stride, nbv2_eff, nbv3_eff,
k_stride, (uint32_t)nbk2, (uint32_t)nbk3,
v_stride, (uint32_t)nbv2, (uint32_t)nbv3,
scale, max_bias, logit_softcap,
mask_n_head_log2, m0, m1,
gqa_ratio, split_kv, split_k };
@@ -11118,10 +11038,6 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
{q_buf, k_buf, v_buf, mask_buf, sinks_buf, dst_buf, mask_opt_buf},
pc, { workgroups_x, workgroups_y, workgroups_z });
}
if (use_dequant_kv) {
ctx->prealloc_x_need_sync = true;
}
}
static vk_conv_shapes ggml_vk_conv_select_shape(ggml_backend_vk_context * ctx, uint32_t K, uint32_t NPQ) {
@@ -12276,16 +12192,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
elements = { ne, 1, 1 };
}
if (pipeline == ctx->device->pipeline_cpy_transpose_02_32 ||
pipeline == ctx->device->pipeline_cpy_transpose_02_16) {
// 32x32 tiles over dims 0 and 2; dim1 and dim3 are the batch
elements[0] = (uint32_t)CEIL_DIV(dst->ne[0], 32);
elements[1] = (uint32_t)CEIL_DIV(dst->ne[2], 32);
elements[2] = (uint32_t)(dst->ne[1]*dst->ne[3]);
elements[0] = std::min(elements[0], ctx->device->properties.limits.maxComputeWorkGroupCount[0]);
elements[1] = std::min(elements[1], ctx->device->properties.limits.maxComputeWorkGroupCount[1]);
elements[2] = std::min(elements[2], ctx->device->properties.limits.maxComputeWorkGroupCount[2]);
} else if (pipeline == ctx->device->pipeline_cpy_transpose_32 ||
if (pipeline == ctx->device->pipeline_cpy_transpose_32 ||
pipeline == ctx->device->pipeline_cpy_transpose_16) {
// 32x32 tiles
elements[0] = (uint32_t)CEIL_DIV(dst->ne[0], 32);
@@ -13213,7 +13120,6 @@ static uint32_t ggml_vk_rms_partials_size(ggml_backend_vk_context * ctx, const g
static vk_op_rope_push_constants ggml_vk_make_rope_constants(const ggml_tensor *dst, const ggml_tensor *src0, const bool has_ff, bool backprop, const uint32_t set_rows_stride) {
const int n_dims = ((const int32_t *) dst->op_params)[1];
const int mode = ((const int32_t *) dst->op_params)[2];
const int n_offs = ((const int32_t *) dst->op_params)[15];
// const int n_ctx = ((const int32_t *) dst->op_params)[3];
const int n_ctx_orig = ((const int32_t *) dst->op_params)[4];
const float freq_base = ((const float *) dst->op_params)[5];
@@ -13243,7 +13149,7 @@ static vk_op_rope_push_constants ggml_vk_make_rope_constants(const ggml_tensor *
uint32_t nb13 = dst->nb[3] / ggml_type_size(dst->type);
vk_op_rope_push_constants rope {
(uint32_t)mode, (uint32_t)ggml_nrows(src0), (uint32_t)n_dims, (uint32_t)n_offs, freq_scale,
(uint32_t)mode, (uint32_t)ggml_nrows(src0), (uint32_t)n_dims, freq_scale,
freq_base, ext_factor, attn_factor, {corr_dims[0], corr_dims[1]}, theta_scale, has_ff,
{ sections[0], sections[1], sections[2], sections[3] }, is_imrope, backprop, set_rows_stride,
@@ -19289,10 +19195,6 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
tensor_clone = ggml_rope_ext_back(ggml_ctx, src_clone[0], src_clone[1], src_clone[2], n_dims, mode, n_ctx_orig_ggml, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
}
}
const int n_offs = ((int32_t *) tensor->op_params)[15];
if (n_offs != 0) {
tensor_clone = ggml_rope_set_offset(tensor_clone, n_offs);
}
} else if (tensor->op == GGML_OP_UNARY) {
switch (ggml_get_unary_op(tensor)) {
case GGML_UNARY_OP_EXP:
@@ -1,61 +0,0 @@
#version 450
#include "types.glsl"
#include "generic_unary_head.glsl"
// workgroup does 32x32 tile, but uses 32x8 threads
#define TILE_DIM 32
layout(local_size_x = 32, local_size_y = 8, local_size_z = 1) in;
// +1 padding avoids shared-memory bank conflicts on the transposed read
shared uint sh[TILE_DIM][TILE_DIM + 1];
void iter(uvec3 wg_id) {
const uint tile_i0 = wg_id.x; // tiles dst ne10 (== src ne00)
const uint tile_i2 = wg_id.y; // tiles dst ne12 (== src ne02)
const uint tid_col = gl_LocalInvocationID.x;
const uint tid_row = gl_LocalInvocationID.y;
const uint i1 = wg_id.z % p.ne11;
const uint i3 = wg_id.z / p.ne11;
const uint i01 = i1;
const uint i03 = i3;
[[unroll]] for (uint y = 0; y < 4; ++y) {
const uint i00 = tile_i0 * TILE_DIM + tid_row + 8 * y;
const uint i02 = tile_i2 * TILE_DIM + tid_col;
if (i00 < p.ne00 && i01 < p.ne01 && i02 < p.ne02 && i03 < p.ne03) {
const uint src_idx = i00 * p.nb00 + i01 * p.nb01 + i02 * p.nb02 + i03 * p.nb03;
sh[tid_row + 8 * y][tid_col] = uint(data_a[get_aoffset() + src_idx]);
}
}
barrier();
[[unroll]] for (uint y = 0; y < 4; ++y) {
const uint i0 = tile_i0 * TILE_DIM + tid_col;
const uint i2 = tile_i2 * TILE_DIM + tid_row + 8 * y;
if (i0 < p.ne10 && i1 < p.ne11 && i2 < p.ne12 && i3 < p.ne13) {
const uint dst_idx = i0 * p.nb10 + i1 * p.nb11 + i2 * p.nb12 + i3 * p.nb13;
data_d[get_doffset() + dst_idx] = D_TYPE(sh[tid_col][tid_row + 8 * y]);
}
}
}
#define CEIL_DIV(a, b) (((a) + (b) - 1) / (b))
void main() {
bool need_barrier = false;
for (uint z = gl_WorkGroupID.z; z < p.ne11 * p.ne13; z += gl_NumWorkGroups.z) {
for (uint y = gl_WorkGroupID.y; y < CEIL_DIV(p.ne12, TILE_DIM); y += gl_NumWorkGroups.y) {
for (uint x = gl_WorkGroupID.x; x < CEIL_DIV(p.ne10, TILE_DIM); x += gl_NumWorkGroups.x) {
if (need_barrier) {
barrier();
}
need_barrier = true;
iter(uvec3(x, y, z));
}
}
}
}
@@ -18,18 +18,7 @@ void main() {
return;
}
#ifdef DEQUANT_TRANSPOSE
// read [HS, NH, KV, NS], write [HS, KV, NH, NS]
const uint HS = p.M, NH = p.K, KVn = p.stride_a;
const uint e0 = ib * 32;
const uint b_idx = (e0 % HS)
+ ((e0 / (HS * NH)) % KVn) * HS
+ ((e0 / HS) % NH) * (HS * KVn)
+ (e0 / (HS * NH * KVn)) * (HS * KVn * NH)
+ 16 * il;
#else
const uint b_idx = 1024*i + 32*ir + 16*il;
#endif
const float d = float(data_a[ib].d);
@@ -50,21 +50,19 @@ void rope_norm(const uint i0, const uint i1, const uint i2, const uint i3, rope_
}
idst += p.d_offset;
if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
if (i0 >= p.n_dims) {
rope_data_d[idst + 0] = ROPE_D_TYPE(rope_data_a[ix + 0]);
rope_data_d[idst + 1] = ROPE_D_TYPE(rope_data_a[ix + 1]);
return;
}
const uint iw = i0 - p.n_offs; // relative idx
const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, i0/2.0f);
const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, iw/2.0f);
const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
float cos_theta, sin_theta;
rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
const float x0 = float(rope_data_a[ix + 0]);
const float x1 = float(rope_data_a[ix + 1]);
@@ -89,28 +87,25 @@ void rope_neox(const uint i0, const uint i1, const uint i2, const uint i3, rope_
}
idst += p.d_offset;
if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
if (i0 >= p.n_dims) {
rope_data_d[idst + i0/2 + 0] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 0]);
rope_data_d[idst + i0/2 + 1] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 1]);
return;
}
const uint iw = i0 - p.n_offs; // relative idx
const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, i0/2.0f);
const float theta_base = rope_data_pos[i2] * pow(p.theta_scale, iw/2.0f);
const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
float cos_theta, sin_theta;
rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
// idst/ix point at channel i0/2; the first channel of the rotated pair is p.n_offs + iw/2 = i0/2 + p.n_offs/2
const float x0 = float(rope_data_a[ix + p.n_offs/2 + 0]);
const float x1 = float(rope_data_a[ix + p.n_offs/2 + p.n_dims/2]);
const float x0 = float(rope_data_a[ix + 0]);
const float x1 = float(rope_data_a[ix + p.n_dims/2]);
rope_data_d[idst + p.n_offs/2 + 0] = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
rope_data_d[idst + p.n_offs/2 + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
rope_data_d[idst + 0] = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
rope_data_d[idst + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
}
@@ -130,56 +125,53 @@ void rope_multi(const uint i0, const uint i1, const uint i2, const uint i3, rope
}
idst += p.d_offset;
if (i0 < p.n_offs || i0 >= p.n_offs + p.n_dims) {
if (i0 >= p.n_dims) {
rope_data_d[idst + i0/2 + 0] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 0]);
rope_data_d[idst + i0/2 + 1] = ROPE_D_TYPE(rope_data_a[ix + i0/2 + 1]);
return;
}
const uint iw = i0 - p.n_offs; // relative idx
const int sect_dims = p.sections[0] + p.sections[1] + p.sections[2] + p.sections[3];
const int sec_w = p.sections[1] + p.sections[0];
const uint sector = (iw / 2) % sect_dims;
const uint sector = (i0 / 2) % sect_dims;
float theta_base = 0.0;
if (p.is_imrope != 0) {
if (sector % 3 == 1 && sector < 3 * p.sections[1]) {
theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, i0/2.0f);
} else if (sector % 3 == 2 && sector < 3 * p.sections[2]) {
theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, i0/2.0f);
} else if (sector % 3 == 0 && sector < 3 * p.sections[0]) {
theta_base = rope_data_pos[i2]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2]*pow(p.theta_scale, i0/2.0f);
} else {
theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, i0/2.0f);
}
} else {
if (sector < p.sections[0]) {
theta_base = rope_data_pos[i2]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2]*pow(p.theta_scale, i0/2.0f);
}
else if (sector >= p.sections[0] && sector < sec_w) {
theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 1]*pow(p.theta_scale, i0/2.0f);
}
else if (sector >= sec_w && sector < sec_w + p.sections[2]) {
theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 2]*pow(p.theta_scale, i0/2.0f);
}
else if (sector >= sec_w + p.sections[2]) {
theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, iw/2.0f);
theta_base = rope_data_pos[i2 + p.ne02 * 3]*pow(p.theta_scale, i0/2.0f);
}
}
const float freq_factor = p.has_ff != 0 ? rope_data_ff[iw/2] : 1.0f;
const float freq_factor = p.has_ff != 0 ? rope_data_ff[i0/2] : 1.0f;
float cos_theta, sin_theta;
rope_yarn(theta_base / freq_factor, iw, cos_theta, sin_theta, p);
rope_yarn(theta_base / freq_factor, i0, cos_theta, sin_theta, p);
// idst/ix point at channel i0/2; the first channel of the rotated pair is p.n_offs + iw/2 = i0/2 + p.n_offs/2
const float x0 = float(rope_data_a[ix + p.n_offs/2 + 0]);
const float x1 = float(rope_data_a[ix + p.n_offs/2 + p.n_dims/2]);
const float x0 = float(rope_data_a[ix + 0]);
const float x1 = float(rope_data_a[ix + p.n_dims/2]);
rope_data_d[idst + p.n_offs/2 + 0] = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
rope_data_d[idst + p.n_offs/2 + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
rope_data_d[idst + 0] = ROPE_D_TYPE(x0*cos_theta - x1*sin_theta);
rope_data_d[idst + p.n_dims/2] = ROPE_D_TYPE(x0*sin_theta + x1*cos_theta);
}
void rope_vision(const uint i0, const uint i1, const uint i2, const uint i3, rope_params p) {
@@ -5,7 +5,6 @@ struct rope_params {
uint rope_mode;
uint nrows;
uint n_dims;
uint n_offs;
float freq_scale;
float freq_base;
float ext_factor;
@@ -780,10 +780,6 @@ void process_shaders() {
if (tname != "f16" && tname != "bf16") {
string_to_spv("dequant_" + tname, "dequant_" + tname + ".comp", merge_maps(base_dict, {{data_a_key, "1"}, {"D_TYPE", "float16_t"}}));
}
// Fused dequant+transpose variant for FA quant-KV (per-head-contiguous f16 scratch).
if (tname == "q8_0") {
string_to_spv("dequant_" + tname + "_transpose", "dequant_" + tname + ".comp", merge_maps(base_dict, {{data_a_key, "1"}, {"D_TYPE", "float16_t"}, {"DEQUANT_TRANSPOSE", "1"}}));
}
shader = (tname == "f32" || tname == "f16" || tname == "bf16") ? "get_rows.comp" : "get_rows_quant.comp";
@@ -830,8 +826,6 @@ void process_shaders() {
string_to_spv("cpy_transpose_16", "copy_transpose.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});
string_to_spv("cpy_transpose_32", "copy_transpose.comp", {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}});
string_to_spv("cpy_transpose_02_16", "copy_transpose_02.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});
string_to_spv("cpy_transpose_02_32", "copy_transpose_02.comp", {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}});
for (std::string t : {"q1_0", "q2_0", "q4_0", "q4_1", "q5_0", "q5_1", "q8_0", "iq4_nl"}) {
string_to_spv("cpy_f32_" + t, "copy_to_quant.comp", {{"DATA_A_" + to_uppercase(t), "1"}, {"S_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
@@ -954,11 +954,10 @@ struct ggml_webgpu_mul_mat_vec_pipeline_key {
int vectorized;
uint32_t num_cols;
bool use_mmvq;
bool src_overlap;
bool operator==(const ggml_webgpu_mul_mat_vec_pipeline_key & other) const {
return src0_type == other.src0_type && src1_type == other.src1_type && vectorized == other.vectorized &&
num_cols == other.num_cols && use_mmvq == other.use_mmvq && src_overlap == other.src_overlap;
num_cols == other.num_cols && use_mmvq == other.use_mmvq;
}
};
@@ -970,7 +969,6 @@ struct ggml_webgpu_mul_mat_vec_pipeline_key_hash {
ggml_webgpu_hash_combine(seed, key.vectorized);
ggml_webgpu_hash_combine(seed, key.num_cols);
ggml_webgpu_hash_combine(seed, key.use_mmvq);
ggml_webgpu_hash_combine(seed, key.src_overlap);
return seed;
}
};
@@ -979,7 +977,6 @@ struct ggml_webgpu_mul_mat_vec_shader_decisions {
uint32_t wg_size;
uint32_t outputs_per_wg;
uint32_t vec_size;
bool src_overlap = false;
};
struct ggml_webgpu_quantize_q8_pipeline_key {
@@ -1001,11 +998,10 @@ struct ggml_webgpu_mul_mat_pipeline_key {
ggml_type src1_type;
int vectorized;
int use_subgroup_matrix;
bool src_overlap;
bool operator==(const ggml_webgpu_mul_mat_pipeline_key & other) const {
return src0_type == other.src0_type && src1_type == other.src1_type && vectorized == other.vectorized &&
use_subgroup_matrix == other.use_subgroup_matrix && src_overlap == other.src_overlap;
use_subgroup_matrix == other.use_subgroup_matrix;
}
};
@@ -1016,7 +1012,6 @@ struct ggml_webgpu_mul_mat_pipeline_key_hash {
ggml_webgpu_hash_combine(seed, key.src1_type);
ggml_webgpu_hash_combine(seed, key.vectorized);
ggml_webgpu_hash_combine(seed, key.use_subgroup_matrix);
ggml_webgpu_hash_combine(seed, key.src_overlap);
return seed;
}
};
@@ -1039,7 +1034,6 @@ struct ggml_webgpu_mul_mat_shader_decisions {
uint32_t subgroup_matrix_n;
uint32_t mul_mat_wg_size;
bool src_overlap = false;
};
/** MUL_MAT_ID **/
@@ -1956,7 +1950,7 @@ class ggml_webgpu_shader_lib {
return quantize_q8_pipelines[key];
}
webgpu_pipeline get_mul_mat_vec_pipeline(const ggml_webgpu_shader_lib_context & context, bool src_overlap) {
webgpu_pipeline get_mul_mat_vec_pipeline(const ggml_webgpu_shader_lib_context & context) {
ggml_webgpu_mul_mat_vec_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
@@ -1967,7 +1961,6 @@ class ggml_webgpu_shader_lib {
key.num_cols = context.dst->ne[1];
key.use_mmvq =
ggml_webgpu_can_use_mmvq(context.src0, context.src1, context.supports_dot_product, context.vendor);
key.src_overlap = src_overlap;
auto it = mul_mat_vec_pipelines.find(key);
if (it != mul_mat_vec_pipelines.end()) {
@@ -2075,11 +2068,6 @@ class ggml_webgpu_shader_lib {
defines.push_back("Q8_1_T");
}
if (key.src_overlap) {
defines.push_back("SRC_OVERLAP");
variant += "_src_overlap";
}
defines.push_back(std::string("WG_SIZE=") + std::to_string(wg_size));
defines.push_back(std::string("OUTPUTS_PER_WG=") + std::to_string(outputs_per_wg));
defines.push_back(context.supports_subgroups ? "USE_SUBGROUP_REDUCTION" : "USE_WORKGROUP_REDUCTION");
@@ -2101,7 +2089,7 @@ class ggml_webgpu_shader_lib {
return mul_mat_vec_pipelines[key];
}
webgpu_pipeline get_mul_mat_fast_pipeline(const ggml_webgpu_shader_lib_context & context, bool src_overlap) {
webgpu_pipeline get_mul_mat_fast_pipeline(const ggml_webgpu_shader_lib_context & context) {
ggml_webgpu_mul_mat_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
@@ -2110,7 +2098,6 @@ class ggml_webgpu_shader_lib {
1 :
0;
key.use_subgroup_matrix = context.supports_subgroup_matrix;
key.src_overlap = src_overlap;
auto it = mul_mat_fast_pipelines.find(key);
if (it != mul_mat_fast_pipelines.end()) {
@@ -2229,11 +2216,6 @@ class ggml_webgpu_shader_lib {
variant += "_vectorized";
}
if (key.src_overlap) {
defines.push_back("SRC_OVERLAP");
variant += "_src_overlap";
}
if (!key.use_subgroup_matrix) {
defines.push_back("WORKGROUP_SIZE_M=" + std::to_string(WEBGPU_MUL_MAT_WG_SIZE_M) + "u");
defines.push_back("WORKGROUP_SIZE_N=" + std::to_string(WEBGPU_MUL_MAT_WG_SIZE_N) + "u");
+24 -43
View File
@@ -1628,65 +1628,48 @@ static webgpu_encoded_op ggml_webgpu_mul_mat(webgpu_context & ctx,
// Get or create pipeline
webgpu_pipeline pipeline;
std::vector<webgpu_dispatch_desc> dispatches;
const bool src_overlap = ggml_webgpu_tensor_binding_overlap(ctx->global_ctx, src0, src1) && !use_mmvq;
if (use_mat_vec) {
if (use_mmvq) {
ggml_webgpu_quantize_q8_dispatch(ctx, src0, src1, dst, dispatches);
}
pipeline = ctx->shader_lib->get_mul_mat_vec_pipeline(shader_lib_ctx, src_overlap);
pipeline = ctx->shader_lib->get_mul_mat_vec_pipeline(shader_lib_ctx);
} else {
pipeline = ctx->shader_lib->get_mul_mat_fast_pipeline(shader_lib_ctx, src_overlap);
}
uint32_t offset_src0 = (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src0) / ggml_type_size(src0->type));
uint32_t offset_src1 = (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src1) / ggml_type_size(src1->type));
size_t merged_offset = 0;
size_t merged_size = 0;
if (src_overlap) {
const ggml_webgpu_merged_binding_range merged_range =
ggml_webgpu_tensor_merged_binding_range(ctx, { src0, src1 });
merged_offset = merged_range.offset;
merged_size = merged_range.size;
offset_src0 = ggml_webgpu_tensor_merged_element_offset(src0, merged_range);
offset_src1 = ggml_webgpu_tensor_merged_element_offset(src1, merged_range);
pipeline = ctx->shader_lib->get_mul_mat_fast_pipeline(shader_lib_ctx);
}
// Build params
std::vector<uint32_t> params = { offset_src0,
offset_src1,
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
(uint32_t) dst->ne[0],
(uint32_t) dst->ne[1],
(uint32_t) src0->ne[0],
(uint32_t) (src0->nb[1] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[1] / ggml_type_size(src1->type)),
(uint32_t) (src0->nb[2] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[2] / ggml_type_size(src1->type)),
(uint32_t) (src0->nb[3] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[3] / ggml_type_size(src1->type)),
(uint32_t) src0->ne[2],
(uint32_t) src0->ne[3],
(uint32_t) (src1->ne[2] / src0->ne[2]),
(uint32_t) (src1->ne[3] / src0->ne[3]) };
std::vector<uint32_t> params = {
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src0) / ggml_type_size(src0->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src1) / ggml_type_size(src1->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
(uint32_t) dst->ne[0],
(uint32_t) dst->ne[1],
(uint32_t) src0->ne[0],
(uint32_t) (src0->nb[1] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[1] / ggml_type_size(src1->type)),
(uint32_t) (src0->nb[2] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[2] / ggml_type_size(src1->type)),
(uint32_t) (src0->nb[3] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[3] / ggml_type_size(src1->type)),
(uint32_t) src0->ne[2],
(uint32_t) src0->ne[3],
(uint32_t) (src1->ne[2] / src0->ne[2]),
(uint32_t) (src1->ne[3] / src0->ne[3])
};
// Build bind group entries
std::vector<wgpu::BindGroupEntry> entries = {};
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
if (use_mmvq) {
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
auto & mmvq_qq8_entry = dispatches[0].bind_group_entries[1];
entries.push_back(ggml_webgpu_make_bind_group_entry(1, ggml_webgpu_tensor_buf(dst), mmvq_qq8_entry.offset,
mmvq_qq8_entry.size));
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
} else if (src_overlap) {
entries.push_back(
ggml_webgpu_make_bind_group_entry(0, ggml_webgpu_tensor_buf(src0), merged_offset, merged_size));
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 1, dst));
} else {
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0));
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 1, src1));
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
}
entries.push_back(ggml_webgpu_make_tensor_bind_group_entry(ctx, 2, dst));
// Calculate workgroup dimensions
uint32_t wg_x = 1;
@@ -4472,9 +4455,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
supports_op = (op->type == GGML_TYPE_F32 && src0->type == GGML_TYPE_F32) && ggml_is_contiguous_rows(src0);
break;
case GGML_OP_ROPE:
// FIXME: support ggml_rope_set_offset
supports_op =
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) && ((const int32_t *) op->op_params)[15] == 0;
supports_op = op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16;
break;
case GGML_OP_GLU:
switch (ggml_get_glu_op(op)) {
@@ -1,7 +1,3 @@
#ifndef SRC0
#define SRC0 src0
#endif
#ifdef BYTE_HELPERS
fn get_byte(value: u32, index: u32) -> u32 {
return (value >> (index * 8)) & 0xFF;
@@ -50,7 +46,7 @@ fn load_f16_as_f32_at_src(byte_offset: u32) -> f32 {
#ifdef DECLARE_BYTE_LOADERS_SRC0
fn load_u16_at_src0(byte_offset: u32) -> u32 {
let word = SRC0[byte_offset / 4u];
let word = src0[byte_offset / 4u];
let shift = (byte_offset & 0x2u) * 8u;
return (word >> shift) & 0xFFFFu;
}
@@ -59,14 +55,14 @@ fn load_u16_at_src0(byte_offset: u32) -> u32 {
// Caller extracts the 16-bit half it needs via & 0xFFFFu or >> 16u.
// this is used in k-quants for better performance
fn load_u32_at_src0_aligned(byte_offset: u32) -> u32 {
return SRC0[(byte_offset & ~3u) / 4u];
return src0[(byte_offset & ~3u) / 4u];
}
fn load_u32_at_src0(byte_offset: u32) -> u32 {
let word_idx = byte_offset / 4u;
let shift = (byte_offset & 0x3u) * 8u;
let lo = SRC0[word_idx];
let hi = SRC0[word_idx + 1u];
let lo = src0[word_idx];
let hi = src0[word_idx + 1u];
let shifted = (lo >> shift) | (hi << (32u - shift));
return select(shifted, lo, shift == 0u);
}
@@ -77,7 +73,7 @@ fn load_f16_at_src0(byte_offset: u32) -> f16 {
}
fn load_f16_as_f32_at_src0(byte_offset: u32) -> f32 {
let word = SRC0[byte_offset / 4u];
let word = src0[byte_offset / 4u];
let shift = (byte_offset & 0x2u) * 8u;
let d_bits = (word >> shift) & 0xFFFFu;
return unpack2x16float(d_bits)[0];
@@ -1,10 +1,3 @@
#ifndef SRC0
#define SRC0 src0
#endif
#ifndef SRC1
#define SRC1 src1
#endif
#ifdef VEC
#define VEC_SIZE 4
#define SHMEM_TYPE vec4<f16>
@@ -46,7 +39,7 @@ fn init_shmem_src0(thread_id: u32, batch_offset: u32, offset_m: u32, k_outer: u3
let src0_idx = batch_offset + global_m * params.stride_01 + global_k;
let src0_val = select( // taking a slight performance hit to avoid oob
SRC0_TYPE(0.0),
SRC0[src0_idx/VEC_SIZE],
src0[src0_idx/VEC_SIZE],
global_m < params.m && global_k < params.k);
store_shmem(SHMEM_TYPE(src0_val), elem_idx);
}
@@ -64,7 +57,7 @@ fn init_shmem_src1(thread_id: u32, batch_offset: u32, offset_n: u32, k_outer: u3
let src1_idx = batch_offset + global_n * params.stride_11 + global_k;
let src1_val = select(
SRC1_TYPE(0.0),
SRC1[src1_idx/VEC_SIZE],
src1[src1_idx/VEC_SIZE],
global_n < params.n && global_k < params.k);
store_shmem(SHMEM_TYPE(src1_val), TILE_SRC0_SHMEM + elem_idx);
}
@@ -1,12 +1,8 @@
enable f16;
#define DECLARE_BYTE_LOADERS_SRC0
#ifdef SRC_OVERLAP
#define SRC0 merged_src
#define SRC1 merged_src
#endif
#include "common_decls.tmpl"
#include "mul_mat_decls.tmpl"
#ifdef VEC
@@ -40,17 +36,11 @@ struct MulMatParams {
broadcast3: u32
};
#ifdef SRC_OVERLAP
@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
#define DST_BINDING 1
#else
@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>; // M rows, K columns
@group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>; // K rows, N columns (transposed)
#define DST_BINDING 2
#endif
@group(0) @binding(2) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
@group(0) @binding(3) var<uniform> params: MulMatParams;
fn get_local_n(thread_id: u32) -> u32 {
return thread_id / WORKGROUP_SIZE_M;
@@ -4,10 +4,6 @@ enable subgroups;
enable chromium_experimental_subgroup_matrix;
#define DECLARE_BYTE_LOADERS_SRC0
#ifdef SRC_OVERLAP
#define SRC0 merged_src
#define SRC1 merged_src
#endif
#include "common_decls.tmpl"
#include "mul_mat_decls.tmpl"
@@ -52,17 +48,11 @@ struct MulMatParams {
};
// SRC0_TYPE and SRC1_TYPE are defined in mul_mat_decls, which is included
#ifdef SRC_OVERLAP
@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
#define DST_BINDING 1
#else
@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>; // M rows, K columns
@group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>; // K rows, N columns (transposed)
#define DST_BINDING 2
#endif
@group(0) @binding(2) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<DST_TYPE>; // M rows, N columns (transposed)
@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
@group(0) @binding(3) var<uniform> params: MulMatParams;
const WG_M_SG_TILE_SIZE = SUBGROUP_M * SUBGROUP_MATRIX_M * SUBGROUP_MATRIX_M_SIZE;
const WG_N_SG_TILE_SIZE = SUBGROUP_N * SUBGROUP_MATRIX_N * SUBGROUP_MATRIX_N_SIZE;
@@ -7,11 +7,6 @@ enable f16;
requires packed_4x8_integer_dot_product;
#endif
#ifdef SRC_OVERLAP
#define SRC0 merged_src
#define SRC1 merged_src
#endif
#define DECLARE_BYTE_LOADERS_SRC0
#include "common_decls.tmpl"
@@ -40,22 +35,17 @@ struct MulMatParams {
broadcast3: u32
};
#if defined(MMVQ)
@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>;
#ifdef MMVQ
@group(0) @binding(1) var<storage, read_write> src1q: array<q8_1>;
#define DST_BINDING 2
#elif defined(SRC_OVERLAP)
@group(0) @binding(0) var<storage, read_write> merged_src: array<SRC0_TYPE>;
#define DST_BINDING 1
#else
@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>;
@group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>;
#define DST_BINDING 2
#endif
@group(0) @binding(DST_BINDING) var<storage, read_write> dst: array<f32>;
@group(0) @binding(2) var<storage, read_write> dst: array<f32>;
// "mul_mat_vec_acc.tmpl" requires params.k, params.m, params.stride_01
@group(0) @binding(DST_BINDING + 1) var<uniform> params: MulMatParams;
@group(0) @binding(3) var<uniform> params: MulMatParams;
// Flattened as [row][thread] to keep each row's reduction contiguous in memory.
var<workgroup> partial_sums: array<f32, OUTPUTS_PER_WG * WG_SIZE>;
@@ -1,10 +1,3 @@
#ifndef SRC0
#define SRC0 src0
#endif
#ifndef SRC1
#define SRC1 src1
#endif
#ifdef U32_DEQUANT_HELPERS
#define SRC0_TYPE u32
@@ -50,13 +43,13 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
for (var k = thread_id; k < k_vec; k += WG_SIZE) {
var x_vals: array<SRC1_TYPE, NUM_COLS>;
for (var col = 0u;col < NUM_COLS;col += 1) {
x_vals[col] = SRC1[src1_idx_base_vec + col * (params.stride_11 / VEC_SIZE) + k];
x_vals[col] = src1[src1_idx_base_vec + col * (params.stride_11 / VEC_SIZE) + k];
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
let output_row = row_base + row;
if (output_row < params.m) {
let src0_idx = (src0_batch_offset + output_row * params.stride_01) / VEC_SIZE + k;
let w = SRC0[src0_idx];
let w = src0[src0_idx];
for (var col = 0u;col < NUM_COLS;col += 1) {
acc[col][row] += inner_dot(w, x_vals[col]);
}
@@ -83,7 +76,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -123,8 +116,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -167,8 +160,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -212,8 +205,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -260,8 +253,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -309,7 +302,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -354,7 +347,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -416,10 +409,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 4u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4u] = f32(SRC1[x_base + col * params.stride_11 + 32u + i]);
x_block[col][i + 8u] = f32(SRC1[x_base + col * params.stride_11 + 64u + i]);
x_block[col][i + 12u] = f32(SRC1[x_base + col * params.stride_11 + 96u + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4u] = f32(src1[x_base + col * params.stride_11 + 32u + i]);
x_block[col][i + 8u] = f32(src1[x_base + col * params.stride_11 + 64u + i]);
x_block[col][i + 12u] = f32(src1[x_base + col * params.stride_11 + 96u + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -525,8 +518,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 8u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 8u] = f32(SRC1[x_base + col * params.stride_11 + 32u + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 8u] = f32(src1[x_base + col * params.stride_11 + 32u + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -617,10 +610,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
for (var col = 0u; col < NUM_COLS;col += 1) {
let col_base = x_base + col * params.stride_11;
for (var i = 0u; i < 4u; i++) {
x_block[col][i] = f32(SRC1[col_base + i]);
x_block[col][i + 4u] = f32(SRC1[col_base + 32u + i]);
x_block[col][i + 8u] = f32(SRC1[col_base + 128u + i]);
x_block[col][i + 12u] = f32(SRC1[col_base + 160u + i]);
x_block[col][i] = f32(src1[col_base + i]);
x_block[col][i + 4u] = f32(src1[col_base + 32u + i]);
x_block[col][i + 8u] = f32(src1[col_base + 128u + i]);
x_block[col][i + 12u] = f32(src1[col_base + 160u + i]);
}
}
@@ -720,10 +713,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
for (var col = 0u; col < NUM_COLS;col += 1) {
let col_base = x_base + col * params.stride_11;
for (var i = 0u; i < 4u; i++) {
x_block[col][i] = f32(SRC1[col_base + i]);
x_block[col][i + 4u] = f32(SRC1[col_base + 32u + i]);
x_block[col][i + 8u] = f32(SRC1[col_base + 128u + i]);
x_block[col][i + 12u] = f32(SRC1[col_base + 160u + i]);
x_block[col][i] = f32(src1[col_base + i]);
x_block[col][i + 4u] = f32(src1[col_base + 32u + i]);
x_block[col][i + 8u] = f32(src1[col_base + 128u + i]);
x_block[col][i + 12u] = f32(src1[col_base + 160u + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -830,10 +823,10 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
for (var col = 0u; col < NUM_COLS;col += 1) {
let col_base = x_base + col * params.stride_11;
for (var l = 0u; l < 4u; l++) {
x_block[col][l] = f32(SRC1[col_base + l]);
x_block[col][l + 4u] = f32(SRC1[col_base + 32u + l]);
x_block[col][l + 8u] = f32(SRC1[col_base + 64u + l]);
x_block[col][l + 12u] = f32(SRC1[col_base + 96u + l]);
x_block[col][l] = f32(src1[col_base + l]);
x_block[col][l + 4u] = f32(src1[col_base + 32u + l]);
x_block[col][l + 8u] = f32(src1[col_base + 64u + l]);
x_block[col][l + 12u] = f32(src1[col_base + 96u + l]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -906,7 +899,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -967,7 +960,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1046,7 +1039,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1108,7 +1101,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1175,7 +1168,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1241,7 +1234,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1309,7 +1302,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1374,8 +1367,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4u] = f32(SRC1[x_base + col * params.stride_11 + i + 16u]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4u] = f32(src1[x_base + col * params.stride_11 + i + 16u]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1425,7 +1418,7 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, 16>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < 16u; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1483,8 +1476,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(SRC1[x_base + col * params.stride_11 + i + 16]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 4] = f32(src1[x_base + col * params.stride_11 + i + 16]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -1528,8 +1521,8 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src
var x_block: array<array<f32, ELEMS_PER_THREAD>, NUM_COLS>;
for (var col = 0u; col < NUM_COLS;col += 1) {
for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) {
x_block[col][i] = f32(SRC1[x_base + col * params.stride_11 + i]);
x_block[col][i + 8] = f32(SRC1[x_base + col * params.stride_11 + i + 8]);
x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]);
x_block[col][i + 8] = f32(src1[x_base + col * params.stride_11 + i + 8]);
}
}
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
+1 -17
View File
@@ -4200,7 +4200,7 @@ static struct ggml_tensor * ggml_rope_impl(
struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
int32_t params[16] = { /*n_past*/ 0, n_dims, mode, /*n_ctx*/ 0, n_ctx_orig };
int32_t params[15] = { /*n_past*/ 0, n_dims, mode, /*n_ctx*/ 0, n_ctx_orig };
memcpy(params + 5, &freq_base, sizeof(float));
memcpy(params + 6, &freq_scale, sizeof(float));
memcpy(params + 7, &ext_factor, sizeof(float));
@@ -4212,8 +4212,6 @@ static struct ggml_tensor * ggml_rope_impl(
} else {
memset(params + 11, 0, sizeof(int32_t) * GGML_MROPE_SECTIONS);
}
params[15] = 0; // n_offs, set via ggml_rope_set_offset()
ggml_set_op_params(result, params, sizeof(params));
result->op = GGML_OP_ROPE;
@@ -4424,20 +4422,6 @@ struct ggml_tensor * ggml_rope_multi_back(
result->op = GGML_OP_ROPE_BACK;
return result;
}
struct ggml_tensor * ggml_rope_set_offset(
struct ggml_tensor * a,
int n_offs) {
GGML_ASSERT(a->op == GGML_OP_ROPE || a->op == GGML_OP_ROPE_BACK);
GGML_ASSERT(n_offs >= 0);
const int32_t mode = ggml_get_op_params_i32(a, 2);
GGML_ASSERT(mode != GGML_ROPE_TYPE_VISION);
ggml_set_op_params_i32(a, 15, n_offs);
return a;
}
// ggml_clamp
struct ggml_tensor * ggml_clamp(
-31
View File
@@ -162,7 +162,6 @@ class Keys:
TARGET_LAYERS = "{arch}.target_layers"
TARGET_HIDDEN_SIZE = "{arch}.target_hidden_size"
BLOCK_SIZE = "{arch}.block_size"
SAMPLE_FROM_ANCHOR = "{arch}.sample_from_anchor"
NORM_BEFORE_RESIDUAL = "{arch}.norm_before_residual"
NORM_BEFORE_FC = "{arch}.norm_before_fc"
@@ -208,7 +207,6 @@ class Keys:
SHARED_KV_LAYERS = "{arch}.attention.shared_kv_layers"
SLIDING_WINDOW_PATTERN = "{arch}.attention.sliding_window_pattern"
TEMPERATURE_SCALE = "{arch}.attention.temperature_scale"
ROPE_PATTERN = "{arch}.attention.rope_pattern"
class Indexer:
HEAD_COUNT = "{arch}.attention.indexer.head_count"
@@ -550,7 +548,6 @@ class MODEL_ARCH(IntEnum):
GRANITE_MOE = auto()
GRANITE_HYBRID = auto()
GRANITE_SWITCH = auto()
GRANITE_SWA = auto()
CHAMELEON = auto()
WAVTOKENIZER_DEC = auto()
PLM = auto()
@@ -1267,7 +1264,6 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
MODEL_ARCH.GRANITE_MOE: "granitemoe",
MODEL_ARCH.GRANITE_HYBRID: "granitehybrid",
MODEL_ARCH.GRANITE_SWITCH: "graniteswitch",
MODEL_ARCH.GRANITE_SWA: "granite_swa",
MODEL_ARCH.CHAMELEON: "chameleon",
MODEL_ARCH.WAVTOKENIZER_DEC: "wavtokenizer-dec",
MODEL_ARCH.PLM: "plm",
@@ -4155,31 +4151,6 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_DOWN,
MODEL_TENSOR.FFN_UP,
],
MODEL_ARCH.GRANITE_SWA: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT_NORM,
MODEL_TENSOR.OUTPUT,
MODEL_TENSOR.ATTN_NORM,
MODEL_TENSOR.ATTN_Q,
MODEL_TENSOR.ATTN_K,
MODEL_TENSOR.ATTN_V,
MODEL_TENSOR.ATTN_OUT,
MODEL_TENSOR.ATTN_SINKS,
MODEL_TENSOR.ROPE_FREQS,
MODEL_TENSOR.FFN_NORM,
MODEL_TENSOR.FFN_GATE,
MODEL_TENSOR.FFN_DOWN,
MODEL_TENSOR.FFN_UP,
# MoE (GraniteMoeSWA)
MODEL_TENSOR.FFN_GATE_INP,
MODEL_TENSOR.FFN_GATE_EXP,
MODEL_TENSOR.FFN_GATE_UP_EXP,
MODEL_TENSOR.FFN_DOWN_EXP,
MODEL_TENSOR.FFN_UP_EXP,
# Shared expert - gate+up kept fused in FFN_UP_SHEXP (LLM_FFN_SWIGLU)
MODEL_TENSOR.FFN_UP_SHEXP,
MODEL_TENSOR.FFN_DOWN_SHEXP,
],
MODEL_ARCH.CHAMELEON: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT_NORM,
@@ -4848,7 +4819,6 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
],
MODEL_ARCH.DFLASH: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT,
MODEL_TENSOR.OUTPUT_NORM,
MODEL_TENSOR.ATTN_NORM,
MODEL_TENSOR.ATTN_Q,
@@ -4888,7 +4858,6 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_UP_SHEXP,
MODEL_TENSOR.FC,
MODEL_TENSOR.ENC_OUTPUT_NORM,
MODEL_TENSOR.D2T,
# optional DSpark heads
MODEL_TENSOR.DSPARK_MARKOV_W1,
MODEL_TENSOR.DSPARK_MARKOV_W2,
-14
View File
@@ -32,10 +32,6 @@ from gguf.constants import (
GGUFEndian,
)
# limits mirroring ggml/src/gguf.cpp (not part of gguf.h)
GGUF_MAX_STRING_LENGTH = 1024 * 1024 * 1024
GGUF_MAX_ARRAY_ELEMENTS = 1024 * 1024 * 1024
logger = logging.getLogger(__name__)
READER_SUPPORTED_VERSIONS = [2, GGUF_VERSION]
@@ -171,10 +167,6 @@ class GGUFReader:
offs += self._push_field(ReaderField(offs, 'GGUF.tensor_count', [temp_counts[:1]], [0], [GGUFValueType.UINT64]))
offs += self._push_field(ReaderField(offs, 'GGUF.kv_count', [temp_counts[1:]], [0], [GGUFValueType.UINT64]))
tensor_count, kv_count = temp_counts
if tensor_count > GGUF_MAX_ARRAY_ELEMENTS:
raise ValueError(f'Tensor count {tensor_count} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
if kv_count > GGUF_MAX_ARRAY_ELEMENTS:
raise ValueError(f'KV count {kv_count} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
offs = self._build_fields(offs, kv_count)
# Build Tensor Info Fields
@@ -225,10 +217,6 @@ class GGUFReader:
def _get_str(self, offset: int) -> tuple[npt.NDArray[np.uint64], npt.NDArray[np.uint8]]:
slen = self._get(offset, np.uint64)
if int(slen[0]) > GGUF_MAX_STRING_LENGTH:
raise ValueError(f'String length {int(slen[0])} exceeds maximum {GGUF_MAX_STRING_LENGTH}')
if offset + 8 + int(slen[0]) > self.data.nbytes:
raise ValueError(f'String length {int(slen[0])} exceeds remaining file size {self.data.nbytes - offset - 8}')
return slen, self._get(offset + 8, np.uint8, slen[0])
def _get_field_parts(
@@ -253,8 +241,6 @@ class GGUFReader:
raw_itype = self._get(offs, np.uint32)
offs += int(raw_itype.nbytes)
alen = self._get(offs, np.uint64)
if int(alen[0]) > GGUF_MAX_ARRAY_ELEMENTS:
raise ValueError(f'Array length {int(alen[0])} exceeds maximum {GGUF_MAX_ARRAY_ELEMENTS}')
offs += int(alen.nbytes)
aparts: list[npt.NDArray[Any]] = [raw_itype, alen]
data_idxs: list[int] = []
-6
View File
@@ -824,9 +824,6 @@ class GGUFWriter:
else:
self.add_array(key, value)
def add_rope_pattern(self, value: Sequence[bool]) -> None:
self.add_array(Keys.Attention.ROPE_PATTERN.format(arch=self.arch), value)
def add_dense_features_dims(self, dense:str, in_f:int, out_f:int) -> None:
self.add_uint32(Keys.LLM.DENSE_FEAT_IN_SIZE.format(arch=self.arch, dense=dense), in_f)
self.add_uint32(Keys.LLM.DENSE_FEAT_OUT_SIZE.format(arch=self.arch, dense=dense), out_f)
@@ -984,9 +981,6 @@ class GGUFWriter:
def add_block_size(self, value: int) -> None:
self.add_uint32(Keys.LLM.BLOCK_SIZE.format(arch=self.arch), value)
def add_sample_from_anchor(self, value: bool) -> None:
self.add_bool(Keys.LLM.SAMPLE_FROM_ANCHOR.format(arch=self.arch), value)
def add_target_layers(self, value: Sequence[int]) -> None:
self.add_array(Keys.LLM.TARGET_LAYERS.format(arch=self.arch), value)
+2 -3
View File
@@ -76,14 +76,14 @@ class TensorNameMap:
# Output
MODEL_TENSOR.OUTPUT: (
"embed_out", # gptneox
"lm_head", # gpt2 mpt falcon llama-hf baichuan qwen mamba dbrx jais nemotron exaone olmoe olmo2 phimoe plamo2 llama4
"lm_head", # gpt2 mpt falcon llama-hf baichuan qwen mamba dbrx jais nemotron exaone olmoe olmo2 phimoe plamo2
"output", # llama-pth bloom internlm2
"word_embeddings_for_head", # persimmon
"lm_head.linear", # phi2
"output_layer", # chatglm
"head", # rwkv
"head.out", # wavtokenizer
"model.lm_head", # dflash
"lm_head", # llama4
"model.transformer.ff_out", # llada
"head.decoder", # modern-bert
),
@@ -458,7 +458,6 @@ class TensorNameMap:
"transformer.decoder_layer.{bid}.router", # Grok
"transformer.blocks.{bid}.ffn.router.layer", # dbrx
"model.layers.{bid}.block_sparse_moe.router.layer", # granitemoe
"model.layers.{bid}.block_sparse_moe.router", # granite_swa
"model.layers.{bid}.feed_forward.router", # llama4 jamba
"encoder.layers.{bid}.mlp.router.layer", # nomic-bert-moe
"model.layers.{bid}.mlp.router", # openai-moe
+3 -38
View File
@@ -4,10 +4,7 @@
# Usage: make-release-checks.sh [--dry-run]
# --dry-run: warn on failures instead of aborting
#
# Env (when running in GitHub Actions):
# GH_TOKEN, GITHUB_REPOSITORY, GITHUB_OUTPUT
# RELEASE_BRANCH: when set, HEAD must belong to origin/RELEASE_BRANCH and must
# not be older than 3 days from the branch HEAD (skipped when unset)
# Env (when running in GitHub Actions): GH_TOKEN, GITHUB_REPOSITORY, GITHUB_OUTPUT
set -euo pipefail
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
@@ -31,39 +28,6 @@ if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
echo "version=${VERSION}" >> "$GITHUB_OUTPUT"
fi
SHA=$(git rev-parse HEAD)
echo "Checking that commit ${SHA} belongs to the release branch..."
if [[ -z "${RELEASE_BRANCH:-}" ]]; then
echo "Warning: RELEASE_BRANCH not set - skipping commit check (local run)"
else
TIP="origin/${RELEASE_BRANCH}"
COMMIT_ERR=""
if ! git rev-parse --verify "${TIP}" >/dev/null 2>&1; then
COMMIT_ERR="branch ${RELEASE_BRANCH} not found on remote"
elif ! git merge-base --is-ancestor "${SHA}" "${TIP}"; then
COMMIT_ERR="commit ${SHA} is not part of branch ${RELEASE_BRANCH}"
else
COMMIT_TS=$(git show -s --format=%ct "${SHA}")
TIP_TS=$(git show -s --format=%ct "${TIP}")
AGE_DAYS=$(( (TIP_TS - COMMIT_TS) / 86400 ))
if (( TIP_TS - COMMIT_TS > 3 * 86400 )); then
COMMIT_ERR="commit ${SHA} is ${AGE_DAYS} day(s) older than the HEAD of ${RELEASE_BRANCH} (max: 3)"
fi
fi
if [[ -n "${COMMIT_ERR}" ]]; then
if [[ "$DRY_RUN" == "true" ]]; then
echo "Warning: ${COMMIT_ERR} (dry run, continuing)."
CHECKS_PASSED=false
else
echo "Error: ${COMMIT_ERR}"
exit 1
fi
else
echo "Commit ${SHA} is on branch ${RELEASE_BRANCH} and within 3 days of its HEAD - OK"
fi
fi
echo "Checking that tag ${VERSION} does not already exist..."
if git ls-remote --tags origin "${VERSION}" | grep -q "${VERSION}"; then
echo "Error: tag ${VERSION} already exists on remote"
@@ -71,11 +35,12 @@ if git ls-remote --tags origin "${VERSION}" | grep -q "${VERSION}"; then
fi
echo "Tag ${VERSION} does not exist on remote - OK"
SHA=$(git rev-parse HEAD)
echo "Checking release.yml status for commit ${SHA}..."
if [[ -z "${GITHUB_REPOSITORY:-}" ]]; then
echo "Warning: GITHUB_REPOSITORY not set - skipping CI check (local run)"
else
RUNS=$(gh api "repos/${GITHUB_REPOSITORY}/actions/workflows/release.yml/runs?per_page=100" \
RUNS=$(gh api "repos/${GITHUB_REPOSITORY}/actions/workflows/release.yml/runs" \
--jq "[.workflow_runs[] | select(.head_sha == \"${SHA}\" and .conclusion == \"success\")] | length")
if [[ "$RUNS" -eq 0 ]]; then
if [[ "$DRY_RUN" == "true" ]]; then
-87
View File
@@ -1,87 +0,0 @@
#!/bin/bash
# Generate the description of a release: the previous release version, the
# change log and the link to the nightly release corresponding to the commit being released.
#
# Usage: make-release-desc.sh <version>
# <version>: current release version (v<maj>.<min>.<pat>, the leading v is optional)
#
# The previous version is the highest plain semver tag (v<maj>.<min>.<pat>)
# strictly below <version>. The change log lists all commits between the
# previous version tag and the release commit, one line per commit.
#
# The release commit is the commit <version> points at when the tag exists,
# HEAD otherwise. The nightly release is the b* tag pointing at that commit
# (release.yml tags the same commit); the link is only generated when that
# tag exists.
#
# Env (when running in GitHub Actions):
# GITHUB_OUTPUT: previous_tag, changelog_title, changelog and nightly are written here
# GITHUB_REPOSITORY: owner/repo, used to build the nightly release URL (skipped when unset)
set -euo pipefail
if [[ $# -ne 1 ]]; then
echo "Usage: $(basename "$0") <version>"
exit 1
fi
VERSION="$1"
# Accept the version with or without the leading v, reject anything else
if [[ "${VERSION}" =~ ^[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
VERSION="v${VERSION}"
elif [[ ! "${VERSION}" =~ ^v[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
echo "Error: invalid version '${VERSION}' (expected v<maj>.<min>.<pat>)"
exit 1
fi
# Make sure all remote tags are available locally (skipped on local runs without origin)
if ! git fetch --tags origin 2>/dev/null; then
echo "Warning: could not fetch tags from origin (local run?)"
fi
# Release commit: the commit <version> points at when the tag exists, HEAD otherwise.
if ! RELEASE_COMMIT="$(git rev-parse -q --verify "refs/tags/${VERSION}^{commit}" 2>/dev/null)"; then
RELEASE_COMMIT="$(git rev-parse HEAD)"
fi
echo "Release commit: $(git rev-parse --short "${RELEASE_COMMIT}")"
PREV="$( { git tag --list; echo "${VERSION}"; } \
| grep -E '^v[0-9]+\.[0-9]+\.[0-9]+$' \
| sort -V \
| awk -v cur="${VERSION}" '$0 == cur { exit } { prev = $0 } END { print prev }')"
if [[ -n "${PREV}" ]]; then
CHANGELOG="$(git log --oneline "${PREV}..${RELEASE_COMMIT}")"
CHANGELOG_TITLE="Change log since ${PREV}"
else
CHANGELOG="(no previous release tag found)"
CHANGELOG_TITLE="Change log"
fi
# Nightly release: the b* tag pointing at the release commit (|| true: no match is not an error)
NIGHTLY_TAG="$(git tag --points-at "${RELEASE_COMMIT}" | grep -E '(^|-)b[0-9]+(-[0-9a-f]{7})?$' | head -n 1 || true)"
NIGHTLY=""
if [[ -n "${NIGHTLY_TAG}" ]]; then
if [[ -n "${GITHUB_REPOSITORY:-}" ]]; then
NIGHTLY_URL="https://github.com/${GITHUB_REPOSITORY}/releases/tag/${NIGHTLY_TAG}"
NIGHTLY="**Nightly build:** [${NIGHTLY_TAG}](${NIGHTLY_URL})"
echo "Nightly release: ${NIGHTLY_URL}"
fi
else
echo "No nightly release found for commit $(git rev-parse --short "${RELEASE_COMMIT}")"
fi
echo "Previous version: ${PREV:-none}"
echo "${CHANGELOG}"
if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
{
echo "previous_tag=${PREV}"
echo "changelog_title=${CHANGELOG_TITLE}"
echo "nightly=${NIGHTLY}"
echo "changelog<<CHANGELOG_EOF"
echo "${CHANGELOG}"
echo "CHANGELOG_EOF"
} >> "${GITHUB_OUTPUT}"
fi
+1 -1
View File
@@ -1 +1 @@
8c63e70982c95ceb862e3a1073a2c1beef75d60a
3834fd814e74e8af277939dabd69ecc780affd21
-89
View File
@@ -7,12 +7,6 @@ import subprocess
HTTPLIB_VERSION = "refs/tags/v0.53.1"
# used by examples/gguf-hash, these repos have no release tag, so we pin a commit
XXHASH_COMMIT = "9f465f1ea932d6ad9a26cd77496311ffa544cd68"
SHA1_COMMIT = "e1e2536fcf6a8f9703be8c85d58724b408552287"
SHA256_COMMIT = "5e637272c13f200872d55ff579f7e2ab6c3f252f"
ROTATE_BITS_COMMIT = "27e784942f67db44abf2115c6638e735b579acd1"
vendor = {
"https://github.com/nlohmann/json/releases/latest/download/json.hpp": "vendor/nlohmann/json.hpp",
"https://github.com/nlohmann/json/releases/latest/download/json_fwd.hpp": "vendor/nlohmann/json_fwd.hpp",
@@ -28,95 +22,12 @@ vendor = {
f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/LICENSE": "vendor/cpp-httplib/LICENSE",
"https://raw.githubusercontent.com/sheredom/subprocess.h/9ce0d701b6fb10f8f8c4445edd31e7c60a1237e3/subprocess.h": "vendor/sheredom/subprocess.h",
f"https://raw.githubusercontent.com/Cyan4973/xxHash/{XXHASH_COMMIT}/xxhash.c": "vendor/hash/xxhash/xxhash.c",
f"https://raw.githubusercontent.com/Cyan4973/xxHash/{XXHASH_COMMIT}/xxhash.h": "vendor/hash/xxhash/xxhash.h",
f"https://raw.githubusercontent.com/Cyan4973/xxHash/{XXHASH_COMMIT}/LICENSE": "vendor/hash/xxhash/LICENSE",
# clibs/sha1 ships no license file, the source header says public domain
f"https://raw.githubusercontent.com/clibs/sha1/{SHA1_COMMIT}/sha1.c": "vendor/hash/sha1/sha1.c",
f"https://raw.githubusercontent.com/clibs/sha1/{SHA1_COMMIT}/sha1.h": "vendor/hash/sha1/sha1.h",
f"https://raw.githubusercontent.com/jb55/sha256.c/{SHA256_COMMIT}/sha256.c": "vendor/hash/sha256/sha256.c",
f"https://raw.githubusercontent.com/jb55/sha256.c/{SHA256_COMMIT}/sha256.h": "vendor/hash/sha256/sha256.h",
f"https://raw.githubusercontent.com/jb55/sha256.c/{SHA256_COMMIT}/LICENSE": "vendor/hash/sha256/LICENSE",
f"https://raw.githubusercontent.com/jb55/rotate-bits.h/{ROTATE_BITS_COMMIT}/rotate-bits.h": "vendor/hash/rotate-bits/rotate-bits.h",
f"https://raw.githubusercontent.com/jb55/rotate-bits.h/{ROTATE_BITS_COMMIT}/LICENSE.md": "vendor/hash/rotate-bits/LICENSE.md",
}
# local changes kept on top of the upstream sources
patches = {
"vendor/hash/xxhash/xxhash.h": [(
'#if defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 201112L) /* >= C11 */\n',
'/* Windows SDK under 10.0.22000 is missing stdalign.h so we add a check\n'
' before allowing the windows compiler to use the C11 form.\n'
' Reference: https://github.com/Cyan4973/xxHash/issues/955 */\n'
'#if defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 201112L) \\\n'
' && (defined(_MSC_VER) && (_MSC_VER >= 1000) || !defined(_MSC_VER)) /* >= C11 */\n'
)],
# sha1 exports a bare "SHA1" symbol, which clashes with the boringssl one at link time.
# we compile it as C++ (see vendor/hash/CMakeLists.txt) and put it in a namespace.
"vendor/hash/sha1/sha1.h": [
(
'#if defined(__cplusplus)\n'
'extern "C" {\n'
'#endif\n',
'namespace vendor_hash {\n'
),
(
'#if defined(__cplusplus)\n'
'}\n'
'#endif\n',
'} // namespace vendor_hash\n'
),
],
"vendor/hash/sha1/sha1.c": [
(
'#include "sha1.h"\n',
'#include "sha1.h"\n'
'\n'
'namespace vendor_hash {\n'
),
(
' SHA1Final((unsigned char *)hash_out, &ctx);\n'
'}\n',
' SHA1Final((unsigned char *)hash_out, &ctx);\n'
'}\n'
'\n'
'} // namespace vendor_hash\n'
),
],
# silence a maybe-uninitialized warning
"vendor/hash/sha256/sha256.c": [(
" uint32_t W[16];\n",
" uint32_t W[16] = {0};\n"
)],
}
for url, filename in vendor.items():
print(f"downloading {url} to {filename}") # noqa: NP100
urllib.request.urlretrieve(url, filename)
for filename, replacements in patches.items():
print(f"patching {filename}") # noqa: NP100
with open(filename, "r", encoding="utf-8", newline="") as f:
content = f.read()
for old, new in replacements:
if content.count(old) != 1:
print(f"Error: cannot apply patch on {filename}, upstream code has changed") # noqa: NP100
sys.exit(1)
content = content.replace(old, new)
with open(filename, "w", encoding="utf-8", newline="") as f:
f.write(content)
print("Splitting httplib.h...") # noqa: NP100
try:
subprocess.check_call([
-3
View File
@@ -102,7 +102,6 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
{ LLM_ARCH_GRANITE_MOE, "granitemoe" },
{ LLM_ARCH_GRANITE_HYBRID, "granitehybrid" },
{ LLM_ARCH_GRANITE_SWITCH, "graniteswitch" },
{ LLM_ARCH_GRANITE_SWA, "granite_swa" },
{ LLM_ARCH_CHAMELEON, "chameleon" },
{ LLM_ARCH_WAVTOKENIZER_DEC, "wavtokenizer-dec" },
{ LLM_ARCH_PLM, "plm" },
@@ -262,8 +261,6 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
{ LLM_KV_ATTENTION_RELATIVE_BUCKETS_COUNT, "%s.attention.relative_buckets_count" },
{ LLM_KV_ATTENTION_SLIDING_WINDOW, "%s.attention.sliding_window" },
{ LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, "%s.attention.sliding_window_pattern" },
{ LLM_KV_ATTENTION_ROPE_PATTERN, "%s.attention.rope_pattern" },
{ LLM_KV_ATTENTION_SCALE, "%s.attention.scale" },
{ LLM_KV_ATTENTION_OUTPUT_SCALE, "%s.attention.output_scale" },
{ LLM_KV_ATTENTION_VALUE_SCALE, "%s.attention.value_scale" },
-3
View File
@@ -107,7 +107,6 @@ enum llm_arch {
LLM_ARCH_GRANITE_MOE,
LLM_ARCH_GRANITE_HYBRID,
LLM_ARCH_GRANITE_SWITCH,
LLM_ARCH_GRANITE_SWA,
LLM_ARCH_CHAMELEON,
LLM_ARCH_WAVTOKENIZER_DEC,
LLM_ARCH_PLM,
@@ -268,8 +267,6 @@ enum llm_kv {
LLM_KV_ATTENTION_SLIDING_WINDOW,
LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN,
LLM_KV_ATTENTION_SCALE,
LLM_KV_ATTENTION_ROPE_PATTERN,
LLM_KV_ATTENTION_OUTPUT_SCALE,
LLM_KV_ATTENTION_VALUE_SCALE,
LLM_KV_ATTENTION_TEMPERATURE_LENGTH,
+1 -5
View File
@@ -291,11 +291,7 @@ bool llama_hparams::has_rope(uint32_t il) const {
return false;
}
if (il < n_layer_all) {
return rope_pattern[il] != 0;
}
GGML_ABORT("%s: il (%u) out of bounds (n_layer_all: %u)\n", __func__, il, n_layer_all);
return true;
}
uint32_t llama_hparams::n_layer() const {

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