mirror of
https://github.com/leejet/stable-diffusion.cpp.git
synced 2026-06-25 07:36:38 +00:00
Compare commits
No commits in common. "18fbb4cdfb3bfd4b111e59db86393470a5faded0" and "22d9a8374b252a8181283519b8b820e93c4b2bd0" have entirely different histories.
18fbb4cdfb
...
22d9a8374b
229
.github/workflows/build.yml
vendored
229
.github/workflows/build.yml
vendored
@ -444,98 +444,12 @@ jobs:
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip
|
||||
|
||||
windows-latest-rocm:
|
||||
runs-on: windows-2022
|
||||
|
||||
env:
|
||||
ROCM_VERSION: "7.12.0"
|
||||
GPU_TARGETS: "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201"
|
||||
|
||||
steps:
|
||||
- uses: actions/checkout@v3
|
||||
with:
|
||||
submodules: recursive
|
||||
|
||||
- name: Cache ROCm Installation
|
||||
id: cache-rocm
|
||||
uses: actions/cache@v4
|
||||
with:
|
||||
path: C:\TheRock\build
|
||||
key: rocm-${{ env.ROCM_VERSION }}-gfx1151-${{ runner.os }}
|
||||
|
||||
- name: ccache
|
||||
uses: ggml-org/ccache-action@v1.2.16
|
||||
with:
|
||||
key: windows-latest-rocm-${{ env.ROCM_VERSION }}-x64
|
||||
evict-old-files: 1d
|
||||
|
||||
- name: Install ROCm
|
||||
if: steps.cache-rocm.outputs.cache-hit != 'true'
|
||||
run: |
|
||||
$ErrorActionPreference = "Stop"
|
||||
write-host "Downloading AMD ROCm ${{ env.ROCM_VERSION }} tarball"
|
||||
Invoke-WebRequest -Uri "https://repo.amd.com/rocm/tarball/therock-dist-windows-gfx1151-${{ env.ROCM_VERSION }}.tar.gz" -OutFile "${env:RUNNER_TEMP}\rocm.tar.gz"
|
||||
write-host "Extracting ROCm tarball"
|
||||
mkdir C:\TheRock\build -Force
|
||||
tar -xzf "${env:RUNNER_TEMP}\rocm.tar.gz" -C C:\TheRock\build --strip-components=1
|
||||
write-host "Completed ROCm extraction"
|
||||
|
||||
- name: Setup ROCm Environment
|
||||
run: |
|
||||
$rocmPath = "C:\TheRock\build"
|
||||
echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
|
||||
echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
|
||||
echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
|
||||
echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
|
||||
echo "$rocmPath\bin" >> $env:GITHUB_PATH
|
||||
echo "$rocmPath\lib\llvm\bin" >> $env:GITHUB_PATH
|
||||
|
||||
- name: Build
|
||||
run: |
|
||||
mkdir build
|
||||
cd build
|
||||
cmake .. `
|
||||
-G "Unix Makefiles" `
|
||||
-DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
|
||||
-DSD_HIPBLAS=ON `
|
||||
-DSD_BUILD_SHARED_LIBS=ON `
|
||||
-DGGML_NATIVE=OFF `
|
||||
-DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
|
||||
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
|
||||
-DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
|
||||
-DHIP_PATH="${env:HIP_PATH}" `
|
||||
-DCMAKE_BUILD_TYPE=Release `
|
||||
-DGPU_TARGETS="${{ env.GPU_TARGETS }}"
|
||||
cmake --build . --config Release --parallel ${env:NUMBER_OF_PROCESSORS}
|
||||
|
||||
- name: Get commit hash
|
||||
id: commit
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: pr-mpt/actions-commit-hash@v2
|
||||
|
||||
- name: Pack artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
run: |
|
||||
cp "${env:HIP_PATH}\bin\hipblas.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\libhipblaslt.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\rocblas.dll" "build\bin\"
|
||||
7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip .\build\bin\*
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip
|
||||
|
||||
windows-latest-cmake-hip:
|
||||
runs-on: windows-2022
|
||||
|
||||
env:
|
||||
HIPSDK_INSTALLER_VERSION: "26.Q1"
|
||||
ROCM_VERSION: "7.1.1"
|
||||
GPU_TARGETS: "gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032"
|
||||
HIPSDK_INSTALLER_VERSION: "25.Q3"
|
||||
GPU_TARGETS: "gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032"
|
||||
|
||||
steps:
|
||||
- uses: actions/checkout@v3
|
||||
@ -570,7 +484,7 @@ jobs:
|
||||
run: |
|
||||
$ErrorActionPreference = "Stop"
|
||||
write-host "Downloading AMD HIP SDK Installer"
|
||||
Invoke-WebRequest -Uri "https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ env.HIPSDK_INSTALLER_VERSION }}-Win11-For-HIP.exe" -OutFile "${env:RUNNER_TEMP}\rocm-install.exe"
|
||||
Invoke-WebRequest -Uri "https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ env.HIPSDK_INSTALLER_VERSION }}-WinSvr2022-For-HIP.exe" -OutFile "${env:RUNNER_TEMP}\rocm-install.exe"
|
||||
write-host "Installing AMD HIP SDK"
|
||||
$proc = Start-Process "${env:RUNNER_TEMP}\rocm-install.exe" -ArgumentList '-install' -NoNewWindow -PassThru
|
||||
$completed = $proc.WaitForExit(600000)
|
||||
@ -623,38 +537,32 @@ jobs:
|
||||
run: |
|
||||
md "build\bin\rocblas\library\"
|
||||
md "build\bin\hipblaslt\library"
|
||||
cp "${env:HIP_PATH}\bin\libhipblas.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\libhipblaslt.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\hipblas.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\hipblaslt.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\rocblas.dll" "build\bin\"
|
||||
cp "${env:HIP_PATH}\bin\rocblas\library\*" "build\bin\rocblas\library\"
|
||||
cp "${env:HIP_PATH}\bin\hipblaslt\library\*" "build\bin\hipblaslt\library\"
|
||||
7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip .\build\bin\*
|
||||
7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-x64.zip .\build\bin\*
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-x64.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-x64.zip
|
||||
|
||||
ubuntu-latest-rocm:
|
||||
runs-on: ubuntu-24.04
|
||||
runs-on: ubuntu-latest
|
||||
container: rocm/dev-ubuntu-24.04:7.2
|
||||
|
||||
env:
|
||||
ROCM_VERSION: "7.2"
|
||||
UBUNTU_VERSION: "24.04"
|
||||
|
||||
strategy:
|
||||
matrix:
|
||||
include:
|
||||
- ROCM_VERSION: "7.2.1"
|
||||
gpu_targets: "gfx908;gfx90a;gfx942;gfx1030;gfx1031;gfx1032;gfx1100;gfx1101;gfx1102;gfx1151;gfx1150;gfx1200;gfx1201"
|
||||
build: 'x64'
|
||||
- ROCM_VERSION: "7.12.0"
|
||||
gpu_targets: "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201"
|
||||
build: x64
|
||||
GPU_TARGETS: "gfx1151;gfx1150;gfx1100;gfx1101;gfx1102;gfx1200;gfx1201"
|
||||
|
||||
steps:
|
||||
- run: apt-get update && apt-get install -y git
|
||||
- name: Clone
|
||||
id: checkout
|
||||
uses: actions/checkout@v6
|
||||
@ -671,38 +579,6 @@ jobs:
|
||||
with:
|
||||
version: 10.15.1
|
||||
|
||||
- name: ccache
|
||||
uses: ggml-org/ccache-action@v1.2.16
|
||||
with:
|
||||
key: ubuntu-rocm-cmake-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
|
||||
evict-old-files: 1d
|
||||
|
||||
- name: Dependencies
|
||||
id: depends
|
||||
run: |
|
||||
sudo apt install -y build-essential cmake wget zip ninja-build
|
||||
|
||||
- name: Setup Legacy ROCm
|
||||
if: matrix.ROCM_VERSION == '7.2.1'
|
||||
id: legacy_env
|
||||
run: |
|
||||
sudo mkdir --parents --mode=0755 /etc/apt/keyrings
|
||||
wget https://repo.radeon.com/rocm/rocm.gpg.key -O - | \
|
||||
gpg --dearmor | sudo tee /etc/apt/keyrings/rocm.gpg > /dev/null
|
||||
|
||||
sudo tee /etc/apt/sources.list.d/rocm.list << EOF
|
||||
deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/${{ matrix.ROCM_VERSION }} noble main
|
||||
EOF
|
||||
|
||||
sudo tee /etc/apt/preferences.d/rocm-pin-600 << EOF
|
||||
Package: *
|
||||
Pin: release o=repo.radeon.com
|
||||
Pin-Priority: 600
|
||||
EOF
|
||||
|
||||
sudo apt update
|
||||
sudo apt-get install -y libssl-dev rocm-hip-sdk
|
||||
|
||||
- name: Free disk space
|
||||
run: |
|
||||
# Remove preinstalled SDKs and caches not needed for this job
|
||||
@ -716,17 +592,51 @@ jobs:
|
||||
sudo rm -rf /var/lib/apt/lists/* || true
|
||||
sudo apt clean
|
||||
|
||||
- name: Setup TheRock
|
||||
if: matrix.ROCM_VERSION != '7.2.1'
|
||||
id: therock_env
|
||||
- name: Dependencies
|
||||
id: depends
|
||||
run: |
|
||||
wget https://repo.amd.com/rocm/tarball/therock-dist-linux-gfx1151-${{ matrix.ROCM_VERSION }}.tar.gz
|
||||
mkdir install
|
||||
tar -xf *.tar.gz -C install
|
||||
export ROCM_PATH=$(pwd)/install
|
||||
echo ROCM_PATH=$ROCM_PATH >> $GITHUB_ENV
|
||||
echo PATH=$PATH:$ROCM_PATH/bin >> $GITHUB_ENV
|
||||
echo LD_LIBRARY_PATH=$ROCM_PATH/lib:$ROCM_PATH/llvm/lib:$ROCM_PATH/lib/rocprofiler-systems >> $GITHUB_ENV
|
||||
sudo apt-get update
|
||||
sudo apt install -y \
|
||||
cmake \
|
||||
hip-dev \
|
||||
hipblas-dev \
|
||||
ninja-build \
|
||||
rocm-dev \
|
||||
zip
|
||||
# Clean apt caches to recover disk space
|
||||
sudo apt clean
|
||||
sudo rm -rf /var/lib/apt/lists/* || true
|
||||
|
||||
- name: Setup ROCm Environment
|
||||
run: |
|
||||
# Add ROCm to PATH for current session
|
||||
echo "/opt/rocm/bin" >> $GITHUB_PATH
|
||||
|
||||
# Build regex pattern from ${{ env.GPU_TARGETS }} (match target as substring)
|
||||
TARGET_REGEX="($(printf '%s' "${{ env.GPU_TARGETS }}" | sed 's/;/|/g'))"
|
||||
|
||||
# Remove library files for architectures we're not building for to save disk space
|
||||
echo "Cleaning up unneeded architecture files..."
|
||||
cd /opt/rocm/lib/rocblas/library
|
||||
# Keep only our target architectures
|
||||
for file in *; do
|
||||
if printf '%s' "$file" | grep -q 'gfx'; then
|
||||
if ! printf '%s' "$file" | grep -Eq "$TARGET_REGEX"; then
|
||||
echo "Removing $file" &&
|
||||
sudo rm -f "$file";
|
||||
fi
|
||||
fi
|
||||
done
|
||||
|
||||
cd /opt/rocm/lib/hipblaslt/library
|
||||
for file in *; do
|
||||
if printf '%s' "$file" | grep -q 'gfx'; then
|
||||
if ! printf '%s' "$file" | grep -Eq "$TARGET_REGEX"; then
|
||||
echo "Removing $file" &&
|
||||
sudo rm -f "$file";
|
||||
fi
|
||||
fi
|
||||
done
|
||||
|
||||
- name: Build
|
||||
id: cmake_build
|
||||
@ -734,12 +644,12 @@ jobs:
|
||||
mkdir build
|
||||
cd build
|
||||
cmake .. -G Ninja \
|
||||
-DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
|
||||
-DCMAKE_HIP_FLAGS="-mllvm --amdgpu-unroll-threshold-local=600" \
|
||||
-DCMAKE_CXX_COMPILER=amdclang++ \
|
||||
-DCMAKE_C_COMPILER=amdclang \
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-DSD_HIPBLAS=ON \
|
||||
-DHIP_PLATFORM=amd \
|
||||
-DGPU_TARGETS="${{ matrix.gpu_targets }}" \
|
||||
-DGPU_TARGETS="${{ env.GPU_TARGETS }}" \
|
||||
-DAMDGPU_TARGETS="${{ env.GPU_TARGETS }}" \
|
||||
-DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
|
||||
-DCMAKE_POSITION_INDEPENDENT_CODE=ON \
|
||||
-DSD_BUILD_SHARED_LIBS=ON
|
||||
@ -758,6 +668,16 @@ jobs:
|
||||
cp ggml/LICENSE ./build/bin/ggml.txt
|
||||
cp LICENSE ./build/bin/stable-diffusion.cpp.txt
|
||||
|
||||
# Move ROCm runtime libraries (to avoid double space consumption)
|
||||
sudo mv /opt/rocm/lib/librocsparse.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/libhsa-runtime64.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/libamdhip64.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/libhipblas.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/libhipblaslt.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/librocblas.so* ./build/bin/
|
||||
sudo mv /opt/rocm/lib/rocblas/ ./build/bin/
|
||||
sudo mv /opt/rocm/lib/hipblaslt/ ./build/bin/
|
||||
|
||||
- name: Fetch system info
|
||||
id: system-info
|
||||
run: |
|
||||
@ -772,15 +692,15 @@ jobs:
|
||||
run: |
|
||||
cp ggml/LICENSE ./build/bin/ggml.txt
|
||||
cp LICENSE ./build/bin/stable-diffusion.cpp.txt
|
||||
zip -y -r sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm-${{ matrix.ROCM_VERSION }}.zip ./build/bin
|
||||
zip -y -r sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm.zip ./build/bin
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm-${{ matrix.ROCM_VERSION }}.zip
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm-${{ matrix.ROCM_VERSION }}.zip
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-Ubuntu-${{ env.UBUNTU_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-rocm.zip
|
||||
|
||||
release:
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
@ -795,7 +715,6 @@ jobs:
|
||||
- macOS-latest-cmake
|
||||
- windows-latest-cmake
|
||||
- windows-latest-cmake-hip
|
||||
- windows-latest-rocm
|
||||
|
||||
steps:
|
||||
- name: Clone
|
||||
|
||||
@ -133,11 +133,9 @@ API and command-line option may change frequently.***
|
||||
## Performance
|
||||
|
||||
If you want to improve performance or reduce VRAM/RAM usage, please refer to [performance guide](./docs/performance.md).
|
||||
For runtime and parameter backend placement, see the [backend selection guide](./docs/backend.md).
|
||||
|
||||
## More Guides
|
||||
|
||||
- [Backend selection](./docs/backend.md)
|
||||
- [SD1.x/SD2.x/SDXL](./docs/sd.md)
|
||||
- [SD3/SD3.5](./docs/sd3.md)
|
||||
- [FLUX.1-dev/FLUX.1-schnell](./docs/flux.md)
|
||||
|
||||
122
docs/backend.md
122
docs/backend.md
@ -1,122 +0,0 @@
|
||||
# Backend selection
|
||||
|
||||
`stable-diffusion.cpp` has two backend assignments:
|
||||
|
||||
- `--backend` selects the runtime backend used to execute model graphs.
|
||||
- `--params-backend` selects the backend used to allocate model parameters.
|
||||
|
||||
If `--params-backend` is not set, parameters use the same backend as their module runtime backend.
|
||||
|
||||
## Syntax
|
||||
|
||||
A backend assignment can be a single backend name:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend cpu
|
||||
```
|
||||
|
||||
This applies to every module that does not have a more specific assignment.
|
||||
|
||||
Assignments can also target individual modules:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend te=cpu,vae=cuda0,diffusion=vulkan0
|
||||
```
|
||||
|
||||
The same syntax is used for parameter placement:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend cuda0 --params-backend te=cpu,vae=cpu
|
||||
```
|
||||
|
||||
Module names are case-insensitive. Hyphens and underscores in module names are ignored, so `clip_vision`, `clip-vision`, and `clipvision` are equivalent.
|
||||
|
||||
`all=`, `default=`, and `*=` can be used to set the default backend inside a mixed assignment:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend all=cuda0,te=cpu
|
||||
```
|
||||
|
||||
## Modules
|
||||
|
||||
| Module | Purpose | Accepted names |
|
||||
| --- | --- | --- |
|
||||
| `diffusion` | UNet, DiT, MMDiT, Flux, Wan, Qwen Image, and other diffusion models | `diffusion`, `model`, `unet`, `dit` |
|
||||
| `te` | Text encoders and conditioners | `te`, `clip`, `text`, `textencoder`, `textencoders`, `conditioner`, `cond`, `llm`, `t5`, `t5xxl` |
|
||||
| `clip_vision` | CLIP vision encoder | `clip_vision`, `clipvision`, `clip-vision`, `vision` |
|
||||
| `vae` | VAE and TAE | `vae`, `firststage`, `autoencoder`, `tae` |
|
||||
| `controlnet` | ControlNet | `controlnet`, `control` |
|
||||
| `photomaker` | PhotoMaker ID encoder and PhotoMaker LoRA | `photomaker`, `photomakerid`, `pmid`, `photo` |
|
||||
| `upscaler` | ESRGAN upscaler | `upscaler`, `esrgan`, `hires` |
|
||||
|
||||
`te` is the preferred module name for text encoders. `clip` is kept as an accepted alias because many existing commands and model names use CLIP terminology.
|
||||
|
||||
## Backend names
|
||||
|
||||
Backend names are resolved against the GGML backend device list. Matching is case-insensitive and accepts exact names or unique prefixes, so common values include names such as:
|
||||
|
||||
- `cpu`
|
||||
- `cuda0`
|
||||
- `vulkan0`
|
||||
- `metal`
|
||||
|
||||
The special values `auto`, `default`, and an empty backend name select the default backend. The default preference is GPU, then integrated GPU, then CPU.
|
||||
|
||||
The special value `gpu` selects the first GPU backend, falling back to the first integrated GPU backend.
|
||||
|
||||
## Runtime backend vs. parameter backend
|
||||
|
||||
The runtime backend controls where graph execution runs. The parameter backend controls where model weights are allocated.
|
||||
|
||||
For example:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend cuda0 --params-backend cpu
|
||||
```
|
||||
|
||||
This runs all modules on `cuda0`, but stores parameters in CPU RAM. During execution, parameters are moved to the runtime backend as needed.
|
||||
|
||||
Per-module assignments can be mixed:
|
||||
|
||||
```shell
|
||||
sd-cli -m model.safetensors -p "a cat" --backend diffusion=cuda0,te=cpu,vae=cpu --params-backend diffusion=cuda0,te=cpu,vae=cpu
|
||||
```
|
||||
|
||||
This keeps text encoding and VAE execution on CPU while the diffusion model runs on GPU.
|
||||
|
||||
## Backend sharing and lifetime
|
||||
|
||||
Backends are managed by `SDBackendManager`.
|
||||
|
||||
Within one manager, backend instances are cached by resolved backend device name. If multiple modules request the same backend, they share the same `ggml_backend_t`.
|
||||
|
||||
For example:
|
||||
|
||||
```shell
|
||||
--backend te=cpu,vae=cpu
|
||||
```
|
||||
|
||||
uses one shared CPU backend for both `te` and `vae` runtime execution.
|
||||
|
||||
Runtime and parameter assignments also share the same backend cache. If `--backend diffusion=cuda0` and `--params-backend diffusion=cuda0` resolve to the same device, both use the same backend instance.
|
||||
|
||||
`SDBackendManager` owns the backend instances and frees them when the context or upscaler is destroyed. Model runners receive non-owning runtime and parameter backend pointers and do not free them.
|
||||
|
||||
## Compatibility flags
|
||||
|
||||
The older CPU placement flags are still supported:
|
||||
|
||||
- `--clip-on-cpu`
|
||||
- `--vae-on-cpu`
|
||||
- `--control-net-cpu`
|
||||
- `--offload-to-cpu`
|
||||
|
||||
`--clip-on-cpu`, `--vae-on-cpu`, and `--control-net-cpu` affect runtime backend assignment only when `--backend` is not set. They map to `te=cpu`, `vae=cpu`, and `controlnet=cpu`.
|
||||
|
||||
`--offload-to-cpu` affects parameter backend assignment only when `--params-backend` is not set. It is equivalent to:
|
||||
|
||||
```shell
|
||||
--params-backend cpu
|
||||
```
|
||||
|
||||
Explicit `--backend` and `--params-backend` assignments are preferred for new commands.
|
||||
@ -26,12 +26,12 @@ Fortunately, `AMD` provides complete help documentation, you can use the help do
|
||||
|
||||
Then we must set `ROCM` as environment variables before running cmake.
|
||||
|
||||
Usually if you install according to the official tutorial and do not modify the ROCM path, then there is a high probability that it is here `C:\Program Files\AMD\ROCm\7.1.1\bin`
|
||||
Usually if you install according to the official tutorial and do not modify the ROCM path, then there is a high probability that it is here `C:\Program Files\AMD\ROCm\5.5\bin`
|
||||
|
||||
This is what I use to set the clang:
|
||||
```Commandline
|
||||
set CC=C:\Program Files\AMD\ROCm\7.1.1\bin\clang.exe
|
||||
set CXX=C:\Program Files\AMD\ROCm\7.1.1\bin\clang++.exe
|
||||
set CC=C:\Program Files\AMD\ROCm\5.5\bin\clang.exe
|
||||
set CXX=C:\Program Files\AMD\ROCm\5.5\bin\clang++.exe
|
||||
```
|
||||
|
||||
## Ninja
|
||||
@ -46,7 +46,7 @@ set ninja=C:\Program Files\ninja\ninja.exe
|
||||
## Building stable-diffusion.cpp
|
||||
|
||||
The thing different from the regular CPU build is `-DSD_HIPBLAS=ON` ,
|
||||
`-G "Ninja"`, `-DCMAKE_C_COMPILER=clang`, `-DCMAKE_CXX_COMPILER=clang++`, `-DAMDGPU_TARGETS=gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032`
|
||||
`-G "Ninja"`, `-DCMAKE_C_COMPILER=clang`, `-DCMAKE_CXX_COMPILER=clang++`, `-DAMDGPU_TARGETS=gfx1100`
|
||||
|
||||
>**Notice**: check the `clang` and `clang++` information:
|
||||
```Commandline
|
||||
@ -59,29 +59,26 @@ If you see like this, we can continue:
|
||||
clang version 17.0.0 (git@github.amd.com:Compute-Mirrors/llvm-project e3201662d21c48894f2156d302276eb1cf47c7be)
|
||||
Target: x86_64-pc-windows-msvc
|
||||
Thread model: posix
|
||||
InstalledDir: C:\Program Files\AMD\ROCm\7.1.1\bin
|
||||
InstalledDir: C:\Program Files\AMD\ROCm\5.5\bin
|
||||
```
|
||||
|
||||
```
|
||||
clang version 17.0.0 (git@github.amd.com:Compute-Mirrors/llvm-project e3201662d21c48894f2156d302276eb1cf47c7be)
|
||||
Target: x86_64-pc-windows-msvc
|
||||
Thread model: posix
|
||||
InstalledDir: C:\Program Files\AMD\ROCm\7.1.1\bin
|
||||
InstalledDir: C:\Program Files\AMD\ROCm\5.5\bin
|
||||
```
|
||||
|
||||
>**Notice** that the GPU targets are now compatible with multiple GPU architectures (ROCm 7.1.1 targets). You can change them to match your GPU architecture. Click here to see your architecture [LLVM Target](https://rocm.docs.amd.com/en/latest/release/windows_support.html#windows-supported-gpus)
|
||||
>**Notice** that the `gfx1100` is the GPU architecture of my GPU, you can change it to your GPU architecture. Click here to see your architecture [LLVM Target](https://rocm.docs.amd.com/en/latest/release/windows_support.html#windows-supported-gpus)
|
||||
|
||||
Examples:
|
||||
- AMD Radeon™ RX 7900 XTX Graphics: `gfx1100`
|
||||
- AMD Radeon™ RX 7900 XT Graphics: `gfx1101`
|
||||
- AMD Radeon™ RX 7900 GRE Graphics: `gfx1102`
|
||||
My GPU is AMD Radeon™ RX 7900 XTX Graphics, so I set it to `gfx1100`.
|
||||
|
||||
option:
|
||||
|
||||
```commandline
|
||||
mkdir build
|
||||
cd build
|
||||
cmake .. -G "Ninja" -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DSD_HIPBLAS=ON -DCMAKE_BUILD_TYPE=Release -DAMDGPU_TARGETS="gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032"
|
||||
cmake .. -G "Ninja" -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DSD_HIPBLAS=ON -DCMAKE_BUILD_TYPE=Release -DAMDGPU_TARGETS=gfx1100
|
||||
cmake --build . --config Release
|
||||
```
|
||||
|
||||
|
||||
@ -21,7 +21,7 @@ You can run Z-Image with stable-diffusion.cpp on GPUs with 4GB of VRAM — or ev
|
||||
### Z-Image-Turbo
|
||||
|
||||
```
|
||||
.\bin\Release\sd-cli.exe --diffusion-model z_image_turbo-Q3_K.gguf --vae ..\..\ComfyUI\models\vae\ae.sft --llm ..\..\ComfyUI\models\text_encoders\Qwen3-4B-Instruct-2507-Q4_K_M.gguf -p "A cinematic, melancholic photograph of a solitary hooded figure walking through a sprawling, rain-slicked metropolis at night. The city lights are a chaotic blur of neon orange and cool blue, reflecting on the wet asphalt. The scene evokes a sense of being a single component in a vast machine. Superimposed over the image in a sleek, modern, slightly glitched font is the philosophical quote: 'THE CITY IS A CIRCUIT BOARD, AND I AM A BROKEN TRANSISTOR.' -- moody, atmospheric, profound, dark academic" --cfg-scale 1.0 -v --offload-to-cpu --diffusion-fa -H 1024 -W 512 --steps 8
|
||||
.\bin\Release\sd-cli.exe --diffusion-model z_image_turbo-Q3_K.gguf --vae ..\..\ComfyUI\models\vae\ae.sft --llm ..\..\ComfyUI\models\text_encoders\Qwen3-4B-Instruct-2507-Q4_K_M.gguf -p "A cinematic, melancholic photograph of a solitary hooded figure walking through a sprawling, rain-slicked metropolis at night. The city lights are a chaotic blur of neon orange and cool blue, reflecting on the wet asphalt. The scene evokes a sense of being a single component in a vast machine. Superimposed over the image in a sleek, modern, slightly glitched font is the philosophical quote: 'THE CITY IS A CIRCUIT BOARD, AND I AM A BROKEN TRANSISTOR.' -- moody, atmospheric, profound, dark academic" --cfg-scale 1.0 -v --offload-to-cpu --diffusion-fa -H 1024 -W 512
|
||||
```
|
||||
|
||||
<img width="256" alt="z-image example" src="../assets/z_image/q3_K.png" />
|
||||
|
||||
@ -800,9 +800,7 @@ int main(int argc, const char* argv[]) {
|
||||
ctx_params.offload_params_to_cpu,
|
||||
ctx_params.diffusion_conv_direct,
|
||||
ctx_params.n_threads,
|
||||
gen_params.upscale_tile_size,
|
||||
ctx_params.backend.c_str(),
|
||||
ctx_params.params_backend.c_str()));
|
||||
gen_params.upscale_tile_size));
|
||||
|
||||
if (upscaler_ctx == nullptr) {
|
||||
LOG_ERROR("new_upscaler_ctx failed");
|
||||
|
||||
@ -388,14 +388,6 @@ ArgOptions SDContextParams::get_options() {
|
||||
"--upscale-model",
|
||||
"path to esrgan model.",
|
||||
&esrgan_path},
|
||||
{"",
|
||||
"--backend",
|
||||
"runtime backend assignment, e.g. cpu or clip=cpu,vae=cuda0,diffusion=vulkan0",
|
||||
&backend},
|
||||
{"",
|
||||
"--params-backend",
|
||||
"parameter backend assignment, e.g. cpu or diffusion=cpu,clip=cpu",
|
||||
¶ms_backend},
|
||||
};
|
||||
|
||||
options.int_options = {
|
||||
@ -694,8 +686,6 @@ std::string SDContextParams::to_string() const {
|
||||
<< " sampler_rng_type: " << sd_rng_type_name(sampler_rng_type) << ",\n"
|
||||
<< " offload_params_to_cpu: " << (offload_params_to_cpu ? "true" : "false") << ",\n"
|
||||
<< " max_vram: " << max_vram << ",\n"
|
||||
<< " backend: \"" << backend << "\",\n"
|
||||
<< " params_backend: \"" << params_backend << "\",\n"
|
||||
<< " enable_mmap: " << (enable_mmap ? "true" : "false") << ",\n"
|
||||
<< " control_net_cpu: " << (control_net_cpu ? "true" : "false") << ",\n"
|
||||
<< " clip_on_cpu: " << (clip_on_cpu ? "true" : "false") << ",\n"
|
||||
@ -773,8 +763,6 @@ sd_ctx_params_t SDContextParams::to_sd_ctx_params_t(bool vae_decode_only, bool f
|
||||
chroma_t5_mask_pad,
|
||||
qwen_image_zero_cond_t,
|
||||
max_vram,
|
||||
backend.c_str(),
|
||||
params_backend.c_str(),
|
||||
};
|
||||
return sd_ctx_params;
|
||||
}
|
||||
|
||||
@ -112,16 +112,14 @@ struct SDContextParams {
|
||||
rng_type_t sampler_rng_type = RNG_TYPE_COUNT;
|
||||
bool offload_params_to_cpu = false;
|
||||
float max_vram = 0.f;
|
||||
std::string backend;
|
||||
std::string params_backend;
|
||||
bool enable_mmap = false;
|
||||
bool control_net_cpu = false;
|
||||
bool clip_on_cpu = false;
|
||||
bool vae_on_cpu = false;
|
||||
bool flash_attn = false;
|
||||
bool diffusion_flash_attn = false;
|
||||
bool diffusion_conv_direct = false;
|
||||
bool vae_conv_direct = false;
|
||||
bool enable_mmap = false;
|
||||
bool control_net_cpu = false;
|
||||
bool clip_on_cpu = false;
|
||||
bool vae_on_cpu = false;
|
||||
bool flash_attn = false;
|
||||
bool diffusion_flash_attn = false;
|
||||
bool diffusion_conv_direct = false;
|
||||
bool vae_conv_direct = false;
|
||||
|
||||
bool circular = false;
|
||||
bool circular_x = false;
|
||||
|
||||
@ -209,8 +209,6 @@ typedef struct {
|
||||
int chroma_t5_mask_pad;
|
||||
bool qwen_image_zero_cond_t;
|
||||
float max_vram; // GiB budget for graph-cut segmented param offload (0 = disabled, -1 = auto free VRAM minus 1 GiB)
|
||||
const char* backend;
|
||||
const char* params_backend;
|
||||
} sd_ctx_params_t;
|
||||
|
||||
typedef struct {
|
||||
@ -445,9 +443,7 @@ SD_API upscaler_ctx_t* new_upscaler_ctx(const char* esrgan_path,
|
||||
bool offload_params_to_cpu,
|
||||
bool direct,
|
||||
int n_threads,
|
||||
int tile_size,
|
||||
const char* backend,
|
||||
const char* params_backend);
|
||||
int tile_size);
|
||||
SD_API void free_upscaler_ctx(upscaler_ctx_t* upscaler_ctx);
|
||||
|
||||
SD_API sd_image_t upscale(upscaler_ctx_t* upscaler_ctx,
|
||||
|
||||
@ -526,10 +526,10 @@ namespace Anima {
|
||||
AnimaNet net;
|
||||
|
||||
AnimaRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model")
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
int64_t num_layers = 0;
|
||||
std::string layer_tag = prefix + ".net.blocks.";
|
||||
for (const auto& kv : tensor_storage_map) {
|
||||
|
||||
@ -664,13 +664,13 @@ struct AutoEncoderKL : public VAE {
|
||||
AutoEncoderKLModel ae;
|
||||
|
||||
AutoEncoderKL(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
bool decode_only = false,
|
||||
bool use_video_decoder = false,
|
||||
SDVersion version = VERSION_SD1)
|
||||
: decode_only(decode_only), VAE(version, backend, params_backend) {
|
||||
: decode_only(decode_only), VAE(version, backend, offload_params_to_cpu) {
|
||||
if (sd_version_is_sd1(version) || sd_version_is_sd2(version)) {
|
||||
scale_factor = 0.18215f;
|
||||
shift_factor = 0.f;
|
||||
|
||||
@ -469,13 +469,13 @@ struct CLIPTextModelRunner : public GGMLRunner {
|
||||
std::vector<float> attention_mask_vec;
|
||||
|
||||
CLIPTextModelRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
CLIPVersion version = OPENAI_CLIP_VIT_L_14,
|
||||
bool with_final_ln = true,
|
||||
bool force_clip_f32 = false)
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
bool proj_in = false;
|
||||
for (const auto& [name, tensor_storage] : tensor_storage_map) {
|
||||
if (!starts_with(name, prefix)) {
|
||||
|
||||
@ -147,7 +147,7 @@ struct FrozenCLIPEmbedderWithCustomWords : public Conditioner {
|
||||
std::map<std::string, std::pair<int, int>> embedding_pos_map;
|
||||
|
||||
FrozenCLIPEmbedderWithCustomWords(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::map<std::string, std::string>& orig_embedding_map,
|
||||
SDVersion version = VERSION_SD1,
|
||||
@ -161,12 +161,12 @@ struct FrozenCLIPEmbedderWithCustomWords : public Conditioner {
|
||||
}
|
||||
bool force_clip_f32 = !embedding_map.empty();
|
||||
if (sd_version_is_sd1(version)) {
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "cond_stage_model.transformer.text_model", OPENAI_CLIP_VIT_L_14, true, force_clip_f32);
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "cond_stage_model.transformer.text_model", OPENAI_CLIP_VIT_L_14, true, force_clip_f32);
|
||||
} else if (sd_version_is_sd2(version)) {
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "cond_stage_model.transformer.text_model", OPEN_CLIP_VIT_H_14, true, force_clip_f32);
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "cond_stage_model.transformer.text_model", OPEN_CLIP_VIT_H_14, true, force_clip_f32);
|
||||
} else if (sd_version_is_sdxl(version)) {
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "cond_stage_model.transformer.text_model", OPENAI_CLIP_VIT_L_14, false, force_clip_f32);
|
||||
text_model2 = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "cond_stage_model.1.transformer.text_model", OPEN_CLIP_VIT_BIGG_14, false, force_clip_f32);
|
||||
text_model = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "cond_stage_model.transformer.text_model", OPENAI_CLIP_VIT_L_14, false, force_clip_f32);
|
||||
text_model2 = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "cond_stage_model.1.transformer.text_model", OPEN_CLIP_VIT_BIGG_14, false, force_clip_f32);
|
||||
}
|
||||
}
|
||||
|
||||
@ -683,9 +683,9 @@ struct FrozenCLIPVisionEmbedder : public GGMLRunner {
|
||||
CLIPVisionModelProjection vision_model;
|
||||
|
||||
FrozenCLIPVisionEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {})
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
std::string prefix = "cond_stage_model.transformer";
|
||||
bool proj_in = false;
|
||||
for (const auto& [name, tensor_storage] : tensor_storage_map) {
|
||||
@ -742,7 +742,7 @@ struct SD3CLIPEmbedder : public Conditioner {
|
||||
std::shared_ptr<T5Runner> t5;
|
||||
|
||||
SD3CLIPEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {})
|
||||
: clip_g_tokenizer(0) {
|
||||
bool use_clip_l = false;
|
||||
@ -762,13 +762,13 @@ struct SD3CLIPEmbedder : public Conditioner {
|
||||
return;
|
||||
}
|
||||
if (use_clip_l) {
|
||||
clip_l = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "text_encoders.clip_l.transformer.text_model", OPENAI_CLIP_VIT_L_14, false);
|
||||
clip_l = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.clip_l.transformer.text_model", OPENAI_CLIP_VIT_L_14, false);
|
||||
}
|
||||
if (use_clip_g) {
|
||||
clip_g = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "text_encoders.clip_g.transformer.text_model", OPEN_CLIP_VIT_BIGG_14, false);
|
||||
clip_g = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.clip_g.transformer.text_model", OPEN_CLIP_VIT_BIGG_14, false);
|
||||
}
|
||||
if (use_t5) {
|
||||
t5 = std::make_shared<T5Runner>(backend, params_backend, tensor_storage_map, "text_encoders.t5xxl.transformer");
|
||||
t5 = std::make_shared<T5Runner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.t5xxl.transformer");
|
||||
}
|
||||
}
|
||||
|
||||
@ -1110,7 +1110,7 @@ struct FluxCLIPEmbedder : public Conditioner {
|
||||
size_t chunk_len = 256;
|
||||
|
||||
FluxCLIPEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {}) {
|
||||
bool use_clip_l = false;
|
||||
bool use_t5 = false;
|
||||
@ -1128,12 +1128,12 @@ struct FluxCLIPEmbedder : public Conditioner {
|
||||
}
|
||||
|
||||
if (use_clip_l) {
|
||||
clip_l = std::make_shared<CLIPTextModelRunner>(backend, params_backend, tensor_storage_map, "text_encoders.clip_l.transformer.text_model", OPENAI_CLIP_VIT_L_14, true);
|
||||
clip_l = std::make_shared<CLIPTextModelRunner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.clip_l.transformer.text_model", OPENAI_CLIP_VIT_L_14, true);
|
||||
} else {
|
||||
LOG_WARN("clip_l text encoder not found! Prompt adherence might be degraded.");
|
||||
}
|
||||
if (use_t5) {
|
||||
t5 = std::make_shared<T5Runner>(backend, params_backend, tensor_storage_map, "text_encoders.t5xxl.transformer");
|
||||
t5 = std::make_shared<T5Runner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.t5xxl.transformer");
|
||||
} else {
|
||||
LOG_WARN("t5xxl text encoder not found! Prompt adherence might be degraded.");
|
||||
}
|
||||
@ -1364,7 +1364,7 @@ struct T5CLIPEmbedder : public Conditioner {
|
||||
bool is_umt5 = false;
|
||||
|
||||
T5CLIPEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
bool use_mask = false,
|
||||
int mask_pad = 0,
|
||||
@ -1381,7 +1381,7 @@ struct T5CLIPEmbedder : public Conditioner {
|
||||
LOG_WARN("IMPORTANT NOTICE: No text encoders provided, cannot process prompts!");
|
||||
return;
|
||||
} else {
|
||||
t5 = std::make_shared<T5Runner>(backend, params_backend, tensor_storage_map, "text_encoders.t5xxl.transformer", is_umt5);
|
||||
t5 = std::make_shared<T5Runner>(backend, offload_params_to_cpu, tensor_storage_map, "text_encoders.t5xxl.transformer", is_umt5);
|
||||
}
|
||||
}
|
||||
|
||||
@ -1566,12 +1566,12 @@ struct AnimaConditioner : public Conditioner {
|
||||
std::shared_ptr<LLM::LLMRunner> llm;
|
||||
|
||||
AnimaConditioner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {}) {
|
||||
qwen_tokenizer = std::make_shared<Qwen2Tokenizer>();
|
||||
llm = std::make_shared<LLM::LLMRunner>(LLM::LLMArch::QWEN3,
|
||||
backend,
|
||||
params_backend,
|
||||
offload_params_to_cpu,
|
||||
tensor_storage_map,
|
||||
"text_encoders.llm",
|
||||
false);
|
||||
@ -1638,11 +1638,10 @@ struct AnimaConditioner : public Conditioner {
|
||||
for (const auto& item : parsed_attention) {
|
||||
const std::string& curr_text = item.first;
|
||||
float curr_weight = item.second;
|
||||
std::vector<int> curr_tokens = t5_tokenizer.encode(curr_text);
|
||||
std::vector<int> curr_tokens = t5_tokenizer.tokenize(curr_text, nullptr, true);
|
||||
t5_tokens.insert(t5_tokens.end(), curr_tokens.begin(), curr_tokens.end());
|
||||
t5_weights.insert(t5_weights.end(), curr_tokens.size(), curr_weight);
|
||||
}
|
||||
t5_tokenizer.pad_tokens(t5_tokens, &t5_weights, nullptr);
|
||||
|
||||
return {qwen_tokens, qwen_weights, t5_tokens, t5_weights};
|
||||
}
|
||||
@ -1685,7 +1684,7 @@ struct LLMEmbedder : public Conditioner {
|
||||
std::shared_ptr<LLM::LLMRunner> llm;
|
||||
|
||||
LLMEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
SDVersion version = VERSION_QWEN_IMAGE,
|
||||
const std::string prefix = "",
|
||||
@ -1706,7 +1705,7 @@ struct LLMEmbedder : public Conditioner {
|
||||
}
|
||||
llm = std::make_shared<LLM::LLMRunner>(arch,
|
||||
backend,
|
||||
params_backend,
|
||||
offload_params_to_cpu,
|
||||
tensor_storage_map,
|
||||
"text_encoders.llm",
|
||||
enable_vision);
|
||||
@ -2070,10 +2069,10 @@ struct LTXAVTextProjectionRunner : public GGMLRunner {
|
||||
LTXAVTextProjection model;
|
||||
|
||||
LTXAVTextProjectionRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& prefix = "")
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
model(tensor_storage_map.find(prefix + ".video_aggregate_embed.weight") != tensor_storage_map.end()) {
|
||||
model.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
@ -2114,20 +2113,20 @@ struct LTXAVEmbedder : public Conditioner {
|
||||
bool dual_projection = false;
|
||||
|
||||
LTXAVEmbedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& llm_prefix = "text_encoders.llm",
|
||||
const std::string& projector_prefix = "text_embedding_projection") {
|
||||
tokenizer = std::make_shared<GemmaTokenizer>();
|
||||
llm = std::make_shared<LLM::LLMRunner>(LLM::LLMArch::GEMMA3_12B,
|
||||
backend,
|
||||
params_backend,
|
||||
offload_params_to_cpu,
|
||||
tensor_storage_map,
|
||||
llm_prefix,
|
||||
false);
|
||||
dual_projection = tensor_storage_map.find(projector_prefix + ".video_aggregate_embed.weight") != tensor_storage_map.end();
|
||||
projector = std::make_shared<LTXAVTextProjectionRunner>(backend,
|
||||
params_backend,
|
||||
offload_params_to_cpu,
|
||||
tensor_storage_map,
|
||||
projector_prefix);
|
||||
}
|
||||
|
||||
@ -319,10 +319,10 @@ struct ControlNet : public GGMLRunner {
|
||||
bool guided_hint_cached = false;
|
||||
|
||||
ControlNet(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
SDVersion version = VERSION_SD1)
|
||||
: GGMLRunner(backend, params_backend), control_net(version) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu), control_net(version) {
|
||||
control_net.init(params_ctx, tensor_storage_map, "");
|
||||
}
|
||||
|
||||
|
||||
@ -68,10 +68,10 @@ struct UNetModel : public DiffusionModel {
|
||||
UNetModelRunner unet;
|
||||
|
||||
UNetModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
SDVersion version = VERSION_SD1)
|
||||
: unet(backend, params_backend, tensor_storage_map, "model.diffusion_model", version) {
|
||||
: unet(backend, offload_params_to_cpu, tensor_storage_map, "model.diffusion_model", version) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -139,9 +139,9 @@ struct MMDiTModel : public DiffusionModel {
|
||||
MMDiTRunner mmdit;
|
||||
|
||||
MMDiTModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {})
|
||||
: mmdit(backend, params_backend, tensor_storage_map, "model.diffusion_model") {
|
||||
: mmdit(backend, offload_params_to_cpu, tensor_storage_map, "model.diffusion_model") {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -206,11 +206,11 @@ struct FluxModel : public DiffusionModel {
|
||||
Flux::FluxRunner flux;
|
||||
|
||||
FluxModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
SDVersion version = VERSION_FLUX,
|
||||
bool use_mask = false)
|
||||
: flux(backend, params_backend, tensor_storage_map, "model.diffusion_model", version, use_mask) {
|
||||
: flux(backend, offload_params_to_cpu, tensor_storage_map, "model.diffusion_model", version, use_mask) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -281,10 +281,10 @@ struct AnimaModel : public DiffusionModel {
|
||||
Anima::AnimaRunner anima;
|
||||
|
||||
AnimaModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model")
|
||||
: prefix(prefix), anima(backend, params_backend, tensor_storage_map, prefix) {
|
||||
: prefix(prefix), anima(backend, offload_params_to_cpu, tensor_storage_map, prefix) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -349,11 +349,11 @@ struct WanModel : public DiffusionModel {
|
||||
WAN::WanRunner wan;
|
||||
|
||||
WanModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model",
|
||||
SDVersion version = VERSION_WAN2)
|
||||
: prefix(prefix), wan(backend, params_backend, tensor_storage_map, prefix, version) {
|
||||
: prefix(prefix), wan(backend, offload_params_to_cpu, tensor_storage_map, prefix, version) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -421,12 +421,12 @@ struct QwenImageModel : public DiffusionModel {
|
||||
Qwen::QwenImageRunner qwen_image;
|
||||
|
||||
QwenImageModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model",
|
||||
SDVersion version = VERSION_QWEN_IMAGE,
|
||||
bool zero_cond_t = false)
|
||||
: prefix(prefix), qwen_image(backend, params_backend, tensor_storage_map, prefix, version, zero_cond_t) {
|
||||
: prefix(prefix), qwen_image(backend, offload_params_to_cpu, tensor_storage_map, prefix, version, zero_cond_t) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -492,10 +492,10 @@ struct HiDreamO1Model : public DiffusionModel {
|
||||
HiDreamO1::HiDreamO1Runner hidream_o1;
|
||||
|
||||
HiDreamO1Model(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& prefix = "model")
|
||||
: prefix(prefix), hidream_o1(backend, params_backend, tensor_storage_map, prefix) {
|
||||
: prefix(prefix), hidream_o1(backend, offload_params_to_cpu, tensor_storage_map, prefix) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -568,11 +568,11 @@ struct ZImageModel : public DiffusionModel {
|
||||
ZImage::ZImageRunner z_image;
|
||||
|
||||
ZImageModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model",
|
||||
SDVersion version = VERSION_Z_IMAGE)
|
||||
: prefix(prefix), z_image(backend, params_backend, tensor_storage_map, prefix, version) {
|
||||
: prefix(prefix), z_image(backend, offload_params_to_cpu, tensor_storage_map, prefix, version) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -638,10 +638,10 @@ struct ErnieImageModel : public DiffusionModel {
|
||||
ErnieImage::ErnieImageRunner ernie_image;
|
||||
|
||||
ErnieImageModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model")
|
||||
: prefix(prefix), ernie_image(backend, params_backend, tensor_storage_map, prefix) {
|
||||
: prefix(prefix), ernie_image(backend, offload_params_to_cpu, tensor_storage_map, prefix) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
@ -704,10 +704,10 @@ struct LTXAVModel : public DiffusionModel {
|
||||
LTXV::LTXAVRunner ltxav;
|
||||
|
||||
LTXAVModel(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "model.diffusion_model")
|
||||
: prefix(prefix), ltxav(backend, params_backend, tensor_storage_map, prefix) {
|
||||
: prefix(prefix), ltxav(backend, offload_params_to_cpu, tensor_storage_map, prefix) {
|
||||
}
|
||||
|
||||
std::string get_desc() override {
|
||||
|
||||
@ -331,10 +331,10 @@ namespace ErnieImage {
|
||||
std::vector<float> pe_vec;
|
||||
|
||||
ErnieImageRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "")
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
ernie_params.num_layers = 0;
|
||||
for (const auto& [name, tensor_storage] : tensor_storage_map) {
|
||||
if (!starts_with(name, prefix)) {
|
||||
|
||||
@ -161,10 +161,10 @@ struct ESRGAN : public GGMLRunner {
|
||||
int tile_size = 128; // avoid cuda OOM for 4gb VRAM
|
||||
|
||||
ESRGAN(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
int tile_size = 128,
|
||||
const String2TensorStorage& tensor_storage_map = {})
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
this->tile_size = tile_size;
|
||||
}
|
||||
|
||||
|
||||
@ -1189,12 +1189,12 @@ namespace Flux {
|
||||
bool use_mask = false;
|
||||
|
||||
FluxRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
SDVersion version = VERSION_FLUX,
|
||||
bool use_mask = false)
|
||||
: GGMLRunner(backend, params_backend), version(version), use_mask(use_mask) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu), version(version), use_mask(use_mask) {
|
||||
flux_params.version = version;
|
||||
flux_params.guidance_embed = false;
|
||||
flux_params.depth = 0;
|
||||
@ -1564,7 +1564,7 @@ namespace Flux {
|
||||
}
|
||||
|
||||
std::shared_ptr<FluxRunner> flux = std::make_shared<FluxRunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"model.diffusion_model",
|
||||
VERSION_FLUX2,
|
||||
|
||||
@ -26,7 +26,7 @@
|
||||
#include "ggml-alloc.h"
|
||||
#include "ggml-backend.h"
|
||||
#include "ggml.h"
|
||||
#include "ggml_extend_backend.h"
|
||||
#include "ggml_extend_backend.hpp"
|
||||
#include "ggml_graph_cut.h"
|
||||
|
||||
#include "model.h"
|
||||
@ -73,6 +73,48 @@ __STATIC_INLINE__ void ggml_log_callback_default(ggml_log_level level, const cha
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ bool backend_name_exists(std::string name) {
|
||||
ggml_backend_load_all_once();
|
||||
const size_t device_count = ggml_backend_dev_count();
|
||||
for (size_t i = 0; i < device_count; ++i) {
|
||||
if (name == ggml_backend_dev_name(ggml_backend_dev_get(i))) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ std::string sanitize_backend_name(std::string name) {
|
||||
if (name == "" || backend_name_exists(name)) {
|
||||
return name;
|
||||
} else {
|
||||
LOG_WARN("Backend %s not found, using default backend", name.c_str());
|
||||
return "";
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ std::string get_default_backend_name() {
|
||||
ggml_backend_load_all_once();
|
||||
// should pick the same backend as ggml_backend_init_best
|
||||
ggml_backend_dev_t dev = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_GPU);
|
||||
dev = dev ? dev : ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_IGPU);
|
||||
dev = dev ? dev : ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU);
|
||||
if (dev == nullptr) {
|
||||
return "";
|
||||
}
|
||||
return ggml_backend_dev_name(dev);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_backend_t init_named_backend(std::string name = "") {
|
||||
ggml_backend_load_all_once();
|
||||
LOG_DEBUG("Initializing backend: %s", name.c_str());
|
||||
if (name.empty()) {
|
||||
return ggml_backend_init_best();
|
||||
} else {
|
||||
return ggml_backend_init_by_name(name.c_str(), nullptr);
|
||||
}
|
||||
}
|
||||
|
||||
static_assert(GGML_MAX_NAME >= 128, "GGML_MAX_NAME must be at least 128");
|
||||
|
||||
// n-mode tensor-matrix product
|
||||
@ -148,7 +190,7 @@ __STATIC_INLINE__ void ggml_ext_im_set_randn_f32(ggml_tensor* tensor, std::share
|
||||
uint32_t n = (uint32_t)ggml_nelements(tensor);
|
||||
std::vector<float> random_numbers = rng->randn(n);
|
||||
for (uint32_t i = 0; i < n; i++) {
|
||||
ggml_ext_im_set_f32_1d(tensor, i, random_numbers[i]);
|
||||
ggml_set_f32_1d(tensor, i, random_numbers[i]);
|
||||
}
|
||||
}
|
||||
|
||||
@ -380,6 +422,39 @@ __STATIC_INLINE__ ggml_tensor* load_tensor_from_file(ggml_context* ctx, const st
|
||||
// file.close();
|
||||
// }
|
||||
|
||||
__STATIC_INLINE__ void copy_ggml_tensor(ggml_tensor* dst, ggml_tensor* src) {
|
||||
if (dst->type == src->type) {
|
||||
dst->nb[0] = src->nb[0];
|
||||
dst->nb[1] = src->nb[1];
|
||||
dst->nb[2] = src->nb[2];
|
||||
dst->nb[3] = src->nb[3];
|
||||
|
||||
memcpy(((char*)dst->data), ((char*)src->data), ggml_nbytes(dst));
|
||||
return;
|
||||
}
|
||||
ggml_init_params params;
|
||||
params.mem_size = 10 * 1024 * 1024; // for padding
|
||||
params.mem_buffer = nullptr;
|
||||
params.no_alloc = false;
|
||||
ggml_context* ctx = ggml_init(params);
|
||||
if (!ctx) {
|
||||
LOG_ERROR("ggml_init() failed");
|
||||
return;
|
||||
}
|
||||
ggml_tensor* final = ggml_cpy(ctx, src, dst);
|
||||
|
||||
ggml_cgraph* graph = ggml_new_graph(ctx);
|
||||
ggml_build_forward_expand(graph, final);
|
||||
ggml_graph_compute_with_ctx(ctx, graph, 1);
|
||||
ggml_free(ctx);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_tensor* ggml_ext_dup_and_cpy_tensor(ggml_context* ctx, ggml_tensor* src) {
|
||||
ggml_tensor* dup = ggml_dup_tensor(ctx, src);
|
||||
copy_ggml_tensor(dup, src);
|
||||
return dup;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ float sigmoid(float x) {
|
||||
return 1 / (1.0f + expf(-x));
|
||||
}
|
||||
@ -1127,33 +1202,18 @@ __STATIC_INLINE__ ggml_tensor* ggml_ext_conv_3d(ggml_context* ctx,
|
||||
ggml_tensor* w,
|
||||
ggml_tensor* b,
|
||||
int64_t IC,
|
||||
int s0 = 1,
|
||||
int s1 = 1,
|
||||
int s2 = 1,
|
||||
int p0 = 0,
|
||||
int p1 = 0,
|
||||
int p2 = 0,
|
||||
int d0 = 1,
|
||||
int d1 = 1,
|
||||
int d2 = 1,
|
||||
bool force_prec_f32 = false) {
|
||||
if (force_prec_f32) {
|
||||
ggml_tensor* im2col = ggml_im2col_3d(ctx, w, x, IC, s0, s1, s2, p0, p1, p2, d0, d1, d2, w->type);
|
||||
|
||||
int64_t OC = w->ne[3] / IC;
|
||||
int64_t N = x->ne[3] / IC;
|
||||
x = ggml_mul_mat(ctx,
|
||||
ggml_reshape_2d(ctx, im2col, im2col->ne[0], im2col->ne[3] * im2col->ne[2] * im2col->ne[1]),
|
||||
ggml_reshape_2d(ctx, w, w->ne[0] * w->ne[1] * w->ne[2] * IC, OC));
|
||||
ggml_mul_mat_set_prec(x, GGML_PREC_F32);
|
||||
|
||||
int64_t OD = im2col->ne[3] / N;
|
||||
x = ggml_reshape_4d(ctx, x, im2col->ne[1] * im2col->ne[2], OD, N, OC);
|
||||
x = ggml_cont(ctx, ggml_permute(ctx, x, 0, 1, 3, 2));
|
||||
x = ggml_reshape_4d(ctx, x, im2col->ne[1], im2col->ne[2], OD, OC * N);
|
||||
} else {
|
||||
x = ggml_conv_3d(ctx, w, x, IC, s0, s1, s2, p0, p1, p2, d0, d1, d2);
|
||||
}
|
||||
int s0 = 1,
|
||||
int s1 = 1,
|
||||
int s2 = 1,
|
||||
int p0 = 0,
|
||||
int p1 = 0,
|
||||
int p2 = 0,
|
||||
int d0 = 1,
|
||||
int d1 = 1,
|
||||
int d2 = 1) {
|
||||
int64_t OC = w->ne[3] / IC;
|
||||
int64_t N = x->ne[3] / IC;
|
||||
x = ggml_conv_3d(ctx, w, x, IC, s0, s1, s2, p0, p1, p2, d0, d1, d2);
|
||||
|
||||
if (b != nullptr) {
|
||||
b = ggml_reshape_4d(ctx, b, 1, 1, 1, b->ne[0]); // [OC, 1, 1, 1]
|
||||
@ -2609,11 +2669,13 @@ protected:
|
||||
public:
|
||||
virtual std::string get_desc() = 0;
|
||||
|
||||
GGMLRunner(ggml_backend_t backend, ggml_backend_t params_backend)
|
||||
: params_backend(params_backend),
|
||||
runtime_backend(backend) {
|
||||
GGML_ASSERT(runtime_backend != nullptr);
|
||||
GGML_ASSERT(params_backend != nullptr);
|
||||
GGMLRunner(ggml_backend_t backend, bool offload_params_to_cpu = false)
|
||||
: runtime_backend(backend) {
|
||||
if (!ggml_backend_is_cpu(runtime_backend) && offload_params_to_cpu) {
|
||||
params_backend = ggml_backend_cpu_init();
|
||||
} else {
|
||||
params_backend = runtime_backend;
|
||||
}
|
||||
alloc_params_ctx();
|
||||
}
|
||||
|
||||
@ -2622,6 +2684,9 @@ public:
|
||||
free_compute_buffer();
|
||||
free_params_ctx();
|
||||
free_compute_ctx();
|
||||
if (params_backend != runtime_backend) {
|
||||
ggml_backend_free(params_backend);
|
||||
}
|
||||
free_cache_ctx_and_buffer();
|
||||
}
|
||||
|
||||
@ -3148,7 +3213,6 @@ protected:
|
||||
std::tuple<int, int, int> padding;
|
||||
std::tuple<int, int, int> dilation;
|
||||
bool bias;
|
||||
bool force_prec_f32;
|
||||
std::string prefix;
|
||||
|
||||
void init_params(ggml_context* ctx, const String2TensorStorage& tensor_storage_map, const std::string prefix = "") override {
|
||||
@ -3172,16 +3236,14 @@ public:
|
||||
std::tuple<int, int, int> stride = {1, 1, 1},
|
||||
std::tuple<int, int, int> padding = {0, 0, 0},
|
||||
std::tuple<int, int, int> dilation = {1, 1, 1},
|
||||
bool bias = true,
|
||||
bool force_prec_f32 = false)
|
||||
bool bias = true)
|
||||
: in_channels(in_channels),
|
||||
out_channels(out_channels),
|
||||
kernel_size(kernel_size),
|
||||
stride(stride),
|
||||
padding(padding),
|
||||
dilation(dilation),
|
||||
bias(bias),
|
||||
force_prec_f32(force_prec_f32) {}
|
||||
bias(bias) {}
|
||||
|
||||
ggml_tensor* forward(GGMLRunnerContext* ctx, ggml_tensor* x) {
|
||||
ggml_tensor* w = params["weight"];
|
||||
@ -3201,8 +3263,7 @@ public:
|
||||
return ggml_ext_conv_3d(ctx->ggml_ctx, x, w, b, in_channels,
|
||||
std::get<2>(stride), std::get<1>(stride), std::get<0>(stride),
|
||||
std::get<2>(padding), std::get<1>(padding), std::get<0>(padding),
|
||||
std::get<2>(dilation), std::get<1>(dilation), std::get<0>(dilation),
|
||||
force_prec_f32);
|
||||
std::get<2>(dilation), std::get<1>(dilation), std::get<0>(dilation));
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
@ -1,600 +0,0 @@
|
||||
#include "ggml_extend_backend.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <cctype>
|
||||
#include <cstdlib>
|
||||
#include <mutex>
|
||||
#include <sstream>
|
||||
#include <stdexcept>
|
||||
#include <vector>
|
||||
|
||||
#include "util.h"
|
||||
|
||||
static std::string trim_copy(const std::string& value) {
|
||||
size_t begin = 0;
|
||||
while (begin < value.size() && std::isspace(static_cast<unsigned char>(value[begin]))) {
|
||||
++begin;
|
||||
}
|
||||
size_t end = value.size();
|
||||
while (end > begin && std::isspace(static_cast<unsigned char>(value[end - 1]))) {
|
||||
--end;
|
||||
}
|
||||
return value.substr(begin, end - begin);
|
||||
}
|
||||
|
||||
static std::string lower_copy(std::string value) {
|
||||
std::transform(value.begin(), value.end(), value.begin(), [](unsigned char c) {
|
||||
return static_cast<char>(std::tolower(c));
|
||||
});
|
||||
return value;
|
||||
}
|
||||
|
||||
static std::vector<std::string> split_copy(const std::string& value, char delimiter) {
|
||||
std::vector<std::string> parts;
|
||||
std::string part;
|
||||
std::istringstream stream(value);
|
||||
while (std::getline(stream, part, delimiter)) {
|
||||
parts.push_back(part);
|
||||
}
|
||||
return parts;
|
||||
}
|
||||
|
||||
static bool is_default_backend_token(const std::string& name) {
|
||||
const std::string lower = lower_copy(trim_copy(name));
|
||||
return lower.empty() || lower == "default" || lower == "auto";
|
||||
}
|
||||
|
||||
static bool parse_backend_module(const std::string& raw_name, SDBackendModule* module) {
|
||||
std::string name = lower_copy(trim_copy(raw_name));
|
||||
name.erase(std::remove(name.begin(), name.end(), '-'), name.end());
|
||||
name.erase(std::remove(name.begin(), name.end(), '_'), name.end());
|
||||
|
||||
if (name == "diffusion" || name == "model" || name == "unet" || name == "dit") {
|
||||
*module = SDBackendModule::DIFFUSION;
|
||||
return true;
|
||||
}
|
||||
if (name == "te" || name == "clip" || name == "text" || name == "textencoder" || name == "textencoders" || name == "conditioner" || name == "cond" || name == "llm" || name == "t5" || name == "t5xxl") {
|
||||
*module = SDBackendModule::TE;
|
||||
return true;
|
||||
}
|
||||
if (name == "clipvision" || name == "vision") {
|
||||
*module = SDBackendModule::CLIP_VISION;
|
||||
return true;
|
||||
}
|
||||
if (name == "vae" || name == "firststage" || name == "autoencoder" || name == "tae") {
|
||||
*module = SDBackendModule::VAE;
|
||||
return true;
|
||||
}
|
||||
if (name == "controlnet" || name == "control") {
|
||||
*module = SDBackendModule::CONTROL_NET;
|
||||
return true;
|
||||
}
|
||||
if (name == "photomaker" || name == "photomakerid" || name == "pmid" || name == "photo") {
|
||||
*module = SDBackendModule::PHOTOMAKER;
|
||||
return true;
|
||||
}
|
||||
if (name == "upscaler" || name == "esrgan" || name == "hires") {
|
||||
*module = SDBackendModule::UPSCALER;
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
static std::string module_assignment_name(const SDBackendAssignment& assignment, SDBackendModule module) {
|
||||
auto it = assignment.module_names.find(module);
|
||||
if (it != assignment.module_names.end()) {
|
||||
return it->second;
|
||||
}
|
||||
return assignment.default_name;
|
||||
}
|
||||
|
||||
static std::string backend_cache_key(ggml_backend_t backend) {
|
||||
if (backend == nullptr) {
|
||||
return "";
|
||||
}
|
||||
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
|
||||
if (dev != nullptr) {
|
||||
return lower_copy(ggml_backend_dev_name(dev));
|
||||
}
|
||||
const char* backend_name = ggml_backend_name(backend);
|
||||
return backend_name != nullptr ? lower_copy(backend_name) : "";
|
||||
}
|
||||
|
||||
static std::string resolve_first_device_by_type(enum ggml_backend_dev_type type) {
|
||||
ggml_backend_dev_t dev = ggml_backend_dev_by_type(type);
|
||||
if (dev == nullptr) {
|
||||
return "";
|
||||
}
|
||||
return ggml_backend_dev_name(dev);
|
||||
}
|
||||
|
||||
static ggml_backend_buffer_t ggml_backend_tensor_buffer(const struct ggml_tensor* tensor) {
|
||||
if (tensor == nullptr) {
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
return tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
|
||||
}
|
||||
|
||||
static bool ggml_backend_tensor_is_host_accessible(const struct ggml_tensor* tensor) {
|
||||
if (tensor == nullptr || tensor->data == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
ggml_backend_buffer_t buffer = ggml_backend_tensor_buffer(tensor);
|
||||
return buffer == nullptr || ggml_backend_buffer_is_host(buffer);
|
||||
}
|
||||
|
||||
static size_t ggml_backend_tensor_offset(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3) {
|
||||
return static_cast<size_t>(i0 * tensor->nb[0] + i1 * tensor->nb[1] + i2 * tensor->nb[2] + i3 * tensor->nb[3]);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void ggml_backend_tensor_write_scalar(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3, T value) {
|
||||
const size_t offset = ggml_backend_tensor_offset(tensor, i0, i1, i2, i3);
|
||||
|
||||
if (ggml_backend_tensor_is_host_accessible(tensor)) {
|
||||
auto* dst = reinterpret_cast<T*>(reinterpret_cast<char*>(tensor->data) + offset);
|
||||
*dst = value;
|
||||
return;
|
||||
}
|
||||
|
||||
ggml_backend_tensor_set(const_cast<struct ggml_tensor*>(tensor), &value, offset, sizeof(T));
|
||||
}
|
||||
|
||||
static void ggml_set_f32_nd(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3, float value) {
|
||||
switch (tensor->type) {
|
||||
case GGML_TYPE_I8:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int8_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int16_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int32_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_F16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, ggml_fp32_to_fp16(value));
|
||||
break;
|
||||
case GGML_TYPE_BF16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, ggml_fp32_to_bf16(value));
|
||||
break;
|
||||
case GGML_TYPE_F32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, value);
|
||||
break;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
}
|
||||
|
||||
void ggml_ext_im_set_f32_1d(const struct ggml_tensor* tensor, int i, float value) {
|
||||
if (!ggml_is_contiguous(tensor)) {
|
||||
int64_t id[4] = {0, 0, 0, 0};
|
||||
ggml_unravel_index(tensor, i, &id[0], &id[1], &id[2], &id[3]);
|
||||
ggml_set_f32_nd(tensor, id[0], id[1], id[2], id[3], value);
|
||||
return;
|
||||
}
|
||||
|
||||
switch (tensor->type) {
|
||||
case GGML_TYPE_I8:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int8_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int16_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int32_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_F16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, ggml_fp32_to_fp16(value));
|
||||
break;
|
||||
case GGML_TYPE_BF16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, ggml_fp32_to_bf16(value));
|
||||
break;
|
||||
case GGML_TYPE_F32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, value);
|
||||
break;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
}
|
||||
|
||||
static void ggml_backend_load_all_once() {
|
||||
// If the registry already has devices and the CPU backend is present,
|
||||
// assume either static registration or explicit host-side preloading has
|
||||
// completed and avoid rescanning the default paths.
|
||||
if (ggml_backend_dev_count() > 0 && ggml_backend_reg_by_name("CPU") != nullptr) {
|
||||
return;
|
||||
}
|
||||
// In dynamic-backend mode the backend modules are discovered at runtime,
|
||||
// so we must load them before asking for the CPU backend or its proc table.
|
||||
// If the host preloaded only a subset of backends, allow one default-path
|
||||
// scan so missing modules can still be discovered.
|
||||
static std::once_flag once;
|
||||
std::call_once(once, []() {
|
||||
if (ggml_backend_dev_count() > 0 && ggml_backend_reg_by_name("CPU") != nullptr) {
|
||||
return;
|
||||
}
|
||||
ggml_backend_load_all();
|
||||
});
|
||||
}
|
||||
|
||||
bool sd_backend_is(ggml_backend_t backend, const std::string& name) {
|
||||
if (!backend) {
|
||||
return false;
|
||||
}
|
||||
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
|
||||
if (!dev) {
|
||||
return false;
|
||||
}
|
||||
std::string dev_name = ggml_backend_dev_name(dev);
|
||||
return lower_copy(dev_name).find(lower_copy(name)) != std::string::npos;
|
||||
}
|
||||
|
||||
static std::string get_default_backend_name() {
|
||||
ggml_backend_load_all_once();
|
||||
// should pick the same backend preference as ggml_backend_init_best
|
||||
std::string name = resolve_first_device_by_type(GGML_BACKEND_DEVICE_TYPE_GPU);
|
||||
if (!name.empty()) {
|
||||
return name;
|
||||
}
|
||||
name = resolve_first_device_by_type(GGML_BACKEND_DEVICE_TYPE_IGPU);
|
||||
if (!name.empty()) {
|
||||
return name;
|
||||
}
|
||||
return resolve_first_device_by_type(GGML_BACKEND_DEVICE_TYPE_CPU);
|
||||
}
|
||||
|
||||
static std::string sd_resolve_backend_name(const std::string& name) {
|
||||
ggml_backend_load_all_once();
|
||||
std::string requested = trim_copy(name);
|
||||
std::string lower = lower_copy(requested);
|
||||
|
||||
if (is_default_backend_token(lower)) {
|
||||
return get_default_backend_name();
|
||||
}
|
||||
if (lower == "gpu") {
|
||||
std::string result = resolve_first_device_by_type(GGML_BACKEND_DEVICE_TYPE_GPU);
|
||||
if (!result.empty()) {
|
||||
return result;
|
||||
}
|
||||
return resolve_first_device_by_type(GGML_BACKEND_DEVICE_TYPE_IGPU);
|
||||
}
|
||||
|
||||
const size_t device_count = ggml_backend_dev_count();
|
||||
for (size_t i = 0; i < device_count; ++i) {
|
||||
ggml_backend_dev_t dev = ggml_backend_dev_get(i);
|
||||
std::string dev_name = ggml_backend_dev_name(dev);
|
||||
if (lower_copy(dev_name) == lower) {
|
||||
return dev_name;
|
||||
}
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < device_count; ++i) {
|
||||
ggml_backend_dev_t dev = ggml_backend_dev_get(i);
|
||||
std::string dev_name = ggml_backend_dev_name(dev);
|
||||
std::string dev_lower = lower_copy(dev_name);
|
||||
if (dev_lower.rfind(lower, 0) == 0) {
|
||||
return dev_name;
|
||||
}
|
||||
}
|
||||
|
||||
return "";
|
||||
}
|
||||
|
||||
static bool backend_name_exists(const std::string& name) {
|
||||
return !sd_resolve_backend_name(name).empty();
|
||||
}
|
||||
|
||||
static ggml_backend_t init_named_backend(const std::string& name) {
|
||||
ggml_backend_load_all_once();
|
||||
LOG_DEBUG("Initializing backend: %s", name.c_str());
|
||||
if (trim_copy(name).empty()) {
|
||||
return ggml_backend_init_best();
|
||||
}
|
||||
|
||||
std::string resolved = sd_resolve_backend_name(name);
|
||||
if (resolved.empty()) {
|
||||
return nullptr;
|
||||
}
|
||||
return ggml_backend_init_by_name(resolved.c_str(), nullptr);
|
||||
}
|
||||
|
||||
static ggml_backend_t sd_get_default_backend() {
|
||||
ggml_backend_load_all_once();
|
||||
static std::once_flag once;
|
||||
std::call_once(once, []() {
|
||||
size_t dev_count = ggml_backend_dev_count();
|
||||
if (dev_count == 0) {
|
||||
LOG_ERROR("No devices found!");
|
||||
} else {
|
||||
LOG_DEBUG("Found %zu backend devices:", dev_count);
|
||||
for (size_t i = 0; i < dev_count; ++i) {
|
||||
auto dev = ggml_backend_dev_get(i);
|
||||
LOG_DEBUG("#%zu: %s", i, ggml_backend_dev_name(dev));
|
||||
}
|
||||
}
|
||||
});
|
||||
|
||||
ggml_backend_t backend = nullptr;
|
||||
const char* SD_VK_DEVICE = getenv("SD_VK_DEVICE");
|
||||
if (SD_VK_DEVICE != nullptr) {
|
||||
std::string sd_vk_device_str = SD_VK_DEVICE;
|
||||
try {
|
||||
unsigned long long device = std::stoull(sd_vk_device_str);
|
||||
std::string vk_device_name = "Vulkan" + std::to_string(device);
|
||||
if (backend_name_exists(vk_device_name)) {
|
||||
LOG_INFO("Selecting %s as main device by env var SD_VK_DEVICE", vk_device_name.c_str());
|
||||
backend = init_named_backend(vk_device_name);
|
||||
if (!backend) {
|
||||
LOG_WARN("Device %s requested by SD_VK_DEVICE failed to init. Falling back to the default device.", vk_device_name.c_str());
|
||||
}
|
||||
} else {
|
||||
LOG_WARN("Device %s requested by SD_VK_DEVICE was not found. Falling back to the default device.", vk_device_name.c_str());
|
||||
}
|
||||
} catch (const std::invalid_argument&) {
|
||||
LOG_WARN("SD_VK_DEVICE environment variable is not a valid integer (%s). Falling back to the default device.", SD_VK_DEVICE);
|
||||
} catch (const std::out_of_range&) {
|
||||
LOG_WARN("SD_VK_DEVICE environment variable value is out of range for `unsigned long long` type (%s). Falling back to the default device.", SD_VK_DEVICE);
|
||||
}
|
||||
}
|
||||
|
||||
if (!backend) {
|
||||
std::string dev_name = get_default_backend_name();
|
||||
backend = init_named_backend(dev_name);
|
||||
if (!backend && !dev_name.empty()) {
|
||||
LOG_WARN("device %s failed to init", dev_name.c_str());
|
||||
}
|
||||
}
|
||||
|
||||
if (!backend) {
|
||||
LOG_WARN("loading CPU backend");
|
||||
backend = ggml_backend_cpu_init();
|
||||
}
|
||||
|
||||
if (ggml_backend_is_cpu(backend)) {
|
||||
LOG_DEBUG("Using CPU backend");
|
||||
}
|
||||
|
||||
return backend;
|
||||
}
|
||||
|
||||
static bool sd_parse_backend_assignment(const std::string& spec, SDBackendAssignment* assignment, std::string* error) {
|
||||
if (assignment == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
*assignment = {};
|
||||
const std::string in = trim_copy(spec);
|
||||
if (in.empty()) {
|
||||
return true;
|
||||
}
|
||||
|
||||
for (const std::string& raw_part : split_copy(in, ',')) {
|
||||
const std::string part = trim_copy(raw_part);
|
||||
if (part.empty()) {
|
||||
continue;
|
||||
}
|
||||
|
||||
const size_t eq = part.find('=');
|
||||
if (eq == std::string::npos) {
|
||||
assignment->set_default(part);
|
||||
continue;
|
||||
}
|
||||
|
||||
const std::string key = trim_copy(part.substr(0, eq));
|
||||
const std::string value = trim_copy(part.substr(eq + 1));
|
||||
if (key.empty() || value.empty()) {
|
||||
if (error != nullptr) {
|
||||
*error = "invalid backend assignment '" + part + "'";
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
const std::string key_lower = lower_copy(key);
|
||||
if (key_lower == "all" || key_lower == "default" || key_lower == "*") {
|
||||
assignment->set_default(value);
|
||||
continue;
|
||||
}
|
||||
|
||||
SDBackendModule module = SDBackendModule::DIFFUSION;
|
||||
if (!parse_backend_module(key, &module)) {
|
||||
if (error != nullptr) {
|
||||
*error = "unknown backend module '" + key + "'";
|
||||
}
|
||||
return false;
|
||||
}
|
||||
assignment->set_module(module, value);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool SDBackendAssignment::empty() const {
|
||||
return default_name.empty() && module_names.empty();
|
||||
}
|
||||
|
||||
std::string SDBackendAssignment::get(SDBackendModule module) const {
|
||||
return module_assignment_name(*this, module);
|
||||
}
|
||||
|
||||
void SDBackendAssignment::set_default(const std::string& name) {
|
||||
default_name = trim_copy(name);
|
||||
}
|
||||
|
||||
void SDBackendAssignment::set_module(SDBackendModule module, const std::string& name) {
|
||||
module_names[module] = trim_copy(name);
|
||||
}
|
||||
|
||||
void SDBackendHandleDeleter::operator()(ggml_backend_t backend) const {
|
||||
ggml_backend_free(backend);
|
||||
}
|
||||
|
||||
SDBackendManager::~SDBackendManager() {
|
||||
reset();
|
||||
}
|
||||
|
||||
void SDBackendManager::reset() {
|
||||
backends_.clear();
|
||||
runtime_assignment_ = {};
|
||||
params_assignment_ = {};
|
||||
}
|
||||
|
||||
ggml_backend_t SDBackendManager::runtime_backend(SDBackendModule module) {
|
||||
return init_cached_backend(runtime_assignment_.get(module));
|
||||
}
|
||||
|
||||
ggml_backend_t SDBackendManager::params_backend(SDBackendModule module) {
|
||||
std::string name = params_assignment_.get(module);
|
||||
if (name.empty()) {
|
||||
return runtime_backend(module);
|
||||
}
|
||||
return init_cached_backend(name);
|
||||
}
|
||||
|
||||
bool SDBackendManager::runtime_backend_is_cpu(SDBackendModule module) {
|
||||
return ggml_backend_is_cpu(runtime_backend(module));
|
||||
}
|
||||
|
||||
bool SDBackendManager::params_backend_is_cpu(SDBackendModule module) {
|
||||
return ggml_backend_is_cpu(params_backend(module));
|
||||
}
|
||||
|
||||
bool SDBackendManager::runtime_backend_supports_host_buffer(SDBackendModule module) {
|
||||
ggml_backend_t backend = runtime_backend(module);
|
||||
if (backend == nullptr) {
|
||||
return false;
|
||||
}
|
||||
if (ggml_backend_is_cpu(backend)) {
|
||||
return true;
|
||||
}
|
||||
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
|
||||
if (dev == nullptr) {
|
||||
return false;
|
||||
}
|
||||
ggml_backend_dev_props props;
|
||||
ggml_backend_dev_get_props(dev, &props);
|
||||
return props.caps.buffer_from_host_ptr;
|
||||
}
|
||||
|
||||
bool SDBackendManager::init(const char* backend_spec,
|
||||
const char* params_backend_spec,
|
||||
bool offload_params_to_cpu,
|
||||
bool keep_clip_on_cpu,
|
||||
bool keep_vae_on_cpu,
|
||||
bool keep_control_net_on_cpu,
|
||||
std::string* error) {
|
||||
reset();
|
||||
|
||||
if (!sd_parse_backend_assignment(SAFE_STR(backend_spec), &runtime_assignment_, error)) {
|
||||
return false;
|
||||
}
|
||||
if (!sd_parse_backend_assignment(SAFE_STR(params_backend_spec), ¶ms_assignment_, error)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (runtime_assignment_.empty()) {
|
||||
if (keep_clip_on_cpu) {
|
||||
runtime_assignment_.set_module(SDBackendModule::TE, "cpu");
|
||||
}
|
||||
if (keep_vae_on_cpu) {
|
||||
runtime_assignment_.set_module(SDBackendModule::VAE, "cpu");
|
||||
}
|
||||
if (keep_control_net_on_cpu) {
|
||||
runtime_assignment_.set_module(SDBackendModule::CONTROL_NET, "cpu");
|
||||
}
|
||||
}
|
||||
|
||||
if (params_assignment_.empty() && offload_params_to_cpu) {
|
||||
params_assignment_.set_default("cpu");
|
||||
}
|
||||
|
||||
return validate(error);
|
||||
}
|
||||
|
||||
bool SDBackendManager::validate(std::string* error) const {
|
||||
auto validate_name = [&](const std::string& name) -> bool {
|
||||
if (is_default_backend_token(name)) {
|
||||
return true;
|
||||
}
|
||||
if (!sd_resolve_backend_name(name).empty()) {
|
||||
return true;
|
||||
}
|
||||
if (error != nullptr) {
|
||||
*error = "backend '" + name + "' was not found";
|
||||
}
|
||||
return false;
|
||||
};
|
||||
|
||||
if (!validate_name(runtime_assignment_.default_name) ||
|
||||
!validate_name(params_assignment_.default_name)) {
|
||||
return false;
|
||||
}
|
||||
for (const auto& kv : runtime_assignment_.module_names) {
|
||||
if (!validate_name(kv.second)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
for (const auto& kv : params_assignment_.module_names) {
|
||||
if (!validate_name(kv.second)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
ggml_backend_t SDBackendManager::init_cached_backend(const std::string& name) {
|
||||
std::string resolved = sd_resolve_backend_name(name);
|
||||
std::string key = lower_copy(resolved);
|
||||
ggml_backend_t backend = nullptr;
|
||||
|
||||
if (!key.empty()) {
|
||||
auto it = backends_.find(key);
|
||||
if (it != backends_.end()) {
|
||||
return it->second.get();
|
||||
}
|
||||
} else if (!is_default_backend_token(name)) {
|
||||
LOG_ERROR("backend '%s' was not found", name.c_str());
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
backend = is_default_backend_token(name) ? sd_get_default_backend() : init_named_backend(resolved);
|
||||
if (backend == nullptr) {
|
||||
LOG_ERROR("failed to initialize backend '%s'", name.c_str());
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
std::string actual_key = backend_cache_key(backend);
|
||||
if (actual_key.empty()) {
|
||||
actual_key = !key.empty() ? key : lower_copy(trim_copy(name));
|
||||
}
|
||||
|
||||
auto it = backends_.find(actual_key);
|
||||
if (it != backends_.end()) {
|
||||
ggml_backend_free(backend);
|
||||
return it->second.get();
|
||||
}
|
||||
|
||||
SDBackendHandle handle(backend);
|
||||
backends_.emplace(actual_key, std::move(handle));
|
||||
return backend;
|
||||
}
|
||||
|
||||
const char* sd_backend_module_name(SDBackendModule module) {
|
||||
switch (module) {
|
||||
case SDBackendModule::DIFFUSION:
|
||||
return "diffusion";
|
||||
case SDBackendModule::TE:
|
||||
return "te";
|
||||
case SDBackendModule::CLIP_VISION:
|
||||
return "clip_vision";
|
||||
case SDBackendModule::VAE:
|
||||
return "vae";
|
||||
case SDBackendModule::CONTROL_NET:
|
||||
return "controlnet";
|
||||
case SDBackendModule::PHOTOMAKER:
|
||||
return "photomaker";
|
||||
case SDBackendModule::UPSCALER:
|
||||
return "upscaler";
|
||||
}
|
||||
return "unknown";
|
||||
}
|
||||
@ -1,77 +0,0 @@
|
||||
#ifndef __SD_GGML_EXTEND_BACKEND_H__
|
||||
#define __SD_GGML_EXTEND_BACKEND_H__
|
||||
|
||||
#include <cstdint>
|
||||
#include <cstring>
|
||||
#include <memory>
|
||||
#include <string>
|
||||
#include <unordered_map>
|
||||
|
||||
#include "ggml-backend.h"
|
||||
#include "ggml-cpu.h"
|
||||
#include "ggml.h"
|
||||
|
||||
enum class SDBackendModule {
|
||||
DIFFUSION,
|
||||
TE,
|
||||
CLIP_VISION,
|
||||
VAE,
|
||||
CONTROL_NET,
|
||||
PHOTOMAKER,
|
||||
UPSCALER,
|
||||
};
|
||||
|
||||
struct SDBackendAssignment {
|
||||
std::string default_name;
|
||||
std::unordered_map<SDBackendModule, std::string> module_names;
|
||||
|
||||
bool empty() const;
|
||||
std::string get(SDBackendModule module) const;
|
||||
void set_default(const std::string& name);
|
||||
void set_module(SDBackendModule module, const std::string& name);
|
||||
};
|
||||
|
||||
struct SDBackendHandleDeleter {
|
||||
void operator()(ggml_backend_t backend) const;
|
||||
};
|
||||
|
||||
using SDBackendHandle = std::unique_ptr<struct ggml_backend, SDBackendHandleDeleter>;
|
||||
|
||||
class SDBackendManager {
|
||||
private:
|
||||
SDBackendAssignment runtime_assignment_;
|
||||
SDBackendAssignment params_assignment_;
|
||||
std::unordered_map<std::string, SDBackendHandle> backends_;
|
||||
|
||||
public:
|
||||
SDBackendManager() = default;
|
||||
~SDBackendManager();
|
||||
|
||||
SDBackendManager(const SDBackendManager&) = delete;
|
||||
SDBackendManager& operator=(const SDBackendManager&) = delete;
|
||||
|
||||
bool init(const char* backend_spec,
|
||||
const char* params_backend_spec,
|
||||
bool offload_params_to_cpu,
|
||||
bool keep_clip_on_cpu,
|
||||
bool keep_vae_on_cpu,
|
||||
bool keep_control_net_on_cpu,
|
||||
std::string* error);
|
||||
void reset();
|
||||
|
||||
ggml_backend_t runtime_backend(SDBackendModule module);
|
||||
ggml_backend_t params_backend(SDBackendModule module);
|
||||
|
||||
bool runtime_backend_is_cpu(SDBackendModule module);
|
||||
bool params_backend_is_cpu(SDBackendModule module);
|
||||
bool runtime_backend_supports_host_buffer(SDBackendModule module);
|
||||
|
||||
private:
|
||||
bool validate(std::string* error) const;
|
||||
ggml_backend_t init_cached_backend(const std::string& name);
|
||||
};
|
||||
|
||||
bool sd_backend_is(ggml_backend_t backend, const std::string& name);
|
||||
const char* sd_backend_module_name(SDBackendModule module);
|
||||
void ggml_ext_im_set_f32_1d(const struct ggml_tensor* tensor, int i, float value);
|
||||
#endif
|
||||
298
src/ggml_extend_backend.hpp
Normal file
298
src/ggml_extend_backend.hpp
Normal file
@ -0,0 +1,298 @@
|
||||
#ifndef __GGML_EXTEND_BACKEND_HPP__
|
||||
#define __GGML_EXTEND_BACKEND_HPP__
|
||||
|
||||
#include <cstring>
|
||||
#include <mutex>
|
||||
|
||||
#include "ggml-backend.h"
|
||||
#include "ggml.h"
|
||||
|
||||
#ifndef __STATIC_INLINE__
|
||||
#define __STATIC_INLINE__ static inline
|
||||
#endif
|
||||
|
||||
inline void ggml_backend_load_all_once() {
|
||||
// If the registry already has devices and the CPU backend is present,
|
||||
// assume either static registration or explicit host-side preloading has
|
||||
// completed and avoid rescanning the default paths.
|
||||
if (ggml_backend_dev_count() > 0 && ggml_backend_reg_by_name("CPU") != nullptr) {
|
||||
return;
|
||||
}
|
||||
// In dynamic-backend mode the backend modules are discovered at runtime,
|
||||
// so we must load them before asking for the CPU backend or its proc table.
|
||||
// If the host preloaded only a subset of backends, allow one default-path
|
||||
// scan so missing modules can still be discovered.
|
||||
static std::once_flag once;
|
||||
std::call_once(once, []() {
|
||||
if (ggml_backend_dev_count() > 0 && ggml_backend_reg_by_name("CPU") != nullptr) {
|
||||
return;
|
||||
}
|
||||
ggml_backend_load_all();
|
||||
});
|
||||
}
|
||||
|
||||
// Do not gate this branch on GGML_CPU or GGML_CPU_ALL_VARIANTS:
|
||||
// those are CMake options used to configure ggml itself, but they are not
|
||||
// exported as PUBLIC compile definitions to stable-diffusion in backend-DL mode.
|
||||
// In practice, this target can reliably see GGML_BACKEND_DL, but not whether
|
||||
// the CPU backend was compiled as a loadable module. We therefore use runtime
|
||||
// backend discovery instead of compile-time assumptions.
|
||||
|
||||
__STATIC_INLINE__ ggml_backend_reg_t ggml_backend_cpu_reg() {
|
||||
ggml_backend_reg_t reg = ggml_backend_reg_by_name("CPU");
|
||||
if (reg != nullptr) {
|
||||
return reg;
|
||||
}
|
||||
|
||||
ggml_backend_load_all_once();
|
||||
return ggml_backend_reg_by_name("CPU");
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_backend_reg_t ggml_backend_reg_from_backend(ggml_backend_t backend) {
|
||||
if (backend != nullptr) {
|
||||
ggml_backend_dev_t device = ggml_backend_get_device(backend);
|
||||
if (device != nullptr) {
|
||||
return ggml_backend_dev_backend_reg(device);
|
||||
}
|
||||
}
|
||||
|
||||
return ggml_backend_cpu_reg();
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_backend_t ggml_backend_cpu_init() {
|
||||
ggml_backend_t backend = ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr);
|
||||
if (backend != nullptr) {
|
||||
return backend;
|
||||
}
|
||||
|
||||
ggml_backend_load_all_once();
|
||||
return ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ bool ggml_backend_is_cpu(ggml_backend_t backend) {
|
||||
if (backend == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
ggml_backend_dev_t device = ggml_backend_get_device(backend);
|
||||
if (device != nullptr) {
|
||||
return ggml_backend_dev_type(device) == GGML_BACKEND_DEVICE_TYPE_CPU;
|
||||
}
|
||||
|
||||
const char* backend_name = ggml_backend_name(backend);
|
||||
return backend_name != nullptr && std::strcmp(backend_name, "CPU") == 0;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void ggml_backend_cpu_set_n_threads(ggml_backend_t backend_cpu, int n_threads) {
|
||||
ggml_backend_reg_t reg = ggml_backend_reg_from_backend(backend_cpu);
|
||||
if (reg == nullptr) {
|
||||
return;
|
||||
}
|
||||
|
||||
auto fn = reinterpret_cast<ggml_backend_set_n_threads_t>(ggml_backend_reg_get_proc_address(reg, "ggml_backend_set_n_threads"));
|
||||
if (fn != nullptr) {
|
||||
fn(backend_cpu, n_threads);
|
||||
}
|
||||
}
|
||||
|
||||
using __ggml_backend_cpu_set_threadpool_t = void (*)(ggml_backend_t backend_cpu, ggml_threadpool_t threadpool);
|
||||
|
||||
__STATIC_INLINE__ void ggml_backend_cpu_set_threadpool(ggml_backend_t backend_cpu, ggml_threadpool_t threadpool) {
|
||||
ggml_backend_reg_t reg = ggml_backend_reg_from_backend(backend_cpu);
|
||||
if (reg == nullptr) {
|
||||
return;
|
||||
}
|
||||
|
||||
auto fn = reinterpret_cast<__ggml_backend_cpu_set_threadpool_t>(ggml_backend_reg_get_proc_address(reg, "ggml_backend_cpu_set_threadpool"));
|
||||
if (fn != nullptr) {
|
||||
fn(backend_cpu, threadpool);
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void ggml_backend_cpu_set_abort_callback(ggml_backend_t backend_cpu, ggml_abort_callback abort_callback, void* abort_callback_data) {
|
||||
ggml_backend_reg_t reg = ggml_backend_reg_from_backend(backend_cpu);
|
||||
if (reg == nullptr) {
|
||||
return;
|
||||
}
|
||||
|
||||
auto fn = reinterpret_cast<ggml_backend_set_abort_callback_t>(ggml_backend_reg_get_proc_address(reg, "ggml_backend_set_abort_callback"));
|
||||
if (fn != nullptr) {
|
||||
fn(backend_cpu, abort_callback, abort_callback_data);
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_backend_buffer_t ggml_backend_tensor_buffer(const struct ggml_tensor* tensor) {
|
||||
if (tensor == nullptr) {
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
return tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ bool ggml_backend_tensor_is_host_accessible(const struct ggml_tensor* tensor) {
|
||||
if (tensor == nullptr || tensor->data == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
ggml_backend_buffer_t buffer = ggml_backend_tensor_buffer(tensor);
|
||||
return buffer == nullptr || ggml_backend_buffer_is_host(buffer);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ size_t ggml_backend_tensor_offset(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3) {
|
||||
return (size_t)(i0 * tensor->nb[0] + i1 * tensor->nb[1] + i2 * tensor->nb[2] + i3 * tensor->nb[3]);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__STATIC_INLINE__ void ggml_backend_tensor_write_scalar(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3, T value) {
|
||||
const size_t offset = ggml_backend_tensor_offset(tensor, i0, i1, i2, i3);
|
||||
|
||||
if (ggml_backend_tensor_is_host_accessible(tensor)) {
|
||||
auto* dst = reinterpret_cast<T*>(reinterpret_cast<char*>(tensor->data) + offset);
|
||||
*dst = value;
|
||||
return;
|
||||
}
|
||||
|
||||
ggml_backend_tensor_set(const_cast<struct ggml_tensor*>(tensor), &value, offset, sizeof(T));
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void ggml_set_f32_nd(const struct ggml_tensor* tensor, int64_t i0, int64_t i1, int64_t i2, int64_t i3, float value) {
|
||||
switch (tensor->type) {
|
||||
case GGML_TYPE_I8:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int8_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int16_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, static_cast<int32_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_F16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, ggml_fp32_to_fp16(value));
|
||||
break;
|
||||
case GGML_TYPE_BF16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, ggml_fp32_to_bf16(value));
|
||||
break;
|
||||
case GGML_TYPE_F32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i0, i1, i2, i3, value);
|
||||
break;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void ggml_set_f32_1d(const struct ggml_tensor* tensor, int i, float value) {
|
||||
if (!ggml_is_contiguous(tensor)) {
|
||||
int64_t id[4] = {0, 0, 0, 0};
|
||||
ggml_unravel_index(tensor, i, &id[0], &id[1], &id[2], &id[3]);
|
||||
ggml_set_f32_nd(tensor, id[0], id[1], id[2], id[3], value);
|
||||
return;
|
||||
}
|
||||
|
||||
switch (tensor->type) {
|
||||
case GGML_TYPE_I8:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int8_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int16_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_I32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, static_cast<int32_t>(value));
|
||||
break;
|
||||
case GGML_TYPE_F16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, ggml_fp32_to_fp16(value));
|
||||
break;
|
||||
case GGML_TYPE_BF16:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, ggml_fp32_to_bf16(value));
|
||||
break;
|
||||
case GGML_TYPE_F32:
|
||||
ggml_backend_tensor_write_scalar(tensor, i, 0, 0, 0, value);
|
||||
break;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ enum ggml_status ggml_graph_compute_with_ctx(struct ggml_context* ctx, struct ggml_cgraph* cgraph, int n_threads) {
|
||||
(void)ctx;
|
||||
|
||||
// The legacy ggml_graph_compute_with_ctx() symbol lives in ggml-cpu, but
|
||||
// the backend proc table does not expose it in GGML_BACKEND_DL mode.
|
||||
// Recreate the old behavior by initializing the CPU backend explicitly and
|
||||
// executing the graph through the generic backend API.
|
||||
ggml_backend_t backend = ggml_backend_cpu_init();
|
||||
if (backend == nullptr) {
|
||||
return GGML_STATUS_ALLOC_FAILED;
|
||||
}
|
||||
|
||||
ggml_backend_cpu_set_n_threads(backend, n_threads);
|
||||
|
||||
const enum ggml_status status = ggml_backend_graph_compute(backend, cgraph);
|
||||
ggml_backend_free(backend);
|
||||
|
||||
return status;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_tensor* ggml_set_f32(struct ggml_tensor* tensor, float value) {
|
||||
GGML_ASSERT(tensor != nullptr);
|
||||
|
||||
if (ggml_backend_tensor_is_host_accessible(tensor) && ggml_is_contiguous(tensor)) {
|
||||
const int64_t nelements = ggml_nelements(tensor);
|
||||
|
||||
switch (tensor->type) {
|
||||
case GGML_TYPE_I8: {
|
||||
auto* data = reinterpret_cast<int8_t*>(tensor->data);
|
||||
const int8_t v = static_cast<int8_t>(value);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = v;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_I16: {
|
||||
auto* data = reinterpret_cast<int16_t*>(tensor->data);
|
||||
const int16_t v = static_cast<int16_t>(value);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = v;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_I32: {
|
||||
auto* data = reinterpret_cast<int32_t*>(tensor->data);
|
||||
const int32_t v = static_cast<int32_t>(value);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = v;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_F16: {
|
||||
auto* data = reinterpret_cast<ggml_fp16_t*>(tensor->data);
|
||||
const ggml_fp16_t v = ggml_fp32_to_fp16(value);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = v;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_BF16: {
|
||||
auto* data = reinterpret_cast<ggml_bf16_t*>(tensor->data);
|
||||
const ggml_bf16_t v = ggml_fp32_to_bf16(value);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = v;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_F32: {
|
||||
auto* data = reinterpret_cast<float*>(tensor->data);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
data[i] = value;
|
||||
}
|
||||
} break;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
|
||||
return tensor;
|
||||
}
|
||||
|
||||
const int64_t nelements = ggml_nelements(tensor);
|
||||
for (int64_t i = 0; i < nelements; ++i) {
|
||||
ggml_set_f32_1d(tensor, static_cast<int>(i), value);
|
||||
}
|
||||
|
||||
return tensor;
|
||||
}
|
||||
|
||||
#endif
|
||||
@ -279,10 +279,10 @@ namespace HiDreamO1 {
|
||||
std::array<std::vector<float>, 4> pos_embed_weight_data_;
|
||||
|
||||
HiDreamO1VisionRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& prefix = "model.visual")
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
params(make_hidream_o1_params()),
|
||||
model(std::make_shared<LLM::VisionModel>(false, params.llm.vision)) {
|
||||
model->init(params_ctx, tensor_storage_map, prefix);
|
||||
@ -336,10 +336,10 @@ namespace HiDreamO1 {
|
||||
std::vector<float> attention_mask_vec;
|
||||
|
||||
HiDreamO1Runner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& prefix = "model")
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
params(make_hidream_o1_params()) {
|
||||
model = HiDreamO1Model(params);
|
||||
model.init(params_ctx, tensor_storage_map, prefix);
|
||||
@ -461,9 +461,9 @@ namespace HiDreamO1 {
|
||||
std::shared_ptr<HiDreamO1VisionRunner> vision_runner;
|
||||
|
||||
HiDreamO1Conditioner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {})
|
||||
: vision_runner(std::make_shared<HiDreamO1VisionRunner>(backend, params_backend, tensor_storage_map)) {}
|
||||
: vision_runner(std::make_shared<HiDreamO1VisionRunner>(backend, offload_params_to_cpu, tensor_storage_map)) {}
|
||||
|
||||
void get_param_tensors(std::map<std::string, ggml_tensor*>& tensors) override {
|
||||
vision_runner->get_param_tensors(tensors);
|
||||
|
||||
10
src/llm.hpp
10
src/llm.hpp
@ -1166,11 +1166,11 @@ namespace LLM {
|
||||
public:
|
||||
LLMRunner(LLMArch arch,
|
||||
ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
bool enable_vision_ = false)
|
||||
: GGMLRunner(backend, params_backend), enable_vision(enable_vision_) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu), enable_vision(enable_vision_) {
|
||||
params.arch = arch;
|
||||
if (arch == LLMArch::MISTRAL_SMALL_3_2 || arch == LLMArch::MINISTRAL_3_3B) {
|
||||
params.head_dim = 128;
|
||||
@ -1477,11 +1477,11 @@ namespace LLM {
|
||||
|
||||
LLMEmbedder(LLMArch arch,
|
||||
ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
bool enable_vision = false)
|
||||
: model(arch, backend, params_backend, tensor_storage_map, prefix, enable_vision) {
|
||||
: model(arch, backend, offload_params_to_cpu, tensor_storage_map, prefix, enable_vision) {
|
||||
if (arch == LLMArch::MISTRAL_SMALL_3_2 || arch == LLMArch::MINISTRAL_3_3B) {
|
||||
tokenizer = std::make_shared<MistralTokenizer>();
|
||||
} else {
|
||||
@ -1731,7 +1731,7 @@ namespace LLM {
|
||||
|
||||
std::shared_ptr<LLMEmbedder> llm = std::make_shared<LLMEmbedder>(arch,
|
||||
backend,
|
||||
backend,
|
||||
true,
|
||||
tensor_storage_map,
|
||||
"text_encoders.llm",
|
||||
true);
|
||||
|
||||
@ -22,11 +22,10 @@ struct LoraModel : public GGMLRunner {
|
||||
|
||||
LoraModel(const std::string& lora_id,
|
||||
ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
const std::string& file_path = "",
|
||||
std::string prefix = "",
|
||||
SDVersion version = VERSION_COUNT)
|
||||
: lora_id(lora_id), file_path(file_path), GGMLRunner(backend, params_backend) {
|
||||
: lora_id(lora_id), file_path(file_path), GGMLRunner(backend, false) {
|
||||
prefix = "lora." + prefix;
|
||||
if (!model_loader.init_from_file_and_convert_name(file_path, prefix, version)) {
|
||||
load_failed = true;
|
||||
|
||||
@ -963,10 +963,10 @@ namespace LTXV {
|
||||
LTXAudioVAE model;
|
||||
|
||||
LTXAudioVAERunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& prefix = "")
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
config(LTXAudioVAEConfig::detect_from_weights(tensor_storage_map)),
|
||||
model(config) {
|
||||
model.init(params_ctx, tensor_storage_map, prefix);
|
||||
@ -1086,7 +1086,7 @@ namespace LTXV {
|
||||
|
||||
auto& tensor_storage_map = model_loader.get_tensor_storage_map();
|
||||
auto ltx_audio_vae = std::make_shared<LTXAudioVAERunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
prefix);
|
||||
|
||||
|
||||
132
src/ltx_vae.hpp
132
src/ltx_vae.hpp
@ -89,8 +89,7 @@ namespace LTXVAE {
|
||||
int kernel_size = 3,
|
||||
std::tuple<int, int, int> stride = {1, 1, 1},
|
||||
int dilation = 1,
|
||||
bool bias = true,
|
||||
bool force_prec_f32 = false) {
|
||||
bool bias = true) {
|
||||
time_kernel_size = kernel_size;
|
||||
blocks["conv"] = std::shared_ptr<GGMLBlock>(new Conv3d(in_channels,
|
||||
out_channels,
|
||||
@ -98,8 +97,7 @@ namespace LTXVAE {
|
||||
stride,
|
||||
{0, kernel_size / 2, kernel_size / 2},
|
||||
{dilation, 1, 1},
|
||||
bias,
|
||||
force_prec_f32));
|
||||
bias));
|
||||
}
|
||||
|
||||
ggml_tensor* forward(GGMLRunnerContext* ctx,
|
||||
@ -471,8 +469,7 @@ namespace LTXVAE {
|
||||
SpaceToDepthDownsample(int64_t in_channels,
|
||||
int64_t out_channels,
|
||||
int factor_t,
|
||||
int factor_s,
|
||||
bool force_conv_prec_f32 = false)
|
||||
int factor_s)
|
||||
: in_channels(in_channels),
|
||||
out_channels(out_channels),
|
||||
factor_t(factor_t),
|
||||
@ -480,13 +477,7 @@ namespace LTXVAE {
|
||||
const int64_t factor = static_cast<int64_t>(factor_t) * static_cast<int64_t>(factor_s) * static_cast<int64_t>(factor_s);
|
||||
GGML_ASSERT(out_channels % factor == 0);
|
||||
|
||||
blocks["conv"] = std::make_shared<CausalConv3d>(in_channels,
|
||||
out_channels / factor,
|
||||
3,
|
||||
std::tuple<int, int, int>{1, 1, 1},
|
||||
1,
|
||||
true,
|
||||
force_conv_prec_f32);
|
||||
blocks["conv"] = std::make_shared<CausalConv3d>(in_channels, out_channels / factor, 3);
|
||||
blocks["skip_downsample"] = std::make_shared<WAN::AvgDown3D>(in_channels, out_channels, factor_t, factor_s);
|
||||
blocks["conv_downsample"] = std::make_shared<WAN::AvgDown3D>(out_channels / factor, out_channels, factor_t, factor_s);
|
||||
}
|
||||
@ -501,7 +492,7 @@ namespace LTXVAE {
|
||||
if (factor_t > 1 && x->ne[2] > 0) {
|
||||
auto first_frame = ggml_ext_slice(ctx->ggml_ctx, x, 2, 0, 1);
|
||||
auto first_frame_pad = first_frame;
|
||||
for (int i = 1; i < factor_t - 1; ++i) {
|
||||
for (int i = 1; i < factor_t; ++i) {
|
||||
first_frame_pad = ggml_concat(ctx->ggml_ctx, first_frame_pad, first_frame, 2);
|
||||
}
|
||||
x = ggml_concat(ctx->ggml_ctx, first_frame_pad, x, 2);
|
||||
@ -559,8 +550,6 @@ namespace LTXVAE {
|
||||
std::vector<Block> blocks;
|
||||
};
|
||||
|
||||
static inline EncoderConfig get_default_encoder_config(int version);
|
||||
|
||||
static inline bool has_tensor(const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& name) {
|
||||
return tensor_storage_map.find(name) != tensor_storage_map.end();
|
||||
@ -644,84 +633,6 @@ namespace LTXVAE {
|
||||
return cfg;
|
||||
}
|
||||
|
||||
static inline EncoderConfig infer_encoder_config_from_weights(const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& prefix,
|
||||
int version) {
|
||||
EncoderConfig cfg;
|
||||
const std::string encoder_prefix = prefix + ".encoder.down_blocks.";
|
||||
|
||||
int64_t current_channels = get_tensor_ne0(tensor_storage_map,
|
||||
prefix + ".encoder.conv_in.conv.bias",
|
||||
0);
|
||||
for (int block_idx = 0;; ++block_idx) {
|
||||
const std::string block_prefix = encoder_prefix + std::to_string(block_idx);
|
||||
const std::string res0_bias = block_prefix + ".res_blocks.0.conv1.conv.bias";
|
||||
const std::string conv_bias = block_prefix + ".conv.conv.bias";
|
||||
|
||||
if (has_tensor(tensor_storage_map, res0_bias)) {
|
||||
int num_layers = 0;
|
||||
while (has_tensor(tensor_storage_map,
|
||||
block_prefix + ".res_blocks." + std::to_string(num_layers) + ".conv1.conv.bias")) {
|
||||
num_layers++;
|
||||
}
|
||||
cfg.blocks.push_back({"res_x", num_layers, 1});
|
||||
current_channels = get_tensor_ne0(tensor_storage_map, res0_bias, current_channels);
|
||||
continue;
|
||||
}
|
||||
|
||||
if (!has_tensor(tensor_storage_map, conv_bias)) {
|
||||
break;
|
||||
}
|
||||
|
||||
const int64_t conv_channels = get_tensor_ne0(tensor_storage_map, conv_bias);
|
||||
int64_t next_channels = 0;
|
||||
for (int next_idx = block_idx + 1;; ++next_idx) {
|
||||
const std::string next_res0_bias = encoder_prefix + std::to_string(next_idx) + ".res_blocks.0.conv1.conv.bias";
|
||||
const std::string next_conv_bias = encoder_prefix + std::to_string(next_idx) + ".conv.conv.bias";
|
||||
if (has_tensor(tensor_storage_map, next_res0_bias)) {
|
||||
next_channels = get_tensor_ne0(tensor_storage_map, next_res0_bias);
|
||||
break;
|
||||
}
|
||||
if (!has_tensor(tensor_storage_map, next_conv_bias)) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
const int64_t multiplier = current_channels > 0 && next_channels > 0 && next_channels % current_channels == 0
|
||||
? std::max<int64_t>(1, next_channels / current_channels)
|
||||
: 1;
|
||||
const int64_t factor = conv_channels > 0 && next_channels > 0 && next_channels % conv_channels == 0
|
||||
? next_channels / conv_channels
|
||||
: 0;
|
||||
|
||||
if (factor == 8) {
|
||||
cfg.blocks.push_back({"compress_all_res", 0, static_cast<int>(multiplier)});
|
||||
} else if (factor == 4) {
|
||||
cfg.blocks.push_back({"compress_space_res", 0, static_cast<int>(multiplier)});
|
||||
} else if (factor == 2) {
|
||||
cfg.blocks.push_back({"compress_time_res", 0, static_cast<int>(multiplier)});
|
||||
} else {
|
||||
LOG_WARN("unexpected LTX VAE encoder downsample factor at '%s': conv_out=%lld current=%lld next=%lld, falling back to compress_all_res x%d",
|
||||
block_prefix.c_str(),
|
||||
(long long)conv_channels,
|
||||
(long long)current_channels,
|
||||
(long long)next_channels,
|
||||
(int)multiplier);
|
||||
cfg.blocks.push_back({"compress_all_res", 0, static_cast<int>(multiplier)});
|
||||
}
|
||||
if (next_channels > 0) {
|
||||
current_channels = next_channels;
|
||||
} else if (current_channels > 0) {
|
||||
current_channels *= multiplier;
|
||||
}
|
||||
}
|
||||
|
||||
if (cfg.blocks.empty()) {
|
||||
return get_default_encoder_config(version);
|
||||
}
|
||||
return cfg;
|
||||
}
|
||||
|
||||
static inline int detect_ltx_vae_version(const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& prefix) {
|
||||
const std::string v2_probe = prefix + ".encoder.down_blocks.1.conv.conv.bias";
|
||||
@ -736,7 +647,7 @@ namespace LTXVAE {
|
||||
return tensor_storage_map.find(prefix + ".decoder.timestep_scale_multiplier") != tensor_storage_map.end();
|
||||
}
|
||||
|
||||
static inline EncoderConfig get_default_encoder_config(int version) {
|
||||
static inline EncoderConfig get_encoder_config(int version) {
|
||||
EncoderConfig cfg;
|
||||
if (version < 2) {
|
||||
GGML_ABORT("LTX VAE encoder is only implemented for version >= 2");
|
||||
@ -763,8 +674,6 @@ namespace LTXVAE {
|
||||
int64_t latent_channels;
|
||||
|
||||
Encoder(int version,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& prefix,
|
||||
int patch_size = 4,
|
||||
int64_t in_channels = 3,
|
||||
int64_t latent_channels = 128)
|
||||
@ -772,12 +681,9 @@ namespace LTXVAE {
|
||||
patch_size(patch_size),
|
||||
in_channels(in_channels),
|
||||
latent_channels(latent_channels) {
|
||||
auto cfg = infer_encoder_config_from_weights(tensor_storage_map, prefix, version);
|
||||
int64_t channels = get_tensor_ne0(tensor_storage_map,
|
||||
prefix + ".encoder.conv_in.conv.bias",
|
||||
0);
|
||||
GGML_ASSERT(channels > 0);
|
||||
int64_t in_dim = in_channels * patch_size * patch_size;
|
||||
auto cfg = get_encoder_config(version);
|
||||
int64_t channels = 128;
|
||||
int64_t in_dim = in_channels * patch_size * patch_size;
|
||||
|
||||
blocks["conv_in"] = std::make_shared<CausalConv3d>(in_dim, channels, 3);
|
||||
|
||||
@ -802,14 +708,11 @@ namespace LTXVAE {
|
||||
1);
|
||||
channels = next_channels;
|
||||
} else if (block.type == "compress_all_res") {
|
||||
int64_t next_channels = channels * block.multiplier;
|
||||
// LTX 2.3 encoder down_blocks.7.conv overflows with fp16 accumulation.
|
||||
bool force_conv_prec_f32 = block_idx == 7;
|
||||
int64_t next_channels = channels * block.multiplier;
|
||||
blocks["down_blocks." + std::to_string(block_idx)] = std::make_shared<SpaceToDepthDownsample>(channels,
|
||||
next_channels,
|
||||
2,
|
||||
2,
|
||||
force_conv_prec_f32);
|
||||
2);
|
||||
channels = next_channels;
|
||||
} else {
|
||||
GGML_ABORT("Unsupported LTX VAE encoder block");
|
||||
@ -1053,10 +956,7 @@ namespace LTXVAE {
|
||||
patch_size(patch_size),
|
||||
decode_only(decode_only) {
|
||||
if (!decode_only) {
|
||||
blocks["encoder"] = std::make_shared<Encoder>(version,
|
||||
tensor_storage_map,
|
||||
prefix,
|
||||
patch_size);
|
||||
blocks["encoder"] = std::make_shared<Encoder>(version, patch_size);
|
||||
}
|
||||
blocks["decoder"] = std::make_shared<Decoder>(version,
|
||||
tensor_storage_map,
|
||||
@ -1137,7 +1037,7 @@ struct LTXVideoVAE : public VAE {
|
||||
LTXVAE::VideoVAE vae;
|
||||
|
||||
LTXVideoVAE(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string& prefix,
|
||||
bool decode_only = true,
|
||||
@ -1153,7 +1053,7 @@ struct LTXVideoVAE : public VAE {
|
||||
patch_size,
|
||||
tensor_storage_map,
|
||||
prefix),
|
||||
VAE(version, backend, params_backend) {
|
||||
VAE(version, backend, offload_params_to_cpu) {
|
||||
vae.init(params_ctx, tensor_storage_map, prefix);
|
||||
decode_timestep_tensor.values()[0] = vae.decode_timestep;
|
||||
}
|
||||
@ -1196,7 +1096,7 @@ struct LTXVideoVAE : public VAE {
|
||||
const sd::Tensor<float>& z,
|
||||
bool decode_graph) override {
|
||||
if (!decode_graph && decode_only) {
|
||||
LOG_ERROR("LTX video VAE encode requires encoder weights; create the context with vae_decode_only=false");
|
||||
LOG_ERROR("LTX video VAE encoder is not implemented yet");
|
||||
return {};
|
||||
}
|
||||
sd::Tensor<float> input = z;
|
||||
@ -1276,7 +1176,7 @@ struct LTXVideoVAE : public VAE {
|
||||
|
||||
auto& tensor_storage_map = model_loader.get_tensor_storage_map();
|
||||
std::shared_ptr<LTXVideoVAE> vae = std::make_shared<LTXVideoVAE>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"first_stage_model",
|
||||
true,
|
||||
|
||||
82
src/ltxv.hpp
82
src/ltxv.hpp
@ -23,28 +23,12 @@ namespace LTXV {
|
||||
return ggml_rms_norm(ctx, x, eps);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_tensor* align_token_modulation(ggml_context* ctx,
|
||||
ggml_tensor* x,
|
||||
ggml_tensor* mod) {
|
||||
if (mod != nullptr && x != nullptr && mod->ne[1] == 1 && mod->ne[2] == x->ne[1] && x->ne[2] == 1) {
|
||||
return ggml_permute(ctx, mod, 0, 2, 1, 3);
|
||||
}
|
||||
return mod;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_tensor* modulate(ggml_context* ctx,
|
||||
ggml_tensor* x,
|
||||
ggml_tensor* shift,
|
||||
ggml_tensor* scale) {
|
||||
shift = align_token_modulation(ctx, x, shift);
|
||||
scale = align_token_modulation(ctx, x, scale);
|
||||
return Flux::modulate(ctx, x, shift, scale, true);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ ggml_tensor* apply_gate(ggml_context* ctx,
|
||||
ggml_tensor* x,
|
||||
ggml_tensor* gate) {
|
||||
gate = align_token_modulation(ctx, x, gate);
|
||||
if (gate->ne[1] != 1) {
|
||||
gate = ggml_reshape_3d(ctx, gate, gate->ne[0], 1, gate->ne[2]);
|
||||
}
|
||||
return ggml_mul(ctx, x, gate);
|
||||
}
|
||||
|
||||
@ -554,7 +538,7 @@ namespace LTXV {
|
||||
auto gate_mlp = mods[5];
|
||||
|
||||
auto x_norm = rms_norm(ctx->ggml_ctx, x);
|
||||
x_norm = modulate(ctx->ggml_ctx, x_norm, shift_msa, scale_msa);
|
||||
x_norm = Flux::modulate(ctx->ggml_ctx, x_norm, shift_msa, scale_msa, true);
|
||||
auto msa = attn1->forward(ctx, x_norm, nullptr, self_attention_mask, pe);
|
||||
x = ggml_add(ctx->ggml_ctx, x, apply_gate(ctx->ggml_ctx, msa, gate_msa));
|
||||
|
||||
@ -564,12 +548,12 @@ namespace LTXV {
|
||||
auto gate_q = mods[8];
|
||||
|
||||
auto q = rms_norm(ctx->ggml_ctx, x);
|
||||
q = modulate(ctx->ggml_ctx, q, shift_q, scale_q);
|
||||
q = Flux::modulate(ctx->ggml_ctx, q, shift_q, scale_q, true);
|
||||
|
||||
auto context_mod = context;
|
||||
if (prompt_timestep != nullptr) {
|
||||
auto prompt_mods = get_prompt_scale_shift_values(ctx, prompt_timestep);
|
||||
context_mod = modulate(ctx->ggml_ctx, context_mod, prompt_mods[0], prompt_mods[1]);
|
||||
context_mod = Flux::modulate(ctx->ggml_ctx, context_mod, prompt_mods[0], prompt_mods[1], true);
|
||||
}
|
||||
|
||||
auto mca = attn2->forward(ctx, q, context_mod, attention_mask, nullptr, nullptr);
|
||||
@ -580,7 +564,7 @@ namespace LTXV {
|
||||
}
|
||||
|
||||
auto y = rms_norm(ctx->ggml_ctx, x);
|
||||
y = modulate(ctx->ggml_ctx, y, shift_mlp, scale_mlp);
|
||||
y = Flux::modulate(ctx->ggml_ctx, y, shift_mlp, scale_mlp, true);
|
||||
auto mlp_out = ff->forward(ctx, y);
|
||||
x = ggml_add(ctx->ggml_ctx, x, apply_gate(ctx->ggml_ctx, mlp_out, gate_mlp));
|
||||
return x;
|
||||
@ -963,11 +947,11 @@ namespace LTXV {
|
||||
if (cross_attention_adaln) {
|
||||
auto q_mods = get_ada_values(ctx, table, timestep, dim, 9, 6, 3);
|
||||
auto q = rms_norm(ctx->ggml_ctx, x);
|
||||
q = modulate(ctx->ggml_ctx, q, q_mods[0], q_mods[1]);
|
||||
q = Flux::modulate(ctx->ggml_ctx, q, q_mods[0], q_mods[1], true);
|
||||
auto context_mod = context;
|
||||
if (prompt_timestep != nullptr && prompt_table != nullptr) {
|
||||
auto p_mods = get_ada_values(ctx, prompt_table, prompt_timestep, dim, 2);
|
||||
context_mod = modulate(ctx->ggml_ctx, context_mod, p_mods[0], p_mods[1]);
|
||||
context_mod = Flux::modulate(ctx->ggml_ctx, context_mod, p_mods[0], p_mods[1], true);
|
||||
}
|
||||
auto out = attn->forward(ctx, q, context_mod, attention_mask, nullptr, nullptr);
|
||||
return apply_gate(ctx->ggml_ctx, out, q_mods[2]);
|
||||
@ -1014,7 +998,7 @@ namespace LTXV {
|
||||
|
||||
auto v_mods = get_ada_values(ctx, v_table, v_timestep, v_dim, cross_attention_adaln ? 9 : 6);
|
||||
auto v_norm = rms_norm(ctx->ggml_ctx, vx);
|
||||
v_norm = modulate(ctx->ggml_ctx, v_norm, v_mods[0], v_mods[1]);
|
||||
v_norm = Flux::modulate(ctx->ggml_ctx, v_norm, v_mods[0], v_mods[1], true);
|
||||
auto v_sa = attn1->forward(ctx, v_norm, nullptr, self_attention_mask, v_pe);
|
||||
vx = ggml_add(ctx->ggml_ctx, vx, apply_gate(ctx->ggml_ctx, v_sa, v_mods[2]));
|
||||
auto v_txt = apply_text_cross_attention(ctx,
|
||||
@ -1032,7 +1016,7 @@ namespace LTXV {
|
||||
if (run_ax) {
|
||||
auto a_mods = get_ada_values(ctx, a_table, a_timestep, a_dim, cross_attention_adaln ? 9 : 6);
|
||||
auto a_norm = rms_norm(ctx->ggml_ctx, ax);
|
||||
a_norm = modulate(ctx->ggml_ctx, a_norm, a_mods[0], a_mods[1]);
|
||||
a_norm = Flux::modulate(ctx->ggml_ctx, a_norm, a_mods[0], a_mods[1], true);
|
||||
auto a_sa = audio_attn1->forward(ctx, a_norm, nullptr, nullptr, a_pe);
|
||||
ax = ggml_add(ctx->ggml_ctx, ax, apply_gate(ctx->ggml_ctx, a_sa, a_mods[2]));
|
||||
auto a_txt = apply_text_cross_attention(ctx,
|
||||
@ -1055,8 +1039,8 @@ namespace LTXV {
|
||||
auto a2v_video_table = ggml_ext_slice(ctx->ggml_ctx, params["scale_shift_table_a2v_ca_video"], 1, 0, 4);
|
||||
auto a2v_audio = get_ada_values(ctx, a2v_audio_table, a_cross_scale_shift_timestep, a_dim, 4);
|
||||
auto a2v_video = get_ada_values(ctx, a2v_video_table, v_cross_scale_shift_timestep, v_dim, 4);
|
||||
auto vx_scaled = modulate(ctx->ggml_ctx, vx_norm3, a2v_video[1], a2v_video[0]);
|
||||
auto ax_scaled = modulate(ctx->ggml_ctx, ax_norm3, a2v_audio[1], a2v_audio[0]);
|
||||
auto vx_scaled = Flux::modulate(ctx->ggml_ctx, vx_norm3, a2v_video[1], a2v_video[0], true);
|
||||
auto ax_scaled = Flux::modulate(ctx->ggml_ctx, ax_norm3, a2v_audio[1], a2v_audio[0], true);
|
||||
auto a2v_out = audio_to_video_attn->forward(ctx, vx_scaled, ax_scaled, nullptr, v_cross_pe, a_cross_pe);
|
||||
auto a2v_gate_table = ggml_ext_slice(ctx->ggml_ctx, params["scale_shift_table_a2v_ca_video"], 1, 4, 5);
|
||||
auto a2v_gate = get_ada_values(ctx, a2v_gate_table, v_cross_gate_timestep, v_dim, 1)[0];
|
||||
@ -1068,8 +1052,8 @@ namespace LTXV {
|
||||
auto v2a_video_table = ggml_ext_slice(ctx->ggml_ctx, params["scale_shift_table_a2v_ca_video"], 1, 0, 4);
|
||||
auto v2a_audio = get_ada_values(ctx, v2a_audio_table, a_cross_scale_shift_timestep, a_dim, 4);
|
||||
auto v2a_video = get_ada_values(ctx, v2a_video_table, v_cross_scale_shift_timestep, v_dim, 4);
|
||||
auto ax_scaled = modulate(ctx->ggml_ctx, ax_norm3, v2a_audio[3], v2a_audio[2]);
|
||||
auto vx_scaled = modulate(ctx->ggml_ctx, vx_norm3, v2a_video[3], v2a_video[2]);
|
||||
auto ax_scaled = Flux::modulate(ctx->ggml_ctx, ax_norm3, v2a_audio[3], v2a_audio[2], true);
|
||||
auto vx_scaled = Flux::modulate(ctx->ggml_ctx, vx_norm3, v2a_video[3], v2a_video[2], true);
|
||||
auto v2a_out = video_to_audio_attn->forward(ctx, ax_scaled, vx_scaled, nullptr, a_cross_pe, v_cross_pe);
|
||||
auto v2a_gate_table = ggml_ext_slice(ctx->ggml_ctx, params["scale_shift_table_a2v_ca_audio"], 1, 4, 5);
|
||||
auto v2a_gate = get_ada_values(ctx, v2a_gate_table, a_cross_gate_timestep, a_dim, 1)[0];
|
||||
@ -1077,14 +1061,14 @@ namespace LTXV {
|
||||
}
|
||||
auto a_ff_mods = get_ada_values(ctx, a_table, a_timestep, a_dim, cross_attention_adaln ? 9 : 6, 3, 3);
|
||||
auto ax_scaled = rms_norm(ctx->ggml_ctx, ax);
|
||||
ax_scaled = modulate(ctx->ggml_ctx, ax_scaled, a_ff_mods[0], a_ff_mods[1]);
|
||||
ax_scaled = Flux::modulate(ctx->ggml_ctx, ax_scaled, a_ff_mods[0], a_ff_mods[1], true);
|
||||
auto a_ff_out = audio_ff->forward(ctx, ax_scaled);
|
||||
ax = ggml_add(ctx->ggml_ctx, ax, apply_gate(ctx->ggml_ctx, a_ff_out, a_ff_mods[2]));
|
||||
}
|
||||
|
||||
auto v_ff_mods = get_ada_values(ctx, v_table, v_timestep, v_dim, cross_attention_adaln ? 9 : 6, 3, 3);
|
||||
auto vx_scaled = rms_norm(ctx->ggml_ctx, vx);
|
||||
vx_scaled = modulate(ctx->ggml_ctx, vx_scaled, v_ff_mods[0], v_ff_mods[1]);
|
||||
vx_scaled = Flux::modulate(ctx->ggml_ctx, vx_scaled, v_ff_mods[0], v_ff_mods[1], true);
|
||||
auto v_ff_out = ff->forward(ctx, vx_scaled);
|
||||
vx = ggml_add(ctx->ggml_ctx, vx, apply_gate(ctx->ggml_ctx, v_ff_out, v_ff_mods[2]));
|
||||
|
||||
@ -1204,15 +1188,6 @@ namespace LTXV {
|
||||
return ax;
|
||||
}
|
||||
|
||||
ggml_tensor* repeat_scalar_timestep_like(GGMLRunnerContext* ctx, ggml_tensor* timestep, ggml_tensor* like) {
|
||||
GGML_ASSERT(timestep != nullptr && like != nullptr);
|
||||
if (timestep->ne[0] == like->ne[0]) {
|
||||
return timestep;
|
||||
}
|
||||
GGML_ASSERT(timestep->ne[0] == 1);
|
||||
return ggml_repeat(ctx->ggml_ctx, timestep, ggml_new_tensor_1d(ctx->ggml_ctx, timestep->type, like->ne[0]));
|
||||
}
|
||||
|
||||
ggml_tensor* unpatchify_audio(GGMLRunnerContext* ctx, ggml_tensor* ax, int64_t audio_length) {
|
||||
if (ax == nullptr) {
|
||||
return nullptr;
|
||||
@ -1392,24 +1367,21 @@ namespace LTXV {
|
||||
if (cfg.cross_attention_adaln) {
|
||||
auto prompt_adaln_single = std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["prompt_adaln_single"]);
|
||||
auto audio_prompt_adaln_single = std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["audio_prompt_adaln_single"]);
|
||||
v_prompt_timestep_mod = prompt_adaln_single->forward(ctx, a_timestep_scaled).first;
|
||||
v_prompt_timestep_mod = prompt_adaln_single->forward(ctx, v_timestep_scaled).first;
|
||||
a_prompt_timestep_mod = audio_prompt_adaln_single->forward(ctx, a_timestep_scaled).first;
|
||||
}
|
||||
|
||||
auto av_ca_video_timestep = repeat_scalar_timestep_like(ctx, effective_audio_timestep, timestep);
|
||||
auto av_ca_audio_timestep = effective_audio_timestep;
|
||||
auto av_ca_factor = cfg.av_ca_timestep_scale_multiplier / cfg.timestep_scale_multiplier;
|
||||
auto av_ca_video_scale_shift_timestep =
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_video_scale_shift_adaln_single"])->forward(ctx, av_ca_video_timestep).first;
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_video_scale_shift_adaln_single"])->forward(ctx, a_timestep_scaled).first;
|
||||
auto av_ca_a2v_gate_noise_timestep =
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_a2v_gate_adaln_single"])
|
||||
->forward(ctx, ggml_ext_scale(ctx->ggml_ctx, av_ca_video_timestep, av_ca_factor))
|
||||
->forward(ctx, ggml_ext_scale(ctx->ggml_ctx, a_timestep_scaled, cfg.av_ca_timestep_scale_multiplier / cfg.timestep_scale_multiplier))
|
||||
.first;
|
||||
auto av_ca_audio_scale_shift_timestep =
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_audio_scale_shift_adaln_single"])->forward(ctx, av_ca_audio_timestep).first;
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_audio_scale_shift_adaln_single"])->forward(ctx, v_timestep_scaled).first;
|
||||
auto av_ca_v2a_gate_noise_timestep =
|
||||
std::dynamic_pointer_cast<AdaLayerNormSingle>(blocks["av_ca_v2a_gate_adaln_single"])
|
||||
->forward(ctx, ggml_ext_scale(ctx->ggml_ctx, av_ca_audio_timestep, av_ca_factor))
|
||||
->forward(ctx, ggml_ext_scale(ctx->ggml_ctx, v_timestep_scaled, cfg.av_ca_timestep_scale_multiplier / cfg.timestep_scale_multiplier))
|
||||
.first;
|
||||
|
||||
for (int i = 0; i < cfg.num_layers; i++) {
|
||||
@ -1438,14 +1410,14 @@ namespace LTXV {
|
||||
|
||||
auto v_shift_scale = get_output_scale_shift(ctx, params["scale_shift_table"], v_embedded_time, cfg.hidden_size);
|
||||
vx = norm_out->forward(ctx, vx);
|
||||
vx = modulate(ctx->ggml_ctx, vx, v_shift_scale[0], v_shift_scale[1]);
|
||||
vx = Flux::modulate(ctx->ggml_ctx, vx, v_shift_scale[0], v_shift_scale[1], true);
|
||||
vx = proj_out->forward(ctx, vx);
|
||||
vx = unpatchify_video(ctx, vx, width, height, frames);
|
||||
|
||||
if (ax != nullptr && audio_time > 0) {
|
||||
auto a_shift_scale = get_output_scale_shift(ctx, params["audio_scale_shift_table"], a_embedded_time, cfg.audio_hidden_size);
|
||||
ax = audio_norm_out->forward(ctx, ax);
|
||||
ax = modulate(ctx->ggml_ctx, ax, a_shift_scale[0], a_shift_scale[1]);
|
||||
ax = Flux::modulate(ctx->ggml_ctx, ax, a_shift_scale[0], a_shift_scale[1], true);
|
||||
ax = audio_proj_out->forward(ctx, ax);
|
||||
ax = unpatchify_audio(ctx, ax, audio_time);
|
||||
}
|
||||
@ -1478,10 +1450,10 @@ namespace LTXV {
|
||||
}
|
||||
|
||||
LTXAVRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string& prefix = "model.diffusion_model")
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
prefix(prefix),
|
||||
params(),
|
||||
model(params) {
|
||||
@ -1892,7 +1864,7 @@ namespace LTXV {
|
||||
|
||||
auto& tensor_storage_map = model_loader.get_tensor_storage_map();
|
||||
std::shared_ptr<LTXAVRunner> ltxav = std::make_shared<LTXAVRunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"model.diffusion_model");
|
||||
|
||||
|
||||
@ -828,10 +828,10 @@ struct MMDiTRunner : public GGMLRunner {
|
||||
MMDiT mmdit;
|
||||
|
||||
MMDiTRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "")
|
||||
: GGMLRunner(backend, params_backend), mmdit(tensor_storage_map) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu), mmdit(tensor_storage_map) {
|
||||
mmdit.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
|
||||
@ -934,7 +934,7 @@ struct MMDiTRunner : public GGMLRunner {
|
||||
// ggml_backend_t backend = ggml_backend_cuda_init(0);
|
||||
ggml_backend_t backend = ggml_backend_cpu_init();
|
||||
ggml_type model_data_type = GGML_TYPE_F16;
|
||||
std::shared_ptr<MMDiTRunner> mmdit = std::make_shared<MMDiTRunner>(backend, backend);
|
||||
std::shared_ptr<MMDiTRunner> mmdit = std::make_shared<MMDiTRunner>(backend, false);
|
||||
{
|
||||
LOG_INFO("loading from '%s'", file_path.c_str());
|
||||
|
||||
|
||||
@ -24,7 +24,7 @@
|
||||
#include "ggml-alloc.h"
|
||||
#include "ggml-backend.h"
|
||||
#include "ggml.h"
|
||||
#include "ggml_extend_backend.h"
|
||||
#include "ggml_extend_backend.hpp"
|
||||
#include "zip.h"
|
||||
|
||||
#include "name_conversion.h"
|
||||
|
||||
@ -411,13 +411,13 @@ public:
|
||||
|
||||
public:
|
||||
PhotoMakerIDEncoder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
SDVersion version = VERSION_SDXL,
|
||||
PMVersion pm_v = PM_VERSION_1,
|
||||
float sty = 20.f)
|
||||
: GGMLRunner(backend, params_backend),
|
||||
: GGMLRunner(backend, offload_params_to_cpu),
|
||||
version(version),
|
||||
pm_version(pm_v),
|
||||
style_strength(sty) {
|
||||
@ -568,11 +568,11 @@ struct PhotoMakerIDEmbed : public GGMLRunner {
|
||||
bool applied = false;
|
||||
|
||||
PhotoMakerIDEmbed(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
ModelLoader* ml,
|
||||
const std::string& file_path = "",
|
||||
const std::string& prefix = "")
|
||||
: file_path(file_path), GGMLRunner(backend, params_backend), model_loader(ml) {
|
||||
: file_path(file_path), GGMLRunner(backend, offload_params_to_cpu), model_loader(ml) {
|
||||
if (!model_loader->init_from_file_and_convert_name(file_path, prefix)) {
|
||||
load_failed = true;
|
||||
}
|
||||
|
||||
@ -488,12 +488,12 @@ namespace Qwen {
|
||||
SDVersion version;
|
||||
|
||||
QwenImageRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
SDVersion version = VERSION_QWEN_IMAGE,
|
||||
bool zero_cond_t = false)
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
qwen_image_params.num_layers = 0;
|
||||
qwen_image_params.zero_cond_t = zero_cond_t;
|
||||
for (auto pair : tensor_storage_map) {
|
||||
@ -686,7 +686,7 @@ namespace Qwen {
|
||||
}
|
||||
|
||||
std::shared_ptr<QwenImageRunner> qwen_image = std::make_shared<QwenImageRunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"model.diffusion_model",
|
||||
VERSION_QWEN_IMAGE);
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
10
src/t5.hpp
10
src/t5.hpp
@ -321,11 +321,11 @@ struct T5Runner : public GGMLRunner {
|
||||
std::vector<int> relative_position_bucket_vec;
|
||||
|
||||
T5Runner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
bool is_umt5 = false)
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
if (is_umt5) {
|
||||
params.vocab_size = 256384;
|
||||
params.relative_attention = false;
|
||||
@ -464,11 +464,11 @@ struct T5Embedder {
|
||||
T5Runner model;
|
||||
|
||||
T5Embedder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
bool is_umt5 = false)
|
||||
: model(backend, params_backend, tensor_storage_map, prefix, is_umt5), tokenizer(is_umt5) {
|
||||
: model(backend, offload_params_to_cpu, tensor_storage_map, prefix, is_umt5), tokenizer(is_umt5) {
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, ggml_tensor*>& tensors, const std::string prefix) {
|
||||
@ -576,7 +576,7 @@ struct T5Embedder {
|
||||
}
|
||||
}
|
||||
|
||||
std::shared_ptr<T5Embedder> t5 = std::make_shared<T5Embedder>(backend, backend, tensor_storage_map, "", true);
|
||||
std::shared_ptr<T5Embedder> t5 = std::make_shared<T5Embedder>(backend, false, tensor_storage_map, "", true);
|
||||
|
||||
t5->alloc_params_buffer();
|
||||
std::map<std::string, ggml_tensor*> tensors;
|
||||
|
||||
@ -541,14 +541,14 @@ struct TinyImageAutoEncoder : public VAE {
|
||||
bool decode_only = false;
|
||||
|
||||
TinyImageAutoEncoder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
bool decoder_only = true,
|
||||
SDVersion version = VERSION_SD1)
|
||||
: decode_only(decoder_only),
|
||||
taesd(decoder_only, version),
|
||||
VAE(version, backend, params_backend) {
|
||||
VAE(version, backend, offload_params_to_cpu) {
|
||||
scale_input = false;
|
||||
taesd.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
@ -603,14 +603,14 @@ struct TinyVideoAutoEncoder : public VAE {
|
||||
bool decode_only = false;
|
||||
|
||||
TinyVideoAutoEncoder(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
bool decoder_only = true,
|
||||
SDVersion version = VERSION_WAN2)
|
||||
: decode_only(decoder_only),
|
||||
taehv(decoder_only, version),
|
||||
VAE(version, backend, params_backend) {
|
||||
VAE(version, backend, offload_params_to_cpu) {
|
||||
scale_input = false;
|
||||
taehv.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
|
||||
@ -603,11 +603,11 @@ struct UNetModelRunner : public GGMLRunner {
|
||||
UnetModelBlock unet;
|
||||
|
||||
UNetModelRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map,
|
||||
const std::string prefix,
|
||||
SDVersion version = VERSION_SD1)
|
||||
: GGMLRunner(backend, params_backend), unet(version, tensor_storage_map) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu), unet(version, tensor_storage_map) {
|
||||
unet.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
|
||||
|
||||
@ -4,18 +4,12 @@
|
||||
#include "stable-diffusion.h"
|
||||
#include "util.h"
|
||||
|
||||
#include <utility>
|
||||
|
||||
UpscalerGGML::UpscalerGGML(int n_threads,
|
||||
bool direct,
|
||||
int tile_size,
|
||||
std::string backend_spec,
|
||||
std::string params_backend_spec)
|
||||
int tile_size)
|
||||
: n_threads(n_threads),
|
||||
direct(direct),
|
||||
tile_size(tile_size),
|
||||
backend_spec(std::move(backend_spec)),
|
||||
params_backend_spec(std::move(params_backend_spec)) {
|
||||
tile_size(tile_size) {
|
||||
}
|
||||
|
||||
void UpscalerGGML::set_max_graph_vram_bytes(size_t max_vram_bytes) {
|
||||
@ -30,51 +24,19 @@ bool UpscalerGGML::load_from_file(const std::string& esrgan_path,
|
||||
int n_threads) {
|
||||
ggml_log_set(ggml_log_callback_default, nullptr);
|
||||
|
||||
std::string error;
|
||||
if (!backend_manager.init(backend_spec.c_str(),
|
||||
params_backend_spec.c_str(),
|
||||
offload_params_to_cpu,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
&error)) {
|
||||
LOG_ERROR("upscaler backend config failed: %s", error.c_str());
|
||||
return false;
|
||||
}
|
||||
auto backend_for = [&](SDBackendModule module) {
|
||||
ggml_backend_t module_backend = backend_manager.runtime_backend(module);
|
||||
if (module_backend == nullptr) {
|
||||
LOG_ERROR("failed to initialize %s backend", sd_backend_module_name(module));
|
||||
}
|
||||
return module_backend;
|
||||
};
|
||||
auto params_backend_for = [&](SDBackendModule module) {
|
||||
ggml_backend_t module_backend = backend_manager.params_backend(module);
|
||||
if (module_backend == nullptr) {
|
||||
LOG_ERROR("failed to initialize %s params backend", sd_backend_module_name(module));
|
||||
}
|
||||
return module_backend;
|
||||
};
|
||||
auto ensure_backend_pair = [&](SDBackendModule module) {
|
||||
if (backend_for(module) == nullptr) {
|
||||
return false;
|
||||
}
|
||||
return params_backend_for(module) != nullptr;
|
||||
};
|
||||
if (!ensure_backend_pair(SDBackendModule::UPSCALER)) {
|
||||
return false;
|
||||
}
|
||||
backend = sd_get_default_backend();
|
||||
|
||||
ModelLoader model_loader;
|
||||
if (!model_loader.init_from_file_and_convert_name(esrgan_path)) {
|
||||
LOG_ERROR("init model loader from file failed: '%s'", esrgan_path.c_str());
|
||||
}
|
||||
model_loader.set_wtype_override(model_data_type);
|
||||
if (!backend) {
|
||||
LOG_DEBUG("Using CPU backend");
|
||||
backend = ggml_backend_cpu_init();
|
||||
}
|
||||
LOG_INFO("Upscaler weight type: %s", ggml_type_name(model_data_type));
|
||||
esrgan_upscaler = std::make_shared<ESRGAN>(backend_for(SDBackendModule::UPSCALER),
|
||||
params_backend_for(SDBackendModule::UPSCALER),
|
||||
tile_size,
|
||||
model_loader.get_tensor_storage_map());
|
||||
esrgan_upscaler = std::make_shared<ESRGAN>(backend, offload_params_to_cpu, tile_size, model_loader.get_tensor_storage_map());
|
||||
esrgan_upscaler->set_max_graph_vram_bytes(max_graph_vram_bytes);
|
||||
if (direct) {
|
||||
esrgan_upscaler->set_conv2d_direct_enabled(true);
|
||||
@ -148,16 +110,14 @@ upscaler_ctx_t* new_upscaler_ctx(const char* esrgan_path_c_str,
|
||||
bool offload_params_to_cpu,
|
||||
bool direct,
|
||||
int n_threads,
|
||||
int tile_size,
|
||||
const char* backend,
|
||||
const char* params_backend) {
|
||||
int tile_size) {
|
||||
upscaler_ctx_t* upscaler_ctx = (upscaler_ctx_t*)malloc(sizeof(upscaler_ctx_t));
|
||||
if (upscaler_ctx == nullptr) {
|
||||
return nullptr;
|
||||
}
|
||||
std::string esrgan_path(esrgan_path_c_str);
|
||||
|
||||
upscaler_ctx->upscaler = new UpscalerGGML(n_threads, direct, tile_size, SAFE_STR(backend), SAFE_STR(params_backend));
|
||||
upscaler_ctx->upscaler = new UpscalerGGML(n_threads, direct, tile_size);
|
||||
if (upscaler_ctx->upscaler == nullptr) {
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
@ -2,7 +2,6 @@
|
||||
#define __SD_UPSCALER_H__
|
||||
|
||||
#include "esrgan.hpp"
|
||||
#include "ggml_extend_backend.h"
|
||||
#include "stable-diffusion.h"
|
||||
#include "tensor.hpp"
|
||||
|
||||
@ -10,7 +9,7 @@
|
||||
#include <string>
|
||||
|
||||
struct UpscalerGGML {
|
||||
SDBackendManager backend_manager;
|
||||
ggml_backend_t backend = nullptr; // general backend
|
||||
ggml_type model_data_type = GGML_TYPE_F16;
|
||||
std::shared_ptr<ESRGAN> esrgan_upscaler;
|
||||
std::string esrgan_path;
|
||||
@ -18,14 +17,10 @@ struct UpscalerGGML {
|
||||
bool direct = false;
|
||||
int tile_size = 128;
|
||||
size_t max_graph_vram_bytes = 0;
|
||||
std::string backend_spec;
|
||||
std::string params_backend_spec;
|
||||
|
||||
UpscalerGGML(int n_threads,
|
||||
bool direct = false,
|
||||
int tile_size = 128,
|
||||
std::string backend_spec = "",
|
||||
std::string params_backend_spec = "");
|
||||
bool direct = false,
|
||||
int tile_size = 128);
|
||||
|
||||
bool load_from_file(const std::string& esrgan_path,
|
||||
bool offload_params_to_cpu,
|
||||
|
||||
72
src/util.cpp
72
src/util.cpp
@ -25,7 +25,7 @@
|
||||
|
||||
#include "ggml-backend.h"
|
||||
#include "ggml.h"
|
||||
#include "ggml_extend_backend.h"
|
||||
#include "ggml_extend_backend.hpp"
|
||||
#include "stable-diffusion.h"
|
||||
|
||||
bool ends_with(const std::string& str, const std::string& ending) {
|
||||
@ -758,6 +758,76 @@ std::vector<std::pair<std::string, float>> parse_prompt_attention(const std::str
|
||||
return res;
|
||||
}
|
||||
|
||||
// test if the backend is a specific one, e.g. "CUDA", "ROCm", "Vulkan" etc.
|
||||
bool sd_backend_is(ggml_backend_t backend, const std::string& name) {
|
||||
if (!backend) {
|
||||
return false;
|
||||
}
|
||||
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
|
||||
if (!dev)
|
||||
return false;
|
||||
std::string dev_name = ggml_backend_dev_name(dev);
|
||||
return dev_name.find(name) != std::string::npos;
|
||||
}
|
||||
|
||||
ggml_backend_t sd_get_default_backend() {
|
||||
ggml_backend_load_all_once();
|
||||
static std::once_flag once;
|
||||
std::call_once(once, []() {
|
||||
size_t dev_count = ggml_backend_dev_count();
|
||||
if (dev_count == 0) {
|
||||
LOG_ERROR("No devices found!");
|
||||
} else {
|
||||
LOG_DEBUG("Found %zu backend devices:", dev_count);
|
||||
for (size_t i = 0; i < dev_count; ++i) {
|
||||
auto dev = ggml_backend_dev_get(i);
|
||||
LOG_DEBUG("#%zu: %s", i, ggml_backend_dev_name(dev));
|
||||
}
|
||||
}
|
||||
});
|
||||
ggml_backend_t backend = nullptr;
|
||||
const char* SD_VK_DEVICE = getenv("SD_VK_DEVICE");
|
||||
if (SD_VK_DEVICE != nullptr) {
|
||||
std::string sd_vk_device_str = SD_VK_DEVICE;
|
||||
try {
|
||||
unsigned long long device = std::stoull(sd_vk_device_str);
|
||||
std::string vk_device_name = "Vulkan" + std::to_string(device);
|
||||
if (backend_name_exists(vk_device_name)) {
|
||||
LOG_INFO("Selecting %s as main device by env var SD_VK_DEVICE", vk_device_name.c_str());
|
||||
backend = init_named_backend(vk_device_name);
|
||||
if (!backend) {
|
||||
LOG_WARN("Device %s requested by SD_VK_DEVICE failed to init. Falling back to the default device.", vk_device_name.c_str());
|
||||
}
|
||||
} else {
|
||||
LOG_WARN("Device %s requested by SD_VK_DEVICE was not found. Falling back to the default device.", vk_device_name.c_str());
|
||||
}
|
||||
} catch (const std::invalid_argument&) {
|
||||
LOG_WARN("SD_VK_DEVICE environment variable is not a valid integer (%s). Falling back to the default device.", SD_VK_DEVICE);
|
||||
} catch (const std::out_of_range&) {
|
||||
LOG_WARN("SD_VK_DEVICE environment variable value is out of range for `unsigned long long` type (%s). Falling back to the default device.", SD_VK_DEVICE);
|
||||
}
|
||||
}
|
||||
|
||||
if (!backend) {
|
||||
std::string dev_name = get_default_backend_name();
|
||||
backend = init_named_backend(dev_name);
|
||||
if (!backend && !dev_name.empty()) {
|
||||
LOG_WARN("device %s failed to init", dev_name.c_str());
|
||||
}
|
||||
}
|
||||
|
||||
if (!backend) {
|
||||
LOG_WARN("loading CPU backend");
|
||||
backend = ggml_backend_cpu_init();
|
||||
}
|
||||
|
||||
if (ggml_backend_is_cpu(backend)) {
|
||||
LOG_DEBUG("Using CPU backend");
|
||||
}
|
||||
|
||||
return backend;
|
||||
}
|
||||
|
||||
// namespace is needed to avoid conflicts with ggml_backend_extend.hpp
|
||||
namespace ggml_cpu {
|
||||
#include "ggml-cpu.h"
|
||||
|
||||
@ -86,6 +86,7 @@ bool sd_should_preview_noisy();
|
||||
|
||||
// test if the backend is a specific one, e.g. "CUDA", "ROCm", "Vulkan" etc.
|
||||
bool sd_backend_is(ggml_backend_t backend, const std::string& name);
|
||||
ggml_backend_t sd_get_default_backend();
|
||||
|
||||
#define LOG_DEBUG(format, ...) log_printf(SD_LOG_DEBUG, __FILE__, __LINE__, format, ##__VA_ARGS__)
|
||||
#define LOG_INFO(format, ...) log_printf(SD_LOG_INFO, __FILE__, __LINE__, format, ##__VA_ARGS__)
|
||||
|
||||
@ -62,8 +62,8 @@ protected:
|
||||
}
|
||||
|
||||
public:
|
||||
VAE(SDVersion version, ggml_backend_t backend, ggml_backend_t params_backend)
|
||||
: version(version), GGMLRunner(backend, params_backend) {}
|
||||
VAE(SDVersion version, ggml_backend_t backend, bool offload_params_to_cpu)
|
||||
: version(version), GGMLRunner(backend, offload_params_to_cpu) {}
|
||||
|
||||
int get_scale_factor() {
|
||||
int scale_factor = 8;
|
||||
@ -219,8 +219,8 @@ public:
|
||||
};
|
||||
|
||||
struct FakeVAE : public VAE {
|
||||
FakeVAE(SDVersion version, ggml_backend_t backend, ggml_backend_t params_backend)
|
||||
: VAE(version, backend, params_backend) {}
|
||||
FakeVAE(SDVersion version, ggml_backend_t backend, bool offload_params_to_cpu)
|
||||
: VAE(version, backend, offload_params_to_cpu) {}
|
||||
|
||||
int get_encoder_output_channels(int input_channels) {
|
||||
return input_channels;
|
||||
|
||||
12
src/wan.hpp
12
src/wan.hpp
@ -1126,12 +1126,12 @@ namespace WAN {
|
||||
WanVAE ae;
|
||||
|
||||
WanVAERunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
bool decode_only = false,
|
||||
SDVersion version = VERSION_WAN2)
|
||||
: decode_only(decode_only), ae(decode_only, version == VERSION_WAN2_2_TI2V), VAE(version, backend, params_backend) {
|
||||
: decode_only(decode_only), ae(decode_only, version == VERSION_WAN2_2_TI2V), VAE(version, backend, offload_params_to_cpu) {
|
||||
ae.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
|
||||
@ -1329,7 +1329,7 @@ namespace WAN {
|
||||
// ggml_backend_t backend = ggml_backend_cuda_init(0);
|
||||
ggml_backend_t backend = ggml_backend_cpu_init();
|
||||
ggml_type model_data_type = GGML_TYPE_F16;
|
||||
std::shared_ptr<WanVAERunner> vae = std::make_shared<WanVAERunner>(backend, backend, String2TensorStorage{}, "", false, VERSION_WAN2_2_TI2V);
|
||||
std::shared_ptr<WanVAERunner> vae = std::make_shared<WanVAERunner>(backend, false, String2TensorStorage{}, "", false, VERSION_WAN2_2_TI2V);
|
||||
{
|
||||
LOG_INFO("loading from '%s'", file_path.c_str());
|
||||
|
||||
@ -2094,11 +2094,11 @@ namespace WAN {
|
||||
SDVersion version;
|
||||
|
||||
WanRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
SDVersion version = VERSION_WAN2)
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
wan_params.num_layers = 0;
|
||||
for (auto pair : tensor_storage_map) {
|
||||
std::string tensor_name = pair.first;
|
||||
@ -2346,7 +2346,7 @@ namespace WAN {
|
||||
}
|
||||
|
||||
std::shared_ptr<WanRunner> wan = std::make_shared<WanRunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"model.diffusion_model",
|
||||
VERSION_WAN2_2_TI2V);
|
||||
|
||||
@ -473,11 +473,11 @@ namespace ZImage {
|
||||
SDVersion version;
|
||||
|
||||
ZImageRunner(ggml_backend_t backend,
|
||||
ggml_backend_t params_backend,
|
||||
bool offload_params_to_cpu,
|
||||
const String2TensorStorage& tensor_storage_map = {},
|
||||
const std::string prefix = "",
|
||||
SDVersion version = VERSION_Z_IMAGE)
|
||||
: GGMLRunner(backend, params_backend) {
|
||||
: GGMLRunner(backend, offload_params_to_cpu) {
|
||||
z_image = ZImageModel(z_image_params);
|
||||
z_image.init(params_ctx, tensor_storage_map, prefix);
|
||||
}
|
||||
@ -620,7 +620,7 @@ namespace ZImage {
|
||||
}
|
||||
|
||||
std::shared_ptr<ZImageRunner> z_image = std::make_shared<ZImageRunner>(backend,
|
||||
backend,
|
||||
false,
|
||||
tensor_storage_map,
|
||||
"model.diffusion_model",
|
||||
VERSION_QWEN_IMAGE);
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user