Compare commits

...

4 Commits

Author SHA1 Message Date
Mario Limonciello
78124b6454
ci: update ROCm releases to 7.14.0 (#1802) 2026-07-24 21:25:04 +08:00
Henry Kroll III
b338b4b4b9
docs: add Gimp plugins to UIs section (#1799) 2026-07-24 21:21:17 +08:00
Wagner Bruna
b8bf67615c
fix: correct dangling pointer to empty image reference vector (#1813) 2026-07-24 21:20:33 +08:00
fszontagh
8d377073e3
feat: add IP-Adapter support for SD 1.5 and SDXL (#1803) 2026-07-24 21:19:44 +08:00
14 changed files with 447 additions and 226 deletions

View File

@ -449,8 +449,8 @@ jobs:
runs-on: windows-2022
env:
ROCM_VERSION: "7.13.0"
GPU_TARGETS: "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
ROCM_VERSION: "7.14.0"
GPU_TARGETS: "gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
steps:
- uses: actions/checkout@v3
@ -472,34 +472,68 @@ jobs:
uses: actions/cache@v4
with:
path: C:\TheRock\build
key: rocm-${{ env.ROCM_VERSION }}-gfx1151-${{ runner.os }}
key: rocm-wheels-${{ env.ROCM_VERSION }}-${{ runner.os }}
- name: ccache
uses: ggml-org/ccache-action@v1.2.16
with:
key: windows-latest-rocm-${{ env.ROCM_VERSION }}-x64
key: windows-rocm-${{ env.ROCM_VERSION }}-x64
evict-old-files: 1d
- name: Install ROCm
- name: Install ROCm with Wheels
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"
write-host "Setting up Python virtual environment"
# Create the venv directly at the cache location to avoid relocation issues
New-Item -Path "C:\TheRock\build" -ItemType Directory -Force | Out-Null
python -m venv C:\TheRock\build\.venv
& C:\TheRock\build\.venv\Scripts\Activate.ps1
write-host "Upgrading pip"
python -m pip install --upgrade pip
write-host "Installing ROCm wheels for multi-arch support"
# Install ROCm wheels for multi-arch support (this may take several minutes)
python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{env.ROCM_VERSION}}"
# Pre-expand the devel tree so it is included in the cache
write-host "Initializing ROCm devel tree"
rocm-sdk init
if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
write-host "Completed ROCm wheel installation to C:\TheRock\build"
- name: Setup ROCm Environment
run: |
$rocmPath = "C:\TheRock\build"
$ErrorActionPreference = "Stop"
# Activate venv from cache or fresh install
& C:\TheRock\build\.venv\Scripts\Activate.ps1
# Expand the devel tree (idempotent; no-op if already done during install)
rocm-sdk init
if ($LASTEXITCODE -ne 0) { throw "rocm-sdk init failed with exit code $LASTEXITCODE" }
# Get ROCm installation paths using the rocm-sdk CLI tool
$rocmPath = (rocm-sdk path --root)
if (-not $rocmPath) { throw "rocm-sdk path --root returned empty - devel package may not be installed" }
$rocmPath = $rocmPath.Trim()
$cmakePath = (rocm-sdk path --cmake).Trim()
$binPath = (rocm-sdk path --bin).Trim()
write-host "ROCm root: $rocmPath"
write-host "CMake path: $cmakePath"
write-host "Bin path: $binPath"
echo "HIP_PATH=$rocmPath" >> $env:GITHUB_ENV
echo "CMAKE_PREFIX_PATH=$cmakePath" >> $env:GITHUB_ENV
echo "HIP_DEVICE_LIB_PATH=$rocmPath\lib\llvm\amdgcn\bitcode" >> $env:GITHUB_ENV
echo "HIP_PLATFORM=amd" >> $env:GITHUB_ENV
echo "LLVM_PATH=$rocmPath\lib\llvm" >> $env:GITHUB_ENV
echo "$rocmPath\bin" >> $env:GITHUB_PATH
echo "$rocmPath\lib\llvm\bin" >> $env:GITHUB_PATH
echo "$binPath" >> $env:GITHUB_PATH
# Keep venv in PATH for subsequent steps
echo "C:\TheRock\build\.venv\Scripts" >> $env:GITHUB_PATH
- name: Build
run: |
@ -527,139 +561,6 @@ jobs:
- name: Pack artifacts
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
run: |
$ErrorActionPreference = "Stop"
$dst = "build\bin"
$rocmBin = Join-Path "${env:HIP_PATH}" "bin"
$requiredRocmPaths = @(
(Join-Path $rocmBin "rocblas.dll"),
(Join-Path $rocmBin "rocblas\library")
)
foreach ($path in $requiredRocmPaths) {
if (!(Test-Path $path)) {
throw "Missing ROCm runtime dependency: $path"
}
}
foreach ($pattern in @("rocblas*.dll", "hipblas*.dll", "libhipblas*.dll")) {
Copy-Item -Path (Join-Path $rocmBin $pattern) -Destination $dst -Force -ErrorAction SilentlyContinue
}
foreach ($dir in @("rocblas", "hipblaslt")) {
$src = Join-Path $rocmBin $dir
if (Test-Path $src) {
Copy-Item -Path $src -Destination $dst -Recurse -Force
}
}
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"
steps:
- uses: actions/checkout@v3
with:
submodules: recursive
- name: Setup Node
uses: actions/setup-node@v4
with:
node-version: 20
- name: Setup pnpm
uses: pnpm/action-setup@v4
with:
version: 10.15.1
- name: Cache ROCm Installation
id: cache-rocm
uses: actions/cache@v4
with:
path: C:\Program Files\AMD\ROCm
key: rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }}
- name: ccache
uses: ggml-org/ccache-action@v1.2.16
with:
key: windows-latest-cmake-hip-${{ env.HIPSDK_INSTALLER_VERSION }}-x64
evict-old-files: 1d
- name: Install ROCm
if: steps.cache-rocm.outputs.cache-hit != 'true'
run: |
$ErrorActionPreference = "Stop"
write-host "Downloading AMD HIP SDK Installer"
Invoke-WebRequest -Uri "https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-${{ env.HIPSDK_INSTALLER_VERSION }}-Win11-For-HIP.exe" -OutFile "${env:RUNNER_TEMP}\rocm-install.exe"
write-host "Installing AMD HIP SDK"
$proc = Start-Process "${env:RUNNER_TEMP}\rocm-install.exe" -ArgumentList '-install' -NoNewWindow -PassThru
$completed = $proc.WaitForExit(600000)
if (-not $completed) {
Write-Error "ROCm installation timed out after 10 minutes. Killing the process"
$proc.Kill()
exit 1
}
if ($proc.ExitCode -ne 0) {
Write-Error "ROCm installation failed with exit code $($proc.ExitCode)"
exit 1
}
write-host "Completed AMD HIP SDK installation"
- name: Verify ROCm
run: |
# Find and test ROCm installation
$clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1
if (-not $clangPath) {
Write-Error "ROCm installation not found"
exit 1
}
& $clangPath.FullName --version
# Set HIP_PATH environment variable for later steps
echo "HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path)" >> $env:GITHUB_ENV
- name: Build
run: |
mkdir build
cd build
$env:CMAKE_PREFIX_PATH="${env:HIP_PATH}"
cmake .. `
-G "Unix Makefiles" `
-DSD_HIPBLAS=ON `
-DSD_BUILD_SHARED_LIBS=ON `
-DGGML_NATIVE=OFF `
-DCMAKE_C_COMPILER=clang `
-DCMAKE_CXX_COMPILER=clang++ `
-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: prompt/actions-commit-hash@v2
- name: Pack artifacts
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
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\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\*
- name: Upload artifacts
@ -679,11 +580,8 @@ jobs:
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.13.0"
gpu_targets: "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
- ROCM_VERSION: "7.14.0"
gpu_targets: "gfx900;gfx906;gfx908;gfx90a;gfx90c;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
build: x64
steps:
@ -702,7 +600,7 @@ jobs:
- name: Dependencies
id: depends
run: |
sudo apt install -y build-essential cmake wget zip ninja-build
sudo apt install -y build-essential git cmake wget
- name: Free disk space
run: |
@ -723,38 +621,36 @@ jobs:
sudo apt clean
df -h
- 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: Setup TheRock
if: matrix.ROCM_VERSION != '7.2.1'
- name: Setup TheRock with Wheels
id: therock_env
run: |
wget https://repo.amd.com/rocm/tarball/therock-dist-linux-gfx1151-${{ matrix.ROCM_VERSION }}.tar.gz
mkdir install
tar -xf *.tar.gz -C install
export ROCM_PATH=$(pwd)/install
echo ROCM_PATH=$ROCM_PATH >> $GITHUB_ENV
echo PATH=$PATH:$ROCM_PATH/bin >> $GITHUB_ENV
echo LD_LIBRARY_PATH=$ROCM_PATH/lib:$ROCM_PATH/llvm/lib:$ROCM_PATH/lib/rocprofiler-systems >> $GITHUB_ENV
# Create Python virtual environment
python3 -m venv .venv
source .venv/bin/activate
# Install ROCm wheels for build
# libraries = HIP runtime and CMake configs needed for linking
# devel = compilers, headers, static libs
python -m pip install --upgrade pip
python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{matrix.ROCM_VERSION}}"
# Get ROCm installation paths using the rocm-sdk CLI tool
ROCM_PATH=$(rocm-sdk path --root)
CMAKE_PATH=$(rocm-sdk path --cmake)
BIN_PATH=$(rocm-sdk path --bin)
echo "ROCM_PATH=$ROCM_PATH"
echo "CMAKE_PATH=$CMAKE_PATH"
echo "BIN_PATH=$BIN_PATH"
# Set environment variables
echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
# Keep venv activated for subsequent steps
echo "$(pwd)/.venv/bin" >> $GITHUB_PATH
# setup-node installs into /opt/hostedtoolcache, which is removed above.
# Keep Node/pnpm setup after disk cleanup so the server frontend can be embedded.
@ -839,7 +735,6 @@ jobs:
- build-and-push-docker-images
- macOS-latest-cmake
- windows-latest-cmake
- windows-latest-cmake-hip
- windows-latest-rocm
steps:

View File

@ -70,6 +70,7 @@ API and command-line option may change frequently.***
- [HunyuanVideo 1.5](./docs/hunyuan_video.md)
- [LingBot-Video](./docs/lingbot_video.md)
- [PhotoMaker](./docs/photo_maker.md) support.
- [IP-Adapter](./docs/ip_adapter.md) support (SD 1.5 and SDXL)
- Control Net support with SD 1.5
- [ADetailer](./docs/adetailer.md)
- LoRA support, same as [stable-diffusion-webui](https://github.com/AUTOMATIC1111/stable-diffusion-webui/wiki/Features#lora)
@ -167,6 +168,7 @@ These projects wrap `stable-diffusion.cpp` for easier use in other languages/fra
These projects use `stable-diffusion.cpp` as a backend for their image generation.
- [GIMP Plugins](https://github.com/themanyone/gimp-plugins)
- [Jellybox](https://jellybox.com)
- [Stable Diffusion GUI](https://github.com/fszontagh/sd.cpp.gui.wx)
- [Stable Diffusion CLI-GUI](https://github.com/piallai/stable-diffusion.cpp)

55
docs/ip_adapter.md Normal file
View File

@ -0,0 +1,55 @@
# IP-Adapter
stable-diffusion.cpp supports [IP-Adapter](https://github.com/tencent-ailab/IP-Adapter)
image-prompt conditioning for SD 1.5 and SDXL. Given a reference image,
IP-Adapter transfers the subject and appearance of that image into the
generation, alongside the text prompt.
IP-Adapter encodes the reference image with a CLIP-Vision (ViT-H/14)
encoder, projects the embedding into a few image tokens, and injects them
through a decoupled cross-attention added to every attn2 layer of the
UNet. It composes with Control Net, so a reference image (appearance) and
an OpenPose hint (pose) can be combined in a single generation.
## Required weights
1. A base SD 1.5 or SDXL model.
2. A CLIP-Vision (ViT-H/14) image encoder, passed with `--clip_vision`
(for example `clip_vision_h.safetensors`).
3. An IP-Adapter weight file, passed with `--ip-adapter`. The `vit-h`
variants reuse the same ViT-H encoder as above. From
[h94/IP-Adapter](https://huggingface.co/h94/IP-Adapter):
- SD 1.5: `models/ip-adapter_sd15.safetensors`
- SDXL: `sdxl_models/ip-adapter_sdxl_vit-h.safetensors`
## Options
- `--ip-adapter <path>` path to the IP-Adapter weight file.
- `--ip-adapter-image <path>` path to the reference image.
- `--ip-adapter-strength <float>` strength of the IP-Adapter injection
(default 1.0). Lower values let the text prompt dominate; 0.6 to 0.8 is
a good starting range.
## Example (SD 1.5)
```
sd-cli -m ..\models\sd_v1.5.safetensors --clip_vision ..\models\clip_vision_h.safetensors --ip-adapter ..\models\ip-adapter_sd15.safetensors --ip-adapter-image ..\assets\reference.png --ip-adapter-strength 0.8 -p "a woman, best quality" -n "lowres, bad anatomy" --cfg-scale 7 --steps 30 --sampling-method dpm++2m --scheduler karras -W 512 -H 512
```
## Example (SDXL)
```
sd-cli -m ..\models\sdxl.safetensors --clip_vision ..\models\clip_vision_h.safetensors --ip-adapter ..\models\ip-adapter_sdxl_vit-h.safetensors --ip-adapter-image ..\assets\reference.png --ip-adapter-strength 0.8 -p "a woman, best quality" -n "lowres, bad anatomy" --cfg-scale 6 --steps 25 --sampling-method dpm++2m --scheduler karras -W 1024 -H 1024 --diffusion-fa --vae-tiling
```
The SDXL VAE decode at 1024x1024 is memory heavy; add `--vae-tiling` (and
`--offload-to-cpu`) on GPUs with limited VRAM.
## Combining with Control Net
Add the usual Control Net options to keep the reference appearance while
controlling the pose:
```
sd-cli -m ..\models\sdxl.safetensors --clip_vision ..\models\clip_vision_h.safetensors --ip-adapter ..\models\ip-adapter_sdxl_vit-h.safetensors --ip-adapter-image ..\assets\character.png --ip-adapter-strength 0.9 --control-net ..\models\OpenPoseXL2.safetensors --control-image ..\assets\pose.png --control-strength 0.8 -p "a character, side view" --cfg-scale 6 --steps 25 -W 1024 -H 1024 --diffusion-fa --vae-tiling
```

View File

@ -817,6 +817,16 @@ int main(int argc, const char* argv[]) {
}
}
if (gen_params.ip_adapter_image_path.size() > 0) {
if (!load_sd_image_from_file(gen_params.ip_adapter_image.put(),
gen_params.ip_adapter_image_path.c_str(),
0,
0)) {
LOG_ERROR("load image from '%s' failed", gen_params.ip_adapter_image_path.c_str());
return 1;
}
}
if (!gen_params.control_video_path.empty()) {
gen_params.control_frames.clear();
if (!load_images_from_dir(gen_params.control_video_path,

View File

@ -427,6 +427,11 @@ ArgOptions SDContextParams::get_options() {
"path to control net model",
0,
&control_net_path},
{"",
"--ip-adapter",
"path to IP-Adapter model (requires --clip_vision)",
0,
&ip_adapter_path},
{"",
"--motion-module",
"path to AnimateDiff motion module (SD 1.5); enables video generation on --video-frames > 1",
@ -876,6 +881,7 @@ sd_ctx_params_t SDContextParams::to_sd_ctx_params_t(bool taesd_preview) {
sd_ctx_params.audio_vae_path = audio_vae_path.c_str();
sd_ctx_params.taesd_path = taesd_path.c_str();
sd_ctx_params.control_net_path = control_net_path.c_str();
sd_ctx_params.ip_adapter_path = ip_adapter_path.c_str();
sd_ctx_params.motion_module_path = motion_module_path.c_str();
sd_ctx_params.embeddings = embedding_vec.data();
sd_ctx_params.embedding_count = static_cast<uint32_t>(embedding_vec.size());
@ -966,6 +972,11 @@ ArgOptions SDGenerationParams::get_options() {
"path to control image, control net",
0,
&control_image_path},
{"",
"--ip-adapter-image",
"path to the IP-Adapter reference image",
0,
&ip_adapter_image_path},
{"",
"--control-video",
"path to control video frames, It must be a directory path. The video frames inside should be stored as images in "
@ -1158,6 +1169,10 @@ ArgOptions SDGenerationParams::get_options() {
"--control-strength",
"strength to apply Control Net (default: 0.9). 1.0 corresponds to full destruction of information in init image",
&control_strength},
{"",
"--ip-adapter-strength",
"strength to apply IP-Adapter (default: 1.0)",
&ip_adapter_strength},
{"",
"--moe-boundary",
"timestep boundary for Wan2.2 MoE model. (default: 0.875). Only enabled if `--high-noise-steps` is set to -1",
@ -2498,6 +2513,8 @@ sd_img_gen_params_t SDGenerationParams::to_sd_img_gen_params_t() {
params.qwen_image_layers = qwen_image_layers;
params.control_image = control_image.get();
params.control_strength = control_strength;
params.ip_adapter_image = ip_adapter_image.get();
params.ip_adapter_strength = ip_adapter_strength;
params.pm_params = pm_params;
params.pulid_params = pulid_params;
params.vae_tiling_params = vae_tiling_params;

View File

@ -133,6 +133,7 @@ struct SDContextParams {
std::string taesd_path;
std::string esrgan_path;
std::string control_net_path;
std::string ip_adapter_path;
std::string motion_module_path;
std::string embedding_dir;
std::string photo_maker_path;
@ -200,6 +201,7 @@ struct SDGenerationParams {
int64_t seed = 42;
float strength = 0.75f;
float control_strength = 0.9f;
float ip_adapter_strength = 1.0f;
bool auto_resize_ref_image = true;
bool increase_ref_index = false;
bool embed_image_metadata = true;
@ -208,6 +210,7 @@ struct SDGenerationParams {
std::string end_image_path;
std::string mask_image_path;
std::string control_image_path;
std::string ip_adapter_image_path;
std::vector<std::string> ref_image_paths;
std::string control_video_path;
@ -274,6 +277,7 @@ struct SDGenerationParams {
std::vector<SDImageOwner> ref_images;
SDImageOwner mask_image;
SDImageOwner control_image;
SDImageOwner ip_adapter_image;
std::vector<SDImageOwner> pm_id_images;
std::vector<SDImageOwner> control_frames;

View File

@ -200,6 +200,7 @@ typedef struct {
const char* audio_vae_path;
const char* taesd_path;
const char* control_net_path;
const char* ip_adapter_path;
const char* motion_module_path;
const sd_embedding_t* embeddings;
uint32_t embedding_count;
@ -375,6 +376,8 @@ typedef struct {
int batch_count;
sd_image_t control_image;
float control_strength;
sd_image_t ip_adapter_image;
float ip_adapter_strength;
sd_pm_params_t pm_params;
sd_pulid_params_t pulid_params;
sd_tiling_params_t vae_tiling_params;

View File

@ -1689,6 +1689,8 @@ struct GGMLRunnerContext {
bool conv2d_direct_enabled = false;
bool circular_x_enabled = false;
bool circular_y_enabled = false;
ggml_tensor* ip_context = nullptr;
float ip_scale = 1.0f;
std::shared_ptr<WeightAdapter> weight_adapter = nullptr;
std::vector<std::pair<ggml_tensor*, std::string>>* debug_tensors = nullptr;
std::function<ggml_tensor*(const std::string&)> get_cache_tensor;

View File

@ -0,0 +1,88 @@
#ifndef __SD_MODEL_ADAPTER_IP_ADAPTER_HPP__
#define __SD_MODEL_ADAPTER_IP_ADAPTER_HPP__
#include "core/ggml_extend.hpp"
#include "model/common/block.hpp"
#include "model_loader.h"
namespace IPAdapter {
struct ImageProjModel : public GGMLBlock {
int64_t num_tokens = 4;
int64_t ctx_dim = 768;
int64_t clip_dim = 1024;
ImageProjModel() {}
ImageProjModel(int64_t num_tokens, int64_t ctx_dim, int64_t clip_dim)
: num_tokens(num_tokens), ctx_dim(ctx_dim), clip_dim(clip_dim) {
blocks["proj"] = std::shared_ptr<GGMLBlock>(new Linear(clip_dim, num_tokens * ctx_dim, true));
blocks["norm"] = std::shared_ptr<GGMLBlock>(new LayerNorm(ctx_dim));
}
ggml_tensor* forward(GGMLRunnerContext* ctx, ggml_tensor* image_embeds) {
auto proj = std::dynamic_pointer_cast<Linear>(blocks["proj"]);
auto norm = std::dynamic_pointer_cast<LayerNorm>(blocks["norm"]);
int64_t n = image_embeds->ne[1];
auto x = proj->forward(ctx, image_embeds);
x = ggml_reshape_3d(ctx->ggml_ctx, x, ctx_dim, num_tokens, n);
x = norm->forward(ctx, x);
return x;
}
};
struct IPAdapterRunner : public GGMLRunner {
ImageProjModel image_proj;
int64_t num_tokens = 4;
std::string prefix;
IPAdapterRunner(ggml_backend_t backend,
const String2TensorStorage& tensor_storage_map,
const std::string prefix,
std::shared_ptr<RunnerWeightManager> weight_manager = nullptr)
: GGMLRunner(backend, weight_manager), prefix(prefix) {
int64_t ctx_dim = 768;
int64_t clip_dim = 1024;
int64_t out_dim = 3072;
auto norm_iter = tensor_storage_map.find(prefix + ".image_proj.norm.weight");
if (norm_iter != tensor_storage_map.end()) {
ctx_dim = norm_iter->second.ne[0];
}
auto proj_iter = tensor_storage_map.find(prefix + ".image_proj.proj.weight");
if (proj_iter != tensor_storage_map.end()) {
clip_dim = proj_iter->second.ne[0];
out_dim = proj_iter->second.ne[1];
}
num_tokens = out_dim / ctx_dim;
image_proj = ImageProjModel(num_tokens, ctx_dim, clip_dim);
image_proj.init(params_ctx, tensor_storage_map, prefix + ".image_proj");
}
std::string get_desc() override {
return "ip_adapter";
}
void get_param_tensors(std::map<std::string, ggml_tensor*>& tensors, const std::string = "") {
image_proj.get_param_tensors(tensors, prefix + ".image_proj");
}
ggml_cgraph* build_graph(const sd::Tensor<float>& image_embeds_tensor) {
ggml_cgraph* gf = new_graph_custom(1024);
ggml_tensor* embeds = make_input(image_embeds_tensor);
auto runner_ctx = get_context();
ggml_tensor* out = image_proj.forward(&runner_ctx, embeds);
ggml_build_forward_expand(gf, out);
return gf;
}
sd::Tensor<float> compute(int n_threads, const sd::Tensor<float>& image_embeds) {
auto get_graph = [&]() -> ggml_cgraph* {
return build_graph(image_embeds);
};
return take_or_empty(GGMLRunner::compute<float>(get_graph, n_threads, true, true, true));
}
};
} // namespace IPAdapter
#endif // __SD_MODEL_ADAPTER_IP_ADAPTER_HPP__

View File

@ -311,16 +311,32 @@ protected:
int64_t n_head;
int64_t d_head;
bool xtra_dim = false;
bool enable_ip = false;
bool has_ip = false;
void init_params(ggml_context* ctx, const String2TensorStorage& tensor_storage_map = {}, const std::string prefix = "") override {
GGMLBlock::init_params(ctx, tensor_storage_map, prefix);
if (enable_ip &&
tensor_storage_map.find(prefix + "to_k_ip.weight") != tensor_storage_map.end()) {
has_ip = true;
int64_t inner_dim = d_head * n_head;
int64_t ip_dim = tensor_storage_map.at(prefix + "to_k_ip.weight").ne[0];
blocks["to_k_ip"] = std::shared_ptr<GGMLBlock>(new Linear(ip_dim, inner_dim, false));
blocks["to_v_ip"] = std::shared_ptr<GGMLBlock>(new Linear(ip_dim, inner_dim, false));
}
}
public:
CrossAttention(int64_t query_dim,
int64_t context_dim,
int64_t n_head,
int64_t d_head)
int64_t d_head,
bool enable_ip = false)
: n_head(n_head),
d_head(d_head),
query_dim(query_dim),
context_dim(context_dim) {
context_dim(context_dim),
enable_ip(enable_ip) {
int64_t inner_dim = d_head * n_head;
if (context_dim == 320 && d_head == 320) {
// LOG_DEBUG("CrossAttention: temp set dim to 1024 for sdxs_09");
@ -363,6 +379,15 @@ public:
}
x = ggml_ext_attention_ext(ctx->ggml_ctx, ctx->backend, q, k, v, n_head, nullptr, false, ctx->flash_attn_enabled); // [N, n_token, inner_dim]
if (has_ip && ctx->ip_context != nullptr && ctx->ip_scale != 0.0f) {
auto to_k_ip = std::dynamic_pointer_cast<Linear>(blocks["to_k_ip"]);
auto to_v_ip = std::dynamic_pointer_cast<Linear>(blocks["to_v_ip"]);
auto k_ip = to_k_ip->forward(ctx, ctx->ip_context);
auto v_ip = to_v_ip->forward(ctx, ctx->ip_context);
auto x_ip = ggml_ext_attention_ext(ctx->ggml_ctx, ctx->backend, q, k_ip, v_ip, n_head, nullptr, false, ctx->flash_attn_enabled);
x = ggml_add(ctx->ggml_ctx, x, ggml_scale(ctx->ggml_ctx, x_ip, ctx->ip_scale));
}
x = to_out_0->forward(ctx, x); // [N, n_token, query_dim]
return x;
}
@ -387,7 +412,7 @@ public:
// inner_dim is always None or equal to dim
// gated_ff is always True
blocks["attn1"] = std::shared_ptr<GGMLBlock>(new CrossAttention(dim, dim, n_head, d_head));
blocks["attn2"] = std::shared_ptr<GGMLBlock>(new CrossAttention(dim, context_dim, n_head, d_head));
blocks["attn2"] = std::shared_ptr<GGMLBlock>(new CrossAttention(dim, context_dim, n_head, d_head, true));
blocks["ff"] = std::shared_ptr<GGMLBlock>(new FeedForward(dim, dim));
blocks["norm1"] = std::shared_ptr<GGMLBlock>(new LayerNorm(dim));
blocks["norm2"] = std::shared_ptr<GGMLBlock>(new LayerNorm(dim));

View File

@ -46,6 +46,8 @@ struct UNetDiffusionExtra {
int num_video_frames = -1;
const std::vector<sd::Tensor<float>>* controls = nullptr;
float control_strength = 0.f;
const sd::Tensor<float>* ip_context = nullptr;
float ip_scale = 1.f;
};
struct SkipLayerDiffusionExtra {

View File

@ -775,7 +775,9 @@ struct UNetModelRunner : public DiffusionModelRunner {
const sd::Tensor<float>& y_tensor = {},
int num_video_frames = -1,
const std::vector<sd::Tensor<float>>& controls_tensor = {},
float control_strength = 0.f) {
float control_strength = 0.f,
const sd::Tensor<float>& ip_context_tensor = {},
float ip_scale = 1.f) {
ggml_cgraph* gf = new_graph_custom(UNET_GRAPH_SIZE);
ggml_tensor* x = make_input(x_tensor);
@ -783,6 +785,7 @@ struct UNetModelRunner : public DiffusionModelRunner {
ggml_tensor* context = make_optional_input(context_tensor);
ggml_tensor* c_concat = make_optional_input(c_concat_tensor);
ggml_tensor* y = make_optional_input(y_tensor);
ggml_tensor* ip_context = make_optional_input(ip_context_tensor);
std::vector<ggml_tensor*> controls;
controls.reserve(controls_tensor.size());
for (const auto& control_tensor : controls_tensor) {
@ -794,6 +797,8 @@ struct UNetModelRunner : public DiffusionModelRunner {
}
auto runner_ctx = get_context();
runner_ctx.ip_context = ip_context;
runner_ctx.ip_scale = ip_scale;
ggml_tensor* out = unet.forward(&runner_ctx,
x,
@ -818,14 +823,16 @@ struct UNetModelRunner : public DiffusionModelRunner {
const sd::Tensor<float>& y = {},
int num_video_frames = -1,
const std::vector<sd::Tensor<float>>& controls = {},
float control_strength = 0.f) {
float control_strength = 0.f,
const sd::Tensor<float>& ip_context = {},
float ip_scale = 1.f) {
// x: [N, in_channels, h, w]
// timesteps: [N, ]
// context: [N, max_position, hidden_size]([N, 77, 768]) or [1, max_position, hidden_size]
// c_concat: [N, in_channels, h, w] or [1, in_channels, h, w]
// y: [N, adm_in_channels] or [1, adm_in_channels]
auto get_graph = [&]() -> ggml_cgraph* {
return build_graph(x, timesteps, context, c_concat, y, num_video_frames, controls, control_strength);
return build_graph(x, timesteps, context, c_concat, y, num_video_frames, controls, control_strength, ip_context, ip_scale);
};
return restore_trailing_singleton_dims(GGMLRunner::compute<float>(get_graph, n_threads, false, false, false), x.dim());
@ -845,7 +852,9 @@ struct UNetModelRunner : public DiffusionModelRunner {
tensor_or_empty(diffusion_params.y),
extra->num_video_frames,
extra->controls ? *extra->controls : empty_controls,
extra->control_strength);
extra->control_strength,
extra->ip_context ? *extra->ip_context : sd::Tensor<float>{},
extra->ip_scale);
}
void test() {

View File

@ -1285,11 +1285,54 @@ static std::string convert_esrgan_tensor_name(std::string name) {
return name;
}
static const std::map<int, std::string>& ip_adapter_index_map(SDVersion version) {
static const std::map<int, std::string> sd15_map = {
{1, "input_blocks.1.1.transformer_blocks.0"}, {3, "input_blocks.2.1.transformer_blocks.0"}, {5, "input_blocks.4.1.transformer_blocks.0"}, {7, "input_blocks.5.1.transformer_blocks.0"}, {9, "input_blocks.7.1.transformer_blocks.0"}, {11, "input_blocks.8.1.transformer_blocks.0"}, {13, "output_blocks.3.1.transformer_blocks.0"}, {15, "output_blocks.4.1.transformer_blocks.0"}, {17, "output_blocks.5.1.transformer_blocks.0"}, {19, "output_blocks.6.1.transformer_blocks.0"}, {21, "output_blocks.7.1.transformer_blocks.0"}, {23, "output_blocks.8.1.transformer_blocks.0"}, {25, "output_blocks.9.1.transformer_blocks.0"}, {27, "output_blocks.10.1.transformer_blocks.0"}, {29, "output_blocks.11.1.transformer_blocks.0"}, {31, "middle_block.1.transformer_blocks.0"}};
static std::map<int, std::string> sdxl_map;
if (sdxl_map.empty()) {
std::vector<std::pair<std::string, int>> order = {
{"input_blocks.4.1", 2}, {"input_blocks.5.1", 2}, {"input_blocks.7.1", 10}, {"input_blocks.8.1", 10}, {"output_blocks.0.1", 10}, {"output_blocks.1.1", 10}, {"output_blocks.2.1", 10}, {"output_blocks.3.1", 2}, {"output_blocks.4.1", 2}, {"output_blocks.5.1", 2}, {"middle_block.1", 10}};
int idx = 1;
for (const auto& [block, depth] : order) {
for (int m = 0; m < depth; m++) {
sdxl_map[idx] = block + ".transformer_blocks." + std::to_string(m);
idx += 2;
}
}
}
return sd_version_is_sdxl(version) ? sdxl_map : sd15_map;
}
static std::string convert_ip_adapter_name(std::string name, SDVersion version) {
if (starts_with(name, "image_proj.")) {
return "ip_adapter." + name;
}
if (starts_with(name, "ip_adapter.")) {
auto items = split_string(name, '.');
if (items.size() < 4) {
return name;
}
int idx = atoi(items[1].c_str());
const auto& mp = ip_adapter_index_map(version);
auto blk = mp.find(idx);
if (blk == mp.end()) {
return name;
}
return "model.diffusion_model." + blk->second + ".attn2." + items[2] + "." + items[3];
}
return name;
}
std::string convert_tensor_name(std::string name, SDVersion version) {
if (version == VERSION_ESRGAN) {
return convert_esrgan_tensor_name(std::move(name));
}
if (starts_with(name, "ip_adapter.") || starts_with(name, "image_proj.")) {
return convert_ip_adapter_name(std::move(name), version);
}
bool is_lora = false;
bool is_lycoris_underline = false;
bool is_underline = false;

View File

@ -22,6 +22,7 @@
#include "conditioning/conditioner.hpp"
#include "core/backend_fit.h"
#include "extensions/generation_extension.h"
#include "model/adapter/ip_adapter.hpp"
#include "model/adapter/lora.hpp"
#include "model/diffusion/anima.hpp"
#include "model/diffusion/animatediff.hpp"
@ -221,6 +222,9 @@ public:
std::shared_ptr<VAE> preview_vae;
std::shared_ptr<LTXV::LTXAudioVAERunner> audio_vae_model;
std::shared_ptr<ControlNet> control_net;
std::shared_ptr<IPAdapter::IPAdapterRunner> ip_adapter;
sd::Tensor<float> ip_adapter_tokens;
float ip_adapter_strength = 1.0f;
std::vector<std::shared_ptr<GenerationExtension>> generation_extensions;
std::vector<std::shared_ptr<LoraModel>> runtime_lora_models;
bool apply_lora_immediately = false;
@ -841,6 +845,13 @@ public:
}
}
if (strlen(SAFE_STR(sd_ctx_params->ip_adapter_path)) > 0) {
if (!model_loader.init_from_file(sd_ctx_params->ip_adapter_path)) {
LOG_ERROR("init ip-adapter model loader from file failed: '%s'", sd_ctx_params->ip_adapter_path);
return false;
}
}
model_loader.convert_tensors_name();
version = model_loader.get_sd_version();
@ -1294,6 +1305,33 @@ public:
}
}
if (strlen(SAFE_STR(sd_ctx_params->ip_adapter_path)) > 0 && clip_vision == nullptr) {
if (!ensure_backend_pair(SDBackendModule::CLIP_VISION)) {
return false;
}
clip_vision = std::make_shared<FrozenCLIPVisionEmbedder>(backend_for(SDBackendModule::CLIP_VISION),
tensor_storage_map,
model_manager);
clip_vision->set_max_graph_vram_bytes(max_graph_vram_bytes_for_module(SDBackendModule::CLIP_VISION));
if (!register_runner_params("CLIP vision",
clip_vision,
SDBackendModule::CLIP_VISION)) {
return false;
}
}
if (strlen(SAFE_STR(sd_ctx_params->ip_adapter_path)) > 0) {
ip_adapter = std::make_shared<IPAdapter::IPAdapterRunner>(backend_for(SDBackendModule::DIFFUSION),
tensor_storage_map,
"ip_adapter",
model_manager);
if (!register_runner_params("IP-Adapter",
ip_adapter,
SDBackendModule::DIFFUSION)) {
return false;
}
}
if (!ensure_backend_pair(SDBackendModule::VAE)) {
return false;
}
@ -2061,6 +2099,24 @@ public:
return output;
}
void compute_ip_adapter_tokens(const sd_image_t& image, float strength) {
ip_adapter_tokens = {};
ip_adapter_strength = strength;
if (ip_adapter == nullptr || clip_vision == nullptr || image.data == nullptr) {
return;
}
auto image_tensor = sd_image_to_tensor(image);
auto embed = get_clip_vision_output(image_tensor, true, -1);
if (embed.empty()) {
return;
}
ip_adapter_tokens = ip_adapter->compute(n_threads, embed);
if (!ip_adapter_tokens.empty()) {
LOG_INFO("IP-Adapter: %lld image tokens, strength %.2f",
(long long)ip_adapter_tokens.shape()[1], strength);
}
}
std::vector<float> process_timesteps(const std::vector<float>& timesteps,
const sd::Tensor<float>& init_latent,
const sd::Tensor<float>& denoise_mask,
@ -2549,7 +2605,8 @@ public:
auto run_condition = [&](const SDCondition& condition,
const sd::Tensor<float>* c_concat_override = nullptr,
const std::vector<int>* local_skip_layers = nullptr,
const std::vector<sd::Tensor<float>>* ref_latents_override = nullptr) -> sd::Tensor<float> {
const std::vector<sd::Tensor<float>>* ref_latents_override = nullptr,
bool apply_ip = true) -> sd::Tensor<float> {
diffusion_params.context = condition.c_crossattn.empty() ? nullptr : &condition.c_crossattn;
diffusion_params.c_concat = c_concat_override != nullptr ? c_concat_override : (condition.c_concat.empty() ? nullptr : &condition.c_concat);
diffusion_params.y = condition.c_vector.empty() ? nullptr : &condition.c_vector;
@ -2560,7 +2617,12 @@ public:
if (animatediff_loaded && noised_input.dim() >= 4 && noised_input.shape()[3] > 1) {
nvf = static_cast<int>(noised_input.shape()[3]);
}
diffusion_params.extra = UNetDiffusionExtra{nvf, &controls, control_strength};
UNetDiffusionExtra unet_extra{nvf, &controls, control_strength};
if (apply_ip && !ip_adapter_tokens.empty()) {
unet_extra.ip_context = &ip_adapter_tokens;
unet_extra.ip_scale = ip_adapter_strength;
}
diffusion_params.extra = unet_extra;
} else if (sd_version_is_sd3(version)) {
diffusion_params.extra = SkipLayerDiffusionExtra{local_skip_layers};
} else if (sd_version_is_flux(version) || sd_version_is_flux2(version) || sd_version_is_longcat(version) || sd_version_is_sefi_image(version)) {
@ -2651,7 +2713,9 @@ public:
}
uncond_out = run_condition(uncond,
uncond.c_concat.empty() ? nullptr : &uncond.c_concat,
uncond_skip_layers);
uncond_skip_layers,
nullptr,
false);
if (uncond_out.empty()) {
return {};
}
@ -2660,7 +2724,8 @@ public:
img_uncond_out = run_condition(img_uncond,
img_uncond.c_concat.empty() ? nullptr : &img_uncond.c_concat,
nullptr,
uncond_without_ref_latents ? &empty_ref_latents : nullptr);
uncond_without_ref_latents ? &empty_ref_latents : nullptr,
false);
if (img_uncond_out.empty()) {
return {};
}
@ -4954,6 +5019,7 @@ static std::optional<ImageGenerationEmbeds> prepare_image_generation_embeds(sd_c
request->pulid_params,
condition_params,
plan->total_steps);
sd_ctx->sd->compute_ip_adapter_tokens(sd_img_gen_params->ip_adapter_image, sd_img_gen_params->ip_adapter_strength);
int64_t prepare_start_ms = ggml_time_ms();
condition_params.zero_out_masked = false;
auto cond = sd_ctx->sd->cond_stage_model->get_learned_condition(sd_ctx->sd->n_threads,
@ -5005,8 +5071,8 @@ static std::optional<ImageGenerationEmbeds> prepare_image_generation_embeds(sd_c
}
condition_params.text = request->negative_prompt;
condition_params.zero_out_masked = zero_out_masked;
if (use_ref_latent_img_cfg) {
std::vector<sd::Tensor<float>> empty_ref_images;
if (use_ref_latent_img_cfg) {
condition_params.ref_images = &empty_ref_images;
}
img_uncond = sd_ctx->sd->cond_stage_model->get_learned_condition(sd_ctx->sd->n_threads,