feat: add MiniMax-H3-Image support (#2324 )

perf(cuda): fuse RMSNorm and RoPE in Sol-H3-VAE (#2323 )
fix: avoid structured binding capture in MiniMax H3 SOL attention
2026-11-10 19:06:33 +00:00 · 2026-11-11 01:24:25 +08:00 · 2026-11-11 01:06:53 +08:00 · 2026-11-11 01:02:14 +08:00 · 2026-11-10 23:56:53 +08:00 · 2026-11-10 23:45:01 +08:00
675 changed files with 205838 additions and 26835 deletions
@@ -1,7 +1,8 @@
 build*/
 docs/
 test/
+tests/

 .cache/
 *.swp
-models/
+models/
@@ -0,0 +1,2 @@
+/start-server.sh text eol=lf
+/start-server.bat text eol=crlf
@@ -6,7 +6,9 @@ body:
  - type: markdown
    attributes:
      value: |
-        Please use this template and include as many details as possible to help us reproduce and fix the issue.
+        Before submitting a bug report, please read the [Troubleshooting guide](https://github.com/leejet/stable-diffusion.cpp/blob/master/docs/troubleshooting.md) and try the steps relevant to your problem.
+
+        If the problem persists, complete this form and include what you tried and the results, along with enough details to help us reproduce and fix the issue.
  - type: textarea
    id: commit
    attributes:
@@ -0,0 +1,4 @@
+contact_links:
+  - name: Troubleshooting
+    url: https://github.com/leejet/stable-diffusion.cpp/blob/master/docs/troubleshooting.md
+    about: Read the troubleshooting guide first. If the problem persists, submit a bug report.
@@ -15,7 +15,8 @@ on:
      [
        ".github/workflows/**",
        ".dockerignore",
-        "Dockerfile*",
+        "docker/Dockerfile*",
+        "scripts/ci-cuda-cache.sh",
        "**/CMakeLists.txt",
        "**/Makefile",
        "**/*.h",
@@ -23,8 +24,13 @@ on:
        "**/*.c",
        "**/*.cpp",
        "**/*.cu",
+        "**/*.cuh",
+        "ggml",
        "examples/server/frontend",
        "examples/server/frontend/**",
+        "examples/server/recipes/*.json",
+        "start-server.sh",
+        "start-server.bat",
      ]
  pull_request:
    types: [opened, synchronize, reopened]
@@ -32,7 +38,8 @@ on:
      [
        ".github/workflows/**",
        ".dockerignore",
-        "Dockerfile*",
+        "docker/Dockerfile*",
+        "scripts/ci-cuda-cache.sh",
        "**/CMakeLists.txt",
        "**/Makefile",
        "**/*.h",
@@ -40,13 +47,21 @@ on:
        "**/*.c",
        "**/*.cpp",
        "**/*.cu",
+        "**/*.cuh",
+        "ggml",
        "examples/server/frontend",
        "examples/server/frontend/**",
+        "examples/server/recipes/*.json",
+        "start-server.sh",
+        "start-server.bat",
      ]

 env:
  BRANCH_NAME: ${{ github.head_ref || github.ref_name }}

+permissions:
+  contents: read
+
 concurrency:
  group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
  cancel-in-progress: true
@@ -58,17 +73,17 @@ jobs:
    steps:
      - name: Clone
        id: checkout
-        uses: actions/checkout@v3
+        uses: actions/checkout@v6
        with:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

@@ -83,13 +98,14 @@ jobs:
        run: |
          mkdir build
          cd build
-          cmake .. -DGGML_AVX2=ON -DSD_BUILD_SHARED_LIBS=ON
+          cmake .. -DSD_BUILD_SHARED_LIBS=ON -DGGML_NATIVE=OFF -DSD_BUILD_SHARED_GGML_LIB=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON -DCMAKE_INSTALL_RPATH='$ORIGIN'
          cmake --build . --config Release

      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Fetch system info
        id: system-info
@@ -105,11 +121,12 @@ jobs:
        run: |
          cp ggml/LICENSE ./build/bin/ggml.txt
          cp LICENSE ./build/bin/stable-diffusion.cpp.txt
-          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip ./build/bin/*
+          chmod +x start-server.sh
+          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip ./build/bin/* ./start-server.sh ./start-server.bat

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
          path: |
@@ -121,17 +138,17 @@ jobs:
    steps:
      - name: Clone
        id: checkout
-        uses: actions/checkout@v3
+        uses: actions/checkout@v6
        with:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

@@ -146,13 +163,14 @@ jobs:
        run: |
          mkdir build
          cd build
-          cmake .. -DSD_BUILD_SHARED_LIBS=ON -DSD_VULKAN=ON
+          cmake .. -DSD_BUILD_SHARED_LIBS=ON -DSD_VULKAN=ON -DGGML_NATIVE=OFF -DSD_BUILD_SHARED_GGML_LIB=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON -DCMAKE_INSTALL_RPATH='$ORIGIN'
          cmake --build . --config Release

      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Fetch system info
        id: system-info
@@ -168,11 +186,12 @@ jobs:
        run: |
          cp ggml/LICENSE ./build/bin/ggml.txt
          cp LICENSE ./build/bin/stable-diffusion.cpp.txt
-          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-vulkan.zip ./build/bin/*
+          chmod +x start-server.sh
+          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-vulkan.zip ./build/bin/* ./start-server.sh ./start-server.bat

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}-vulkan.zip
          path: |
@@ -207,6 +226,7 @@ jobs:
              UBUNTU_VERSION=24.04
              CUDA_ARCHITECTURES=121
              GGML_CUDA_FA_ALL_QUANTS=ON
+              GGML_CUDA_ENABLE_DYNAMIC_CPU_BACKENDS=OFF

    env:
      REGISTRY: ghcr.io
@@ -219,25 +239,27 @@ jobs:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Set up Docker Buildx
-        uses: docker/setup-buildx-action@v3
+        id: buildx
+        uses: docker/setup-buildx-action@v4

      - name: Log in to the container registry
-        uses: docker/login-action@v3
+        uses: docker/login-action@v4
        with:
          registry: ${{ env.REGISTRY }}
          username: ${{ github.actor }}
@@ -245,7 +267,7 @@ jobs:

      - name: Extract metadata for Docker
        id: meta
-        uses: docker/metadata-action@v5
+        uses: docker/metadata-action@v6
        with:
          images: ${{ env.REGISTRY }}/${{ env.IMAGE_NAME }}

@@ -256,36 +278,57 @@ jobs:
          # if set to "true" but frees about 6 GB
          tool-cache: false

+      - name: Cache CUDA compilation
+        if: matrix.variant == 'cuda'
+        uses: actions/cache@v5
+        with:
+          path: ${{ runner.temp }}/cuda-ccache
+          key: cuda-ccache-${{ runner.os }}-${{ runner.arch }}-${{ hashFiles('docker/Dockerfile.cuda') }}-${{ github.run_id }}-${{ github.run_attempt }}
+          restore-keys: |
+            cuda-ccache-${{ runner.os }}-${{ runner.arch }}-${{ hashFiles('docker/Dockerfile.cuda') }}-
+
+      - name: Restore CUDA compiler cache mount
+        if: matrix.variant == 'cuda'
+        env:
+          CUDA_CACHE_BUILDER: ${{ steps.buildx.outputs.name }}
+        run: bash scripts/ci-cuda-cache.sh restore
+
      - name: Build and push Docker image
        id: build-push
-        uses: docker/build-push-action@v6
+        uses: docker/build-push-action@v7
        with:
          context: .
          platforms: ${{ matrix.platform }}
          push: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-          file: Dockerfile.${{ matrix.variant }}
+          file: docker/Dockerfile.${{ matrix.variant }}
          tags: ${{ env.REGISTRY }}/${{ env.IMAGE_NAME }}:${{ env.BRANCH_NAME }}-${{ matrix.variant }}${{ matrix.tag-suffix }}
          labels: ${{ steps.meta.outputs.labels }}
          annotations: ${{ steps.meta.outputs.annotations }}
          build-args: ${{ matrix.build-args }}

+      - name: Export CUDA compiler cache mount
+        if: matrix.variant == 'cuda'
+        env:
+          CUDA_CACHE_BUILDER: ${{ steps.buildx.outputs.name }}
+        run: bash scripts/ci-cuda-cache.sh save
+
  macOS-latest-cmake:
    runs-on: macos-latest

    steps:
      - name: Clone
        id: checkout
-        uses: actions/checkout@v3
+        uses: actions/checkout@v6
        with:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

@@ -306,7 +349,8 @@ jobs:
      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Fetch system info
        id: system-info
@@ -322,11 +366,12 @@ jobs:
        run: |
          cp ggml/LICENSE ./build/bin/ggml.txt
          cp LICENSE ./build/bin/stable-diffusion.cpp.txt
-          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip ./build/bin/*
+          chmod +x start-server.sh
+          zip -j sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip ./build/bin/* ./start-server.sh ./start-server.bat

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
          path: |
@@ -341,44 +386,47 @@ jobs:
    strategy:
      matrix:
        include:
-          - build: "noavx"
-            defines: "-DGGML_NATIVE=OFF -DGGML_AVX=OFF -DGGML_AVX2=OFF -DGGML_FMA=OFF -DSD_BUILD_SHARED_LIBS=ON"
-          - build: "avx2"
-            defines: "-DGGML_NATIVE=OFF -DGGML_AVX2=ON -DSD_BUILD_SHARED_LIBS=ON"
-          - build: "avx"
-            defines: "-DGGML_NATIVE=OFF -DGGML_AVX=ON -DGGML_AVX2=OFF -DSD_BUILD_SHARED_LIBS=ON"
-          - build: "avx512"
-            defines: "-DGGML_NATIVE=OFF -DGGML_AVX512=ON -DGGML_AVX=ON -DGGML_AVX2=ON -DSD_BUILD_SHARED_LIBS=ON"
+          - build: "cpu"
+            defines: "-DGGML_NATIVE=OFF -DSD_BUILD_SHARED_LIBS=ON -DSD_BUILD_SHARED_GGML_LIB=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON"
          - build: "cuda12"
-            defines: "-DSD_CUDA=ON -DSD_BUILD_SHARED_LIBS=ON -DCMAKE_CUDA_ARCHITECTURES='61;70;75;80;86;89;90;100;120' -DCMAKE_CUDA_FLAGS='-Xcudafe \"--diag_suppress=177\" -Xcudafe \"--diag_suppress=550\"'"
+            defines: "-DSD_CUDA=ON -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DCMAKE_CUDA_COMPILER_LAUNCHER=ccache -DSD_BUILD_SHARED_LIBS=ON -DCMAKE_CUDA_ARCHITECTURES='61;70;75;80;86;89;90;100;120' -DCMAKE_CUDA_FLAGS='-Xcudafe \"--diag_suppress=177\" -Xcudafe \"--diag_suppress=550\"' -DGGML_NATIVE=OFF -DSD_BUILD_SHARED_GGML_LIB=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON"
          - build: "vulkan"
-            defines: "-DSD_VULKAN=ON -DSD_BUILD_SHARED_LIBS=ON"
+            defines: "-DSD_VULKAN=ON -DSD_BUILD_SHARED_LIBS=ON -DGGML_NATIVE=OFF -DSD_BUILD_SHARED_GGML_LIB=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON"
    steps:
      - name: Clone
        id: checkout
-        uses: actions/checkout@v3
+        uses: actions/checkout@v6
        with:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

      - name: Install cuda-toolkit
        id: cuda-toolkit
        if: ${{ matrix.build == 'cuda12' }}
-        uses: Jimver/cuda-toolkit@v0.2.22
+        uses: Jimver/cuda-toolkit@v0.2.35
        with:
          cuda: "12.8.1"
          method: "network"
          sub-packages: '["nvcc", "cudart", "cublas", "cublas_dev", "thrust", "visual_studio_integration"]'

+      - name: Cache CUDA compilation
+        if: matrix.build == 'cuda12'
+        uses: ggml-org/ccache-action@v1.2.21
+        with:
+          variant: ccache
+          install: binary
+          key: windows-cuda-12.8.1-x64
+          max-size: "2G"
+
      - name: Install Vulkan SDK
        id: get_vulkan
        if: ${{ matrix.build == 'vulkan' }}
@@ -388,35 +436,25 @@ jobs:
          Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}"
          Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin"

-      - name: Activate MSVC environment
-        id: msvc_dev_cmd
-        uses: ilammy/msvc-dev-cmd@v1
-
      - name: Build
        id: cmake_build
+        shell: pwsh
        run: |
+          $vswhere = "${env:ProgramFiles(x86)}\Microsoft Visual Studio\Installer\vswhere.exe"
+          $vsPath = & $vswhere -latest -products * -requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 -property installationPath
+          if (-not $vsPath) { throw "Visual Studio C++ tools not found" }
+          & "$vsPath\Common7\Tools\Launch-VsDevShell.ps1" -Arch amd64 -HostArch amd64 -SkipAutomaticLocation
+
          mkdir build
          cd build
          cmake .. -DCMAKE_CXX_FLAGS='/bigobj' -G Ninja -DCMAKE_C_COMPILER=cl.exe -DCMAKE_CXX_COMPILER=cl.exe -DCMAKE_BUILD_TYPE=Release ${{ matrix.defines }}
          cmake --build .

-      - name: Check AVX512F support
-        id: check_avx512f
-        if: ${{ matrix.build == 'avx512' }}
-        continue-on-error: true
-        run: |
-          cd build
-          $vcdir = $(vswhere -latest -products * -requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 -property installationPath)
-          $msvc = $(join-path $vcdir $('VC\Tools\MSVC\'+$(gc -raw $(join-path $vcdir 'VC\Auxiliary\Build\Microsoft.VCToolsVersion.default.txt')).Trim()))
-          $cl =  $(join-path $msvc 'bin\Hostx64\x64\cl.exe')
-          echo 'int main(void){unsigned int a[4];__cpuid(a,7);return !(a[1]&65536);}' >> avx512f.c
-          & $cl /O2 /GS- /kernel avx512f.c /link /nodefaultlib /entry:main
-          .\avx512f.exe && echo "AVX512F: YES" && ( echo HAS_AVX512F=1 >> $env:GITHUB_ENV ) || echo "AVX512F: NO"
-
      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Pack artifacts
        id: pack_artifacts
@@ -436,7 +474,10 @@ jobs:
            ls .\build\bin
            throw "Can't find stable-diffusion.dll"
          }
-          7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip $filePath
+          if ("${{ matrix.build }}" -eq "cuda12") {
+            Copy-Item thirdparty/cutlass/LICENSE.txt (Join-Path (Split-Path -Parent $filePath) 'cutlass.txt')
+          }
+          7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip $filePath .\start-server.sh .\start-server.bat

      - name: Copy and pack Cuda runtime
        id: pack_cuda_runtime
@@ -449,7 +490,7 @@ jobs:

      - name: Upload Cuda runtime
        if: ${{ matrix.build == 'cuda12' && (github.event_name == 'push' && github.ref == 'refs/heads/master' || github.event.inputs.create_release == 'true') }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-cudart-sd-bin-win-cu12-x64.zip
          path: |
@@ -457,7 +498,7 @@ jobs:

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip
          path: |
@@ -467,57 +508,91 @@ 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: "10.1.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
+      - uses: actions/checkout@v6
        with:
          submodules: recursive

      - name: Setup Node
-        uses: actions/setup-node@v4
+        uses: actions/setup-node@v5
        with:
-          node-version: 20
+          node-version: 24

      - name: Setup pnpm
-        uses: pnpm/action-setup@v4
+        uses: pnpm/action-setup@v5
        with:
          version: 10.15.1

      - name: Cache ROCm Installation
        id: cache-rocm
-        uses: actions/cache@v4
+        uses: actions/cache@v5
        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
+        uses: ggml-org/ccache-action@v1.2.21
        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://stable.repo.amd.com/rocm/whl-next/ "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: |
@@ -540,125 +615,29 @@ jobs:
      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

-      - name: Pack artifacts
+      - name: Bundle HIP runtime libraries
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        run: |
-          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
+          $binPath = (rocm-sdk path --bin).Trim()
+          if (-not $binPath) { throw "rocm-sdk path --bin returned empty" }
+          foreach ($dll in @("amdhip64_7.dll", "rocm_kpack.dll", "amd_comgr.dll")) {
+            $f = Get-ChildItem -Path $binPath -Filter $dll -ErrorAction SilentlyContinue
+            if (-not $f) { throw "no match for $dll in $binPath" }
+            Copy-Item $f.FullName -Destination .\build\bin\ -Force
          }
-          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\*
+          7z a sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip .\build\bin\* .\start-server.sh .\start-server.bat

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        with:
          name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-rocm-${{ env.ROCM_VERSION }}-x64.zip
          path: |
@@ -673,12 +652,9 @@ jobs:
    strategy:
      matrix:
        include:
-          - ROCM_VERSION: "7.2.1"
-            gpu_targets: "gfx908;gfx90a;gfx942;gfx1030;gfx1031;gfx1032;gfx1100;gfx1101;gfx1102;gfx1151;gfx1150;gfx1200;gfx1201"
+          - ROCM_VERSION: "10.1.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'
-          - 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"
-            build: x64

    steps:
      - name: Clone
@@ -687,72 +663,88 @@ jobs:
        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: ccache
-        uses: ggml-org/ccache-action@v1.2.16
+        uses: ggml-org/ccache-action@v1.2.21
        with:
          key: ubuntu-rocm-cmake-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
          evict-old-files: 1d
+          max-size: "2G"
+
+      - name: Tune ccache for reinstalled ROCm toolchain
+        run: |
+          # ROCm is pip-installed fresh each run, so the clang binary's mtime
+          # changes every time. With the default compiler_check=mtime that
+          # invalidates the cache; hash compiler contents instead so warm
+          # builds hit.
+          ccache --set-config=compiler_check=content
+          ccache --set-config=sloppiness=time_macros,include_file_mtime,include_file_ctime

      - 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
+          sudo apt install -y build-essential git cmake wget

      - name: Free disk space
        run: |
-          # Remove preinstalled SDKs and caches not needed for this job
+          df -h
+
+          # Remove preinstalled SDKs and caches not needed for this job before
+          # installing ROCm. The legacy ROCm apt packages are large enough to
+          # exhaust ubuntu-latest if cleanup runs after installation.
          sudo rm -rf /usr/share/dotnet || true
          sudo rm -rf /usr/local/lib/android || true
          sudo rm -rf /opt/ghc || true
          sudo rm -rf /usr/local/.ghcup || true
          sudo rm -rf /opt/hostedtoolcache || true
+          sudo rm -rf /usr/share/swift || true
+          sudo rm -rf /usr/local/share/boost || true
+          docker system prune -af || true

-          # Remove old package lists and caches
-          sudo rm -rf /var/lib/apt/lists/* || true
          sudo apt clean
+          df -h

-      - 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://stable.repo.amd.com/rocm/whl-next/ "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.
+      - name: Setup Node
+        uses: actions/setup-node@v5
+        with:
+          node-version: 24
+
+      - name: Setup pnpm
+        uses: pnpm/action-setup@v5
+        with:
+          version: 10.15.1

      - name: Build
        id: cmake_build
@@ -768,13 +760,19 @@ jobs:
            -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
            -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
            -DCMAKE_POSITION_INDEPENDENT_CODE=ON \
+            -DCMAKE_INSTALL_RPATH='$ORIGIN' \
+            -DGGML_NATIVE=OFF \
+            -DSD_BUILD_SHARED_GGML_LIB=ON \
+            -DGGML_BACKEND_DL=ON \
+            -DGGML_CPU_ALL_VARIANTS=ON \
            -DSD_BUILD_SHARED_LIBS=ON
          cmake --build . --config Release

      - 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
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Prepare artifacts
        id: prepare_artifacts
@@ -798,11 +796,12 @@ 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
+          chmod +x start-server.sh
+          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 ./start-server.sh ./start-server.bat

      - name: Upload artifacts
        if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v6
        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
          path: |
@@ -813,6 +812,9 @@ jobs:

    runs-on: ubuntu-latest

+    permissions:
+      contents: write
+
    needs:
      - ubuntu-latest-cmake
      - ubuntu-latest-cmake-vulkan
@@ -820,18 +822,17 @@ jobs:
      - build-and-push-docker-images
      - macOS-latest-cmake
      - windows-latest-cmake
-      - windows-latest-cmake-hip
      - windows-latest-rocm

    steps:
      - name: Clone
-        uses: actions/checkout@v3
+        uses: actions/checkout@v6
        with:
          fetch-depth: 0

      - name: Download artifacts
        id: download-artifact
-        uses: actions/download-artifact@v4
+        uses: actions/download-artifact@v7
        with:
          path: ./artifact
          pattern: sd-*
@@ -844,36 +845,47 @@ jobs:

      - name: Get commit hash
        id: commit
-        uses: prompt/actions-commit-hash@v2
+        shell: bash
+        run: echo "short=${GITHUB_SHA:0:7}" >> "$GITHUB_OUTPUT"

      - name: Create release
        id: create_release
        if: ${{ github.event_name == 'workflow_dispatch' || github.ref_name == 'master' }}
-        uses: anzz1/action-create-release@v1
+        uses: actions/github-script@v9
        env:
-          GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+          RELEASE_TAG: ${{ format('{0}-{1}-{2}', env.BRANCH_NAME, steps.commit_count.outputs.count, steps.commit.outputs.short) }}
        with:
-          tag_name: ${{ format('{0}-{1}-{2}', env.BRANCH_NAME, steps.commit_count.outputs.count, steps.commit.outputs.short) }}
+          script: |
+            const { data: release } = await github.rest.repos.createRelease({
+              ...context.repo,
+              tag_name: process.env.RELEASE_TAG,
+              name: process.env.RELEASE_TAG,
+              target_commitish: context.sha,
+              draft: false,
+              prerelease: false,
+            });
+            core.setOutput('id', release.id);

      - name: Upload release
        id: upload_release
        if: ${{ github.event_name == 'workflow_dispatch' || github.ref_name == 'master' }}
-        uses: actions/github-script@v3
+        uses: actions/github-script@v9
+        env:
+          RELEASE_ID: ${{ steps.create_release.outputs.id }}
        with:
-          github-token: ${{secrets.GITHUB_TOKEN}}
          script: |
            const path = require('path');
-            const fs = require('fs');
-            const release_id = '${{ steps.create_release.outputs.id }}';
-            for (let file of await fs.readdirSync('./artifact')) {
+            const fs = require('fs').promises;
+            const release_id = Number(process.env.RELEASE_ID);
+            for (const file of await fs.readdir('./artifact')) {
              if (path.extname(file) === '.zip') {
                console.log('uploadReleaseAsset', file);
-                await github.repos.uploadReleaseAsset({
+                await github.rest.repos.uploadReleaseAsset({
                  owner: context.repo.owner,
                  repo: context.repo.repo,
                  release_id: release_id,
                  name: file,
-                  data: await fs.readFileSync(`./artifact/${file}`)
+                  data: await fs.readFile(`./artifact/${file}`)
                });
              }
            }
@@ -0,0 +1,48 @@
+name: Close inactive issues
+
+on:
+  schedule:
+    # Run daily. GitHub cron schedules use UTC.
+    - cron: "30 1 * * *"
+  workflow_dispatch:
+    inputs:
+      debug_only:
+        description: "Dry run: log intended actions without changing issues"
+        required: false
+        default: false
+        type: boolean
+
+permissions:
+  issues: write
+
+concurrency:
+  group: ${{ github.workflow }}
+  cancel-in-progress: false
+
+jobs:
+  close-inactive-issues:
+    runs-on: ubuntu-latest
+    steps:
+      - name: Comment and close inactive issues
+        uses: actions/stale@v10
+        with:
+          days-before-issue-stale: 365
+          days-before-issue-close: 0
+
+          days-before-pr-stale: -1
+          days-before-pr-close: -1
+
+          stale-issue-label: issue:inactive
+          close-issue-label: issue:auto-closed
+          close-issue-reason: completed
+          stale-issue-message: ""
+          close-issue-message: >
+            This issue has had no activity for one year. The latest version of
+            the code may already have fixed the problem.
+
+            If the issue still exists in the latest version, you can reopen
+            this issue at any time with updated reproduction details.
+
+          remove-issue-stale-when-updated: true
+          operations-per-run: 1000
+          debug-only: ${{ github.event_name == 'workflow_dispatch' && inputs.debug_only || false }}
@@ -0,0 +1,61 @@
+name: Close PRs from organization forks
+
+on:
+  pull_request_target:
+    types: [opened, reopened]
+
+permissions:
+  pull-requests: write
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number }}
+  cancel-in-progress: false
+
+jobs:
+  close-organization-fork-pr:
+    if: >-
+      github.event.pull_request.head.repo.owner.type == 'Organization' &&
+      github.event.pull_request.head.repo.id != github.event.pull_request.base.repo.id
+    runs-on: ubuntu-latest
+    timeout-minutes: 5
+    steps:
+      - name: Explain the contribution policy and close the PR
+        uses: actions/github-script@v9
+        with:
+          script: |
+            const { data: pr } = await github.rest.pulls.get({
+              ...context.repo,
+              pull_number: context.issue.number,
+            });
+            const headRepo = pr.head.repo;
+            if (pr.state !== 'open' || !headRepo ||
+                headRepo.id === pr.base.repo.id || headRepo.owner.type !== 'Organization') {
+              return;
+            }
+
+            const marker = '<!-- organization-fork-policy -->';
+            const comments = await github.paginate(github.rest.issues.listComments, {
+              ...context.repo,
+              issue_number: pr.number,
+              per_page: 100,
+            });
+            const alreadyExplained = comments.some(comment =>
+              comment.user?.login === 'github-actions[bot]' && comment.body?.includes(marker));
+            if (!alreadyExplained) {
+              await github.rest.issues.createComment({
+                ...context.repo,
+                issue_number: pr.number,
+                body: [
+                  marker,
+                  'This repository requires contributions from forks to use a personal fork with **Allow edits from maintainers** enabled.',
+                  'GitHub does not support this option for organization-owned forks, so this PR is being closed automatically.',
+                  'Please open a new PR from a fork in your personal GitHub account and enable **Allow edits from maintainers** so maintainers can help update the branch.',
+                  'See [the GitHub documentation](https://docs.github.com/en/pull-requests/how-tos/work-with-forks/allowing-changes-to-a-pull-request-branch-created-from-a-fork).',
+                ].join('\n\n'),
+              });
+            }
+            await github.rest.pulls.update({
+              ...context.repo,
+              pull_number: pr.number,
+              state: 'closed',
+            });
@@ -1,11 +1,17 @@
 build*/
 cmake-build-*/
 test/
+tests/
+logs/
 .vscode/
 .idea/
 .cache/
 *.swp
 *.bat
+!/start-server.bat
+/start-server.local.sh
+/start-server.local.bat
+/start-server.local.ps1
 *.bin
 *.exe
 *.gguf
@@ -13,3 +19,7 @@ output*.png
 models*
 *.log
 preview.png
+.claude/
+CLAUDE.local.md
+.agents/
+.codex/
@@ -10,3 +10,6 @@
 [submodule "thirdparty/libwebm"]
 	path = thirdparty/libwebm
 	url = https://github.com/webmproject/libwebm.git
+[submodule "thirdparty/cutlass"]
+	path = thirdparty/cutlass
+	url = https://github.com/NVIDIA/cutlass.git
@@ -0,0 +1,186 @@
+# Instructions for stable-diffusion.cpp
+
+This document is for AI coding agents working in this repository. It should
+describe agent-specific workflow, repository routing, editing boundaries, and
+project-specific pitfalls.
+
+For general contribution rules, including PR scope, commit conventions, code
+style, dependency updates, security hygiene, and AI-assisted contribution policy,
+see `CONTRIBUTING.md`.
+
+---
+
+## Agent Operating Rules
+
+Before analyzing or modifying the repository:
+
+1. Read this file.
+2. Use `rg` / `rg --files` or directory listing commands to confirm the current
+   tree before relying on a path.
+3. Start from `src/` and relevant `docs/` for runtime behavior.
+4. Read the relevant code before editing.
+5. Prefer the smallest change that fits the existing architecture.
+6. Report focused verification and mention any tests not run.
+
+Agents must not:
+
+* Run `git push`, create PRs, or submit issue/PR comments on the user's behalf.
+* Create commits unless the user explicitly requests that specific commit.
+* Modify `ggml/`, `thirdparty/`, or `examples/server/frontend/` unless
+  explicitly requested and necessary.
+* Read large local model files or tokenizer vocabulary files.
+* Rewrite unrelated code for style-only reasons.
+* Add secrets, model weights, generated binaries, local absolute paths, or
+  machine-specific output.
+
+When a change is large, architectural, or likely to affect public behavior,
+pause and present a short plan before editing.
+
+---
+
+## Repository Map and Editing Boundaries
+
+This is a routing map for agents, not a full architecture document. The layout
+can change, so verify paths before using them. Do not inspect excluded
+large-data directories while checking the tree.
+
+### Primary Project Code
+
+Core implementation lives under `src/`.
+
+Current source layout includes:
+
+* `src/core/` - shared tensor, ggml integration, backend, graph, RNG, and utility
+  code.
+* `src/model/` - model families and model components.
+* `src/model_io/` - model file loading, GGUF, safetensors, pickle, and related
+  serialization helpers.
+* `src/runtime/` - sampling, denoising, guidance, caching, preprocessing, and
+  runtime execution helpers.
+* `src/tokenizers/` - tokenizer implementations.
+* `src/conditioning/` - conditioning and prompt-related implementation.
+* `src/extensions/` - optional feature extensions.
+* top-level `src/*.cpp` and `src/*.h` files - public implementation entry
+  points, model loading, conversion, versioning, and shared managers.
+
+`src/tokenizers/vocab/` contains large tokenizer vocabulary data. Do not read or
+parse files in this directory; reference the path only when necessary.
+
+### Public API
+
+`include/` contains the C API exposed by the project. Currently the primary
+public header is `include/stable-diffusion.h`.
+
+Treat public headers as stable API. Avoid breaking compatibility unless the user
+explicitly requests it. If public behavior changes, update relevant examples or
+documentation.
+
+### Examples
+
+`examples/` contains programs demonstrating library usage.
+
+* `examples/cli/` - command line program for running models, testing features,
+  and debugging.
+* `examples/common/` - shared example support code.
+* `examples/server/` - server application built on top of the library.
+  Before modifying built-in recipes, read and follow
+  [the built-in recipe policy](examples/server/recipes/README.md).
+* `examples/server/frontend/` - git submodule containing independent frontend
+  code. Avoid modifying it unless explicitly requested.
+
+### Documentation and Tooling
+
+* `docs/` - documentation for supported models, build options, behavior, and
+  workflows.
+* `scripts/` - development, model processing, build automation, formatting, and
+  tooling scripts.
+* `cmake/` - CMake support modules.
+* `docker/` - Docker-related project files.
+* `assets/` - documentation assets; not runtime code.
+
+### External, Local, and Generated State
+
+* `ggml/` - git submodule for the ggml dependency.
+* `thirdparty/` - vendored third-party dependencies.
+* `models/` - local model storage. Ignore this directory and do not read model
+  files.
+* `test/` - local testing scripts. Use only when relevant to the task.
+* `build/`, `build_*`, and similar directories - generated build outputs.
+  Inspect them only when debugging a build result.
+
+---
+
+## Agent Workflow for Code Changes
+
+1. Identify the relevant modules under `src/`.
+2. Check whether the change touches the public API in `include/`.
+3. Consult relevant `docs/` and examples before changing user-facing behavior.
+4. Follow existing local patterns before adding new abstractions.
+5. Keep edits scoped to the requested behavior.
+6. Run the narrowest useful build, test, or inspection command available.
+
+Follow `CONTRIBUTING.md` for formatting, naming, PR expectations, dependency
+update policy, and security rules. For tokenizer additions, follow its embedded-data
+allowlist and default to an external `tokenizer.json`.
+
+---
+
+## Code Comments
+
+Keep comments rare and useful.
+
+Do not add comments that only describe what the code does. Add comments only
+when the code cannot fully express the logic, the logic is unusually complex, or
+there are historical reasons, invariants, constraints, compatibility concerns,
+or known pitfalls that future maintainers need to understand.
+
+Do not add task-specific comments that will be meaningless after review.
+
+Examples from the current codebase:
+
+```cpp
+// GOOD: explains a safety constraint that is not obvious from the assignment.
+// From src/model_io/pickle_io.cpp.
+// Non-tensor checkpoint metadata can use REDUCE for arbitrary
+// Python objects. Do not execute it; keep stack shape only.
+stack.push_back(make_none_value());
+
+// BAD: describes only what the next line does.
+// Set the token count to zero.
+token_count = 0;
+```
+
+---
+
+## Text File Encoding
+
+When reading or editing repository text files:
+
+* Prefer UTF-8 with LF for Markdown, frontend source, JSON, and other text-first
+  project files unless the file already clearly uses a different encoding.
+* Do not assume terminal output encoding matches file encoding on Windows.
+* A file that looks garbled in PowerShell output may still be valid UTF-8.
+* When inspecting UTF-8 files in PowerShell, prefer explicit UTF-8 reads such as:
+  * `[Console]::OutputEncoding = [System.Text.Encoding]::UTF8`
+  * `Get-Content -Encoding utf8 <path>`
+* Avoid rewriting a file purely because console output looked garbled; verify
+  the actual file encoding first.
+
+---
+
+## Tensor and Layout Notes
+
+Additional tensor/layout rules for this codebase:
+
+* `sd::Tensor` shape order is not PyTorch/NumPy-style. `shape()[0]` is the
+  lowest and most contiguous dimension, and higher indices are higher
+  dimensions.
+* Broadcasting for `sd::Tensor` must align dimensions from low to high dimension
+  indices. If one tensor has fewer dimensions, append implicit `1`s at the
+  higher-dimension end.
+* `ggml_n_dims` / `ggml_n_dims(tensor)` can drop trailing singleton high
+  dimensions. Do not assume a logical trailing dimension of `1` will still be
+  counted in ggml metadata.
+* Internal tensor-returning interfaces use an empty `sd::Tensor` to represent
+  null, absent, or failure states. Do not add `std::optional<sd::Tensor<...>>`
+  for internal APIs unless a distinct semantic state is truly required.
@@ -0,0 +1,11 @@
+@AGENTS.md
+
+## Claude Code
+
+Follow `AGENTS.md` as the shared repository instructions.
+
+Do not duplicate contribution, style, PR, dependency, or security policy here;
+use `CONTRIBUTING.md` as the canonical source for those rules.
+
+Keep Claude-specific project notes in this file only when they do not apply to
+other coding agents.
@@ -9,12 +9,28 @@ if (NOT XCODE AND NOT MSVC AND NOT CMAKE_BUILD_TYPE)
 endif()

 if (MSVC)
+    cmake_host_system_information(RESULT SD_LOGICAL_CORES QUERY NUMBER_OF_LOGICAL_CORES)
+    if (SD_LOGICAL_CORES GREATER 8)
+        math(EXPR SD_DEFAULT_MSVC_MP_COUNT "${SD_LOGICAL_CORES} - 2")
+    elseif (SD_LOGICAL_CORES GREATER 2)
+        math(EXPR SD_DEFAULT_MSVC_MP_COUNT "${SD_LOGICAL_CORES} - 1")
+    else()
+        set(SD_DEFAULT_MSVC_MP_COUNT "${SD_LOGICAL_CORES}")
+    endif()
+    if (NOT SD_DEFAULT_MSVC_MP_COUNT GREATER 0)
+        set(SD_DEFAULT_MSVC_MP_COUNT 1)
+    endif()
+    set(SD_MSVC_MP_COUNT "${SD_DEFAULT_MSVC_MP_COUNT}" CACHE STRING "Maximum parallel MSVC compilations per project")
+    if (NOT SD_MSVC_MP_COUNT MATCHES "^[1-9][0-9]*$" OR SD_MSVC_MP_COUNT GREATER 65536)
+        message(FATAL_ERROR "SD_MSVC_MP_COUNT must be an integer between 1 and 65536")
+    endif()
    add_compile_definitions(_CRT_SECURE_NO_WARNINGS)
    add_compile_definitions(_SILENCE_CXX17_CODECVT_HEADER_DEPRECATION_WARNING)
+    # /MP is MSVC-only: icx rejects it outright once offloading is enabled.
    add_compile_options(
-        $<$<COMPILE_LANGUAGE:C>:/MP>
+        $<$<AND:$<COMPILE_LANGUAGE:C>,$<C_COMPILER_ID:MSVC>>:/MP${SD_MSVC_MP_COUNT}>
        $<$<COMPILE_LANGUAGE:C>:/utf-8>
-        $<$<COMPILE_LANGUAGE:CXX>:/MP>
+        $<$<AND:$<COMPILE_LANGUAGE:CXX>,$<CXX_COMPILER_ID:MSVC>>:/MP${SD_MSVC_MP_COUNT}>
        $<$<COMPILE_LANGUAGE:CXX>:/utf-8>
    )
 endif()
@@ -94,6 +110,8 @@ option(SD_MUSA                       "sd: musa backend" OFF)
 option(SD_BUILD_SHARED_LIBS          "sd: build shared libs" OFF)
 option(SD_BUILD_SHARED_GGML_LIB      "sd: build ggml as a separate shared lib" OFF)
 option(SD_USE_SYSTEM_GGML            "sd: use system-installed GGML library" OFF)
+option(SD_USE_UPSTREAM_GGML          "sd: build with upstream GGML instead of the patched GGML extensions" OFF)
+set(SD_GGML_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/ggml" CACHE PATH "sd: ggml source directory (also supplies private headers for system ggml)")
 #option(SD_BUILD_SERVER               "sd: build server example"                           ON)

 set(CMAKE_C_STANDARD 11)
@@ -108,6 +126,12 @@ if(SD_CUDA)
 endif()

 if(SD_METAL)
+    if(NOT SD_USE_UPSTREAM_GGML)
+        if(CMAKE_VERSION VERSION_LESS 3.16)
+            message(FATAL_ERROR "SD_METAL with SD_USE_UPSTREAM_GGML=OFF requires CMake 3.16 or later for OBJC and OBJCXX support")
+        endif()
+        enable_language(OBJC OBJCXX)
+    endif()
    message("-- Use Metal as backend stable-diffusion")
    set(GGML_METAL ON)
 endif()
@@ -125,6 +149,12 @@ endif ()
 if (SD_HIPBLAS)
    message("-- Use HIPBLAS as backend stable-diffusion")
    set(GGML_HIP ON)
+    # ggml-hip's device-stub objects must be position-independent, or the
+    # default-PIE sd-cli link fails with `relocation R_X86_64_32 ... cannot be
+    # used when making a PIE object` on distros that default to PIE
+    # (Ubuntu 24.04, Fedora 40+, Debian 12+). The shared-library branch below
+    # already sets this; the static build (the HIP default) did not.
+    set(CMAKE_POSITION_INDEPENDENT_CODE ON)
 endif ()

 if(SD_MUSA)
@@ -198,12 +228,39 @@ if(SD_WEBM)
    endif()
 endif()

+if (SD_RPC)
+    message("-- Use RPC as backend stable-diffusion")
+    set(GGML_RPC ON)
+    add_definitions(-DSD_USE_RPC)
+endif ()
+
 set(SD_LIB stable-diffusion)

 file(GLOB SD_LIB_SOURCES CONFIGURE_DEPENDS
    "src/*.h"
    "src/*.cpp"
    "src/*.hpp"
+    "src/conditioning/*.h"
+    "src/conditioning/*.cpp"
+    "src/conditioning/*.hpp"
+    "src/core/*.h"
+    "src/core/*.cpp"
+    "src/core/*.hpp"
+    "src/core/cpu/*.h"
+    "src/core/cpu/*.cpp"
+    "src/core/probe/*.h"
+    "src/core/probe/*.cpp"
+    "src/extensions/*.h"
+    "src/extensions/*.cpp"
+    "src/extensions/*.hpp"
+    "src/model/*/*.h"
+    "src/model/*/*.cpp"
+    "src/model/*/*.hpp"
+    "src/pipeline/*.h"
+    "src/pipeline/*.cpp"
+    "src/runtime/*.h"
+    "src/runtime/*.cpp"
+    "src/runtime/*.hpp"
    "src/model_io/*.h"
    "src/model_io/*.cpp"
    "src/tokenizers/*.h"
@@ -262,6 +319,12 @@ else()
    add_library(${SD_LIB} STATIC ${SD_LIB_SOURCES})
 endif()

+if(MSVC)
+    target_compile_options(${SD_LIB} PRIVATE $<$<COMPILE_LANGUAGE:CXX>:/bigobj>)
+    # ggml backends can throw C++ exceptions through their C API.
+    target_compile_options(${SD_LIB} PRIVATE $<$<AND:$<COMPILE_LANGUAGE:CXX>,$<CXX_COMPILER_ID:MSVC>>:/EHsc->)
+endif()
+
 if(APPLE)
    sd_set_macos_rpaths(${SD_LIB})
 endif()
@@ -285,28 +348,81 @@ set(CMAKE_POLICY_DEFAULT_CMP0077 NEW)

 if (NOT SD_USE_SYSTEM_GGML)
    # see https://github.com/ggerganov/ggml/pull/682
-    add_definitions(-DGGML_MAX_NAME=128)
+    add_definitions(-DGGML_MAX_NAME=160)
 endif()

 # deps
-# Only add ggml if it hasn't been added yet
-if (NOT TARGET ggml)
-    if (SD_USE_SYSTEM_GGML)
-        find_package(ggml REQUIRED)
-        if (NOT ggml_FOUND)
-            message(FATAL_ERROR "System-installed GGML library not found.")
-        endif()
-        add_library(ggml ALIAS ggml::ggml)
-    else()
-        add_subdirectory(ggml)
-    endif()
-endif()
+include(cmake/ggml.cmake)

 add_subdirectory(thirdparty)

-target_link_libraries(${SD_LIB} PUBLIC ggml zip)
-target_include_directories(${SD_LIB} PUBLIC . src include)
-target_include_directories(${SD_LIB} PUBLIC . thirdparty)
+target_sources(${SD_LIB} PRIVATE $<TARGET_OBJECTS:zip>)
+target_link_libraries(${SD_LIB} PUBLIC ggml)
+find_package(Threads REQUIRED)
+target_link_libraries(${SD_LIB} PRIVATE Threads::Threads)
+target_link_libraries(${SD_LIB} PRIVATE onig sd-utf8proc)
+if (SD_CUDA)
+    find_package(CUDAToolkit REQUIRED)
+    # Keep the driver stub on downstream link lines when no driver is installed.
+    target_link_libraries(${SD_LIB} PUBLIC CUDA::cuda_driver)
+    set_property(SOURCE src/core/ggml_extend_backend.cpp APPEND PROPERTY COMPILE_DEFINITIONS SD_USE_CUDA)
+    if (NOT SD_USE_UPSTREAM_GGML)
+        include(cmake/cutlass.cmake)
+        enable_language(CUDA)
+        set_target_properties(${SD_LIB} PROPERTIES CUDA_STANDARD 17 CUDA_STANDARD_REQUIRED ON)
+        if (SD_CUDA_CUTLASS_PATH)
+            target_sources(${SD_LIB} PRIVATE src/core/cuda/cute_attention.cu src/core/cuda/i8_projection.cu src/core/cuda/i8_gated_residual.cu src/core/cuda/i8_gemm.cu src/core/cuda/bf16_gemm.cu)
+            set_source_files_properties(src/core/cuda/cute_attention.cu src/core/cuda/i8_projection.cu src/core/cuda/i8_gated_residual.cu src/core/cuda/i8_gemm.cu src/core/cuda/bf16_gemm.cu PROPERTIES
+                INCLUDE_DIRECTORIES "${SD_CUDA_CUTLASS_PATH}/include"
+                COMPILE_OPTIONS "--use_fast_math;--extended-lambda;--expt-relaxed-constexpr")
+            if (MSVC)
+                set_property(SOURCE src/core/cuda/cute_attention.cu src/core/cuda/i8_projection.cu src/core/cuda/i8_gated_residual.cu src/core/cuda/i8_gemm.cu src/core/cuda/bf16_gemm.cu APPEND PROPERTY COMPILE_OPTIONS -Xcompiler=/Zc:preprocessor,/Zc:__cplusplus)
+            endif()
+            if (GGML_CUDA_FORCE_CUBLAS)
+                set_property(SOURCE src/core/cuda/i8_projection.cu src/core/cuda/i8_gated_residual.cu src/core/cuda/i8_gemm.cu src/core/cuda/bf16_gemm.cu APPEND PROPERTY COMPILE_DEFINITIONS SD_CUDA_FORCE_CUBLAS)
+            endif()
+        endif()
+    endif()
+endif()
+set(SD_BACKEND_REGISTRATIONS)
+if (SD_METAL AND NOT SD_USE_UPSTREAM_GGML)
+    add_subdirectory(src/core/metal)
+    target_link_libraries(${SD_LIB} PRIVATE "-framework Foundation" "-framework Metal")
+    list(APPEND SD_BACKEND_REGISTRATIONS sd_ggml_register_metal_i8)
+endif()
+if (NOT SD_USE_UPSTREAM_GGML AND (SD_CUDA OR SD_HIPBLAS OR SD_MUSA))
+    if(SD_HIPBLAS)
+        include(cmake/hip.cmake)
+    endif()
+    add_subdirectory(src/core/cuda)
+    foreach(VENDOR CUDA HIP MUSA)
+        if(SD_${VENDOR} OR (VENDOR STREQUAL "HIP" AND SD_HIPBLAS))
+            string(TOLOWER "${VENDOR}" NAMESPACE)
+            list(APPEND SD_BACKEND_REGISTRATIONS "sd_${NAMESPACE}::register_extensions")
+        endif()
+    endforeach()
+endif()
+if (SD_VULKAN AND NOT SD_USE_UPSTREAM_GGML)
+    add_subdirectory(src/core/vulkan)
+    list(APPEND SD_BACKEND_REGISTRATIONS sd_ggml_register_vulkan_rope sd_ggml_register_vulkan_i8 sd_ggml_register_vulkan_modulate sd_ggml_register_vulkan_rms_affine)
+endif()
+if(SD_BACKEND_REGISTRATIONS)
+    string(REPLACE ";" ",\n    " SD_BACKEND_REGISTRATIONS "${SD_BACKEND_REGISTRATIONS}")
+    configure_file(cmake/backend_extensions.h.in generated/sd_backend_extensions.h @ONLY)
+    set_property(SOURCE src/core/ggml_extensions.cpp APPEND PROPERTY INCLUDE_DIRECTORIES "${CMAKE_CURRENT_BINARY_DIR}/generated")
+    set_property(SOURCE src/core/ggml_extensions.cpp APPEND PROPERTY COMPILE_DEFINITIONS SD_BACKEND_EXTENSIONS)
+endif()
+if (SD_HIPBLAS)
+    find_package(hip REQUIRED CONFIG HINTS "$ENV{ROCM_PATH}" /opt/rocm)
+    target_link_libraries(${SD_LIB} PRIVATE hip::host)
+    set_property(SOURCE src/core/ggml_extend_backend.cpp APPEND PROPERTY COMPILE_DEFINITIONS SD_USE_HIP)
+    if (GGML_HIP_MMQ_MFMA)
+        set_property(SOURCE src/core/ggml_extend_backend.cpp APPEND PROPERTY COMPILE_DEFINITIONS SD_HIP_MFMA)
+    endif()
+endif()
+target_include_directories(${SD_LIB} PUBLIC include)
+target_include_directories(${SD_LIB} PRIVATE src src/core)
+target_include_directories(${SD_LIB} PUBLIC thirdparty)
 target_compile_features(${SD_LIB} PUBLIC c_std_11 cxx_std_17)


@@ -314,7 +430,58 @@ if (SD_BUILD_EXAMPLES)
    add_subdirectory(examples)
 endif()

-set(SD_PUBLIC_HEADERS include/stable-diffusion.h)
-set_target_properties(${SD_LIB} PROPERTIES PUBLIC_HEADER "${SD_PUBLIC_HEADERS}")

-install(TARGETS ${SD_LIB} LIBRARY PUBLIC_HEADER)
+
+#
+# install
+#
+
+include(CMakePackageConfigHelpers)
+include(GNUInstallDirs)
+
+set(SD_INSTALL_VERSION "${SDCPP_BUILD_VERSION}")
+set(SD_INSTALL_COMMIT  "${SDCPP_BUILD_COMMIT}")
+set(SD_SHARED_LIB      ${SD_BUILD_SHARED_LIBS})
+
+set(SD_INCLUDE_INSTALL_DIR ${CMAKE_INSTALL_INCLUDEDIR} CACHE PATH "Location of header  files")
+set(SD_LIB_INSTALL_DIR     ${CMAKE_INSTALL_LIBDIR}     CACHE PATH "Location of library files")
+set(SD_BIN_INSTALL_DIR     ${CMAKE_INSTALL_BINDIR}     CACHE PATH "Location of binary  files")
+
+set(SD_PUBLIC_HEADERS
+    ${CMAKE_CURRENT_SOURCE_DIR}/include/stable-diffusion.h)
+
+set_target_properties(${SD_LIB}
+    PROPERTIES
+        PUBLIC_HEADER "${SD_PUBLIC_HEADERS}")
+
+
+install(TARGETS ${SD_LIB}
+        ARCHIVE
+        LIBRARY
+        RUNTIME
+        PUBLIC_HEADER)
+
+
+configure_package_config_file(
+        ${CMAKE_CURRENT_SOURCE_DIR}/cmake/stable-diffusion-config.cmake.in
+        ${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion-config.cmake
+    INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/stable-diffusion
+    PATH_VARS SD_INCLUDE_INSTALL_DIR
+              SD_LIB_INSTALL_DIR
+              SD_BIN_INSTALL_DIR )
+
+write_basic_package_version_file(
+        ${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion-version.cmake
+    VERSION ${SD_INSTALL_VERSION}
+    COMPATIBILITY SameMajorVersion)
+
+install(FILES ${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion-config.cmake
+              ${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion-version.cmake
+        DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/stable-diffusion)
+
+configure_file(cmake/stable-diffusion.pc.in
+        "${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion.pc"
+        @ONLY)
+
+install(FILES "${CMAKE_CURRENT_BINARY_DIR}/stable-diffusion.pc"
+        DESTINATION ${CMAKE_INSTALL_LIBDIR}/pkgconfig)
@@ -12,8 +12,14 @@ If you want to update a third-party dependency, please open an issue first inste

 ## Pull Requests

+When contributing from a fork, use a fork under your personal GitHub account and enable **Allow edits from maintainers**. This lets maintainers make follow-up fixes directly on the PR branch.
+
+PRs from organization-owned forks are automatically closed when opened or reopened because GitHub does not support this maintainer-edit option for those forks. Submit the changes from a personal fork instead. See [GitHub's documentation](https://docs.github.com/en/pull-requests/how-tos/work-with-forks/allowing-changes-to-a-pull-request-branch-created-from-a-fork).
+
 Keep each PR focused on one clear change. Large or overly complex PRs are harder to review and may not be merged.

+Do not include test code or test scripts in commits or PRs. Keep them local and report verification results in the PR description.
+
 Follow Conventional Commit-style subjects seen in history: `feat:`, `fix:`, `refactor:`, `ci:`, `docs:`, `chore:`. Keep subjects imperative and scoped.

 PRs should include:
@@ -23,11 +29,38 @@ PRs should include:
 - Linked issue/PR context when applicable.
 - Screenshots or sample outputs for UI/visual behavior changes.

+## Documentation
+
+Write model usage guides in `docs/` for users. Focus on supported weights,
+setup, runnable commands, parameter meanings, output behavior, and practical
+limits such as memory requirements. Include technical details only when they
+help users choose settings or troubleshoot a problem.
+
+Keep tensor layouts, internal class names, kernel implementation details,
+algorithm derivations, and development history out of usage guides. Keep task-specific
+profiling logs, numerical comparison procedures, and test artifacts local, and
+summarize relevant verification results in the PR description.
+
+## README UI Listings
+
+The [UIs list in the README](README.md#uis) only includes open-source projects that use `stable-diffusion.cpp` as a backend. Do not submit PRs to add closed-source UIs.
+
+## Server Recipes
+
+Before adding or changing a built-in server recipe, read and follow the
+[built-in recipe policy](examples/server/recipes/README.md). The policy governs
+native editing, community editing LoRAs, distinct task checkpoints, external
+adapters, native components and layer-count variants. Do not add recipes solely
+for style or acceleration LoRAs, separately stored native components, or added
+or removed model layers. Explain the task, component compatibility or capability
+difference that justifies a new recipe, and verify its required components,
+feature declarations and embedded API representation.
+
 ## Code Style

 Format code according to the repository style before submitting changes.

-Formatting follows `.clang-format` (Chromium base, 4-space indent, no tabs). Run `format-code.sh` before opening a PR. Keep C++ standard at C++17-compatible patterns used in this repo.
+Formatting follows `.clang-format` (Chromium base, 4-space indent, no tabs). Run `scripts/format-code.sh` or `scripts/format-code.ps1` before opening a PR. Keep C++ standard at C++17-compatible patterns used in this repo.

 Naming conventions:

@@ -35,15 +68,31 @@ Naming conventions:
 - In `PascalCase` names, preserve common abbreviations in uppercase, for example `SD`, `API`, `HTTP`, `JSON`, `RGB`, `VAE`, `TAE`, `LoRA`, and `WebP`.
 - Use `snake_case` for functions, methods, variables, and file names unless an existing API requires a different style.
 - Use a trailing underscore for private data member names, for example `hidden_size_` or `tokenizer_`.
- Use `.h` for C and C++ header files. Do not introduce new `.hpp` headers.
+- Use `.hpp` for model headers under `src/model/`, including new model headers. Do not rename these headers to `.h`. Use `.h` for other C and C++ header files.
 - Use macro-based header include guards instead of `#pragma once`.
 - Format header include guards as `__SD_{PATH}__`, where `{PATH}` is the header path in uppercase snake case without the file extension. For example, `src/sample.h` should use `__SD_SAMPLE_H__`.
 - Do not introduce anonymous namespaces in new or modified code; prefer `static` file-local functions/variables or an explicit named namespace when scoping is needed.
 - In `class`/`struct` definitions, place data members before member functions unless an existing type already clearly follows a different pattern.
- Keep `test_*.cpp` / `test_*.py` naming for tests.

 Some older code in the project may not fully follow the current conventions. Please do not submit PRs that only rewrite existing code to match style rules.

+When adding or modifying model implementations, follow the model config and weight detection conventions in [docs/model_config.md](docs/model_config.md).
+
+## Tokenizer Data
+
+New model integrations must use an external `tokenizer.json` by default. Do not
+embed new vocabularies or merge tables solely for less widely used models;
+these tables increase the binary size for every user.
+
+The embedded-data allowlist is CLIP, T5/UMT5, Qwen 2/3, Mistral, and Gemma 3/4.
+Models may reuse an existing embedded tokenizer when its vocabulary and behavior
+match their text encoder. Gemma 2 and GPT-OSS require external JSON files.
+
+Adding to this allowlist requires maintainer approval, supported by the model's
+usage, reuse across models, and measured binary-size cost. Document the matching
+JSON and CLI option for models that require an external tokenizer, and fail
+initialization clearly when it is missing.
+
 ## AI-Assisted Contributions

 AI tools may be used to assist development, but contributors are responsible for the quality and correctness of the submitted code.
@@ -1,23 +0,0 @@
-ARG UBUNTU_VERSION=24.04
-
-FROM ubuntu:$UBUNTU_VERSION AS build
-
-RUN apt-get update && apt-get install -y --no-install-recommends build-essential git cmake
-
-WORKDIR /sd.cpp
-
-COPY . .
-
-RUN cmake . -B ./build
-RUN cmake --build ./build --config Release --parallel
-
-FROM ubuntu:$UBUNTU_VERSION AS runtime
-
-RUN apt-get update && \
-    apt-get install --yes --no-install-recommends libgomp1 && \
-    apt-get clean
-
-COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
-COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
-
-ENTRYPOINT [ "/sd-cli" ]
@@ -1,31 +0,0 @@
-ARG CUDA_VERSION=12.6.3
-ARG UBUNTU_VERSION=24.04
-
-FROM nvidia/cuda:${CUDA_VERSION}-cudnn-devel-ubuntu${UBUNTU_VERSION} AS build
-
-RUN apt-get update && apt-get install -y --no-install-recommends build-essential git ccache cmake
-
-WORKDIR /sd.cpp
-
-COPY . .
-
-ARG CUDACXX=/usr/local/cuda/bin/nvcc
-ARG CUDA_ARCHITECTURES=""
-ARG GGML_CUDA_FA_ALL_QUANTS=""
-
-RUN cmake . -B ./build \
-    -DSD_CUDA=ON \
-    ${CUDA_ARCHITECTURES:+-DCMAKE_CUDA_ARCHITECTURES="${CUDA_ARCHITECTURES}"} \
-    ${GGML_CUDA_FA_ALL_QUANTS:+-DGGML_CUDA_FA_ALL_QUANTS=${GGML_CUDA_FA_ALL_QUANTS}}
-RUN cmake --build ./build --config Release -j$(nproc)
-
-FROM nvidia/cuda:${CUDA_VERSION}-cudnn-runtime-ubuntu${UBUNTU_VERSION} AS runtime
-
-RUN apt-get update && \
-    apt-get install --yes --no-install-recommends libgomp1 && \
-    apt-get clean
-
-COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
-COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
-
-ENTRYPOINT [ "/sd-cli" ]
@@ -1,24 +0,0 @@
-ARG MUSA_VERSION=rc4.2.0
-ARG UBUNTU_VERSION=22.04
-
-FROM mthreads/musa:${MUSA_VERSION}-devel-ubuntu${UBUNTU_VERSION}-amd64 as build
-
-RUN apt-get update && apt-get install -y ccache cmake git
-
-WORKDIR /sd.cpp
-
-COPY . .
-
-RUN mkdir build && cd build && \
-    cmake .. -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ \
-        -DCMAKE_C_FLAGS="${CMAKE_C_FLAGS} -fopenmp -I/usr/lib/llvm-14/lib/clang/14.0.0/include -L/usr/lib/llvm-14/lib" \
-        -DCMAKE_CXX_FLAGS="${CMAKE_CXX_FLAGS} -fopenmp -I/usr/lib/llvm-14/lib/clang/14.0.0/include -L/usr/lib/llvm-14/lib" \
-        -DSD_MUSA=ON -DCMAKE_BUILD_TYPE=Release && \
-    cmake --build . --config Release
-
-FROM mthreads/musa:${MUSA_VERSION}-runtime-ubuntu${UBUNTU_VERSION}-amd64 as runtime
-
-COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
-COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
-
-ENTRYPOINT [ "/sd-cli" ]
@@ -1,21 +0,0 @@
-# ggml SYCL hardware detection uses BMG G31/WCL architecture enums added in oneAPI 2025.3.
-ARG SYCL_VERSION=2025.3.2-0
-
-FROM intel/oneapi-basekit:${SYCL_VERSION}-devel-ubuntu24.04 AS build
-
-RUN apt-get update && apt-get install -y cmake
-
-WORKDIR /sd.cpp
-
-COPY . .
-
-RUN mkdir build && cd build && \
-    cmake .. -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DSD_SYCL=ON -DCMAKE_BUILD_TYPE=Release && \
-    cmake --build . --config Release -j$(nproc)
-
-FROM intel/oneapi-basekit:${SYCL_VERSION}-devel-ubuntu24.04 AS runtime
-
-COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
-COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
-
-ENTRYPOINT [ "/sd-cli" ]
@@ -1,23 +0,0 @@
-ARG UBUNTU_VERSION=24.04
-
-FROM ubuntu:$UBUNTU_VERSION AS build
-
-RUN apt-get update && apt-get install -y --no-install-recommends build-essential git cmake libvulkan-dev glslc spirv-headers
-
-WORKDIR /sd.cpp
-
-COPY . .
-
-RUN cmake . -B ./build -DSD_VULKAN=ON
-RUN cmake --build ./build --config Release --parallel
-
-FROM ubuntu:$UBUNTU_VERSION AS runtime
-
-RUN apt-get update && \
-    apt-get install --yes --no-install-recommends libgomp1 libvulkan1 mesa-vulkan-drivers && \
-    apt-get clean
-
-COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
-COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
-
-ENTRYPOINT [ "/sd-cli" ]
@@ -15,6 +15,16 @@ API and command-line option may change frequently.***

 ## 🔥Important News

+* **2026/11/11** 🚀 stable-diffusion.cpp adds **Day-0 support for MiniMax-H3-Image**
+* **2026/10/16** 🚀 stable-diffusion.cpp adds **Day-0 support for MOSS-Image-Base / MOSS-Image-SFT**
+* **2026/10/16** 🚀 stable-diffusion.cpp adds **Day-0 support for FLUX.3 Image Klein / Klein Base**
+* **2026/09/20** 🚀 stable-diffusion.cpp adds **Day-0 support for Qwen-Image-2.1**
+* **2026/08/20** 🚀 stable-diffusion.cpp now supports **LTX-2.5**
+* **2026/08/04** 🚀 stable-diffusion.cpp adds **Day-1 support for MiniMax-H3**
+* **2026/06/25** 🚀 stable-diffusion.cpp now supports **Krea2**
+* **2026/06/04** 🚀 stable-diffusion.cpp now supports **Ideogram4**
+* **2026/05/31** 🚀 stable-diffusion.cpp now supports **PiD**
+* **2026/05/27** 🚀 stable-diffusion.cpp now supports **Lens**
 * **2026/05/17** 🚀 stable-diffusion.cpp now supports **LTX-2.3**
 * **2026/04/11** 🚀 stable-diffusion.cpp now uses a brand-new embedded web UI.  
 * **2026/01/18** 🚀 stable-diffusion.cpp now supports **FLUX.2-klein**  
@@ -31,32 +41,66 @@ API and command-line option may change frequently.***
 - Super lightweight and without external dependencies
 - Supported models
  - Image Models
-    - SD1.x, SD2.x, [SD-Turbo](https://huggingface.co/stabilityai/sd-turbo)
-    - SDXL, [SDXL-Turbo](https://huggingface.co/stabilityai/sdxl-turbo)
+    - [SD1.x, SD2.x, SD-Turbo](./docs/sd.md)
+    - [SDXL, SDXL-Turbo](./docs/sd.md)
    - [Some SD1.x and SDXL distilled models](./docs/distilled_sd.md)
    - [SD3/SD3.5](./docs/sd3.md)
    - [FLUX.1-dev/FLUX.1-schnell](./docs/flux.md)
    - [FLUX.2-dev/FLUX.2-klein](./docs/flux2.md)
+    - [FLUX.3 Image Klein / Klein Base](./docs/flux3.md)
+    - [MOSS-Image-Base / MOSS-Image-SFT](./docs/moss_image.md)
+    - [Fibo 2.0 Dev](./docs/fibo2.md)
+    - [Lens](./docs/lens.md)
    - [Chroma](./docs/chroma.md)
    - [Chroma1-Radiance](./docs/chroma_radiance.md)
    - [Qwen Image](./docs/qwen_image.md)
+    - [Qwen Image 2.1](./docs/qwen_image_2.1.md)
+    - [PiD](./docs/pid.md)
+    - [LongCat Image](./docs/longcat_image.md)
    - [Z-Image](./docs/z_image.md)
+    - [Z-Image L2P](./docs/z_image_l2p.md)
+    - [MiniT2I](./docs/minit2i.md)
+    - [P-H3-Image](./docs/p_h3_image.md)
+    - [MiniMax-H3-Image](./docs/minimax_h3_image.md)
+    - [SenseNova U1.5](./docs/sensenova_u1.md)
    - [Ovis-Image](./docs/ovis_image.md)
    - [Anima](./docs/anima.md)
    - [ERNIE-Image](./docs/ernie_image.md)
+    - [Boogu Image](./docs/boogu_image.md)
+    - [Krea2](./docs/krea2.md)
+    - [Mage-Flow](./docs/mage_flow.md)
+    - [SeFi-Image](./docs/sefi_image.md)
    - [HiDream-O1-Image](./docs/hidream_o1_image.md)
-  - Image Edit Models
+    - [Ideogram4](./docs/ideogram4.md)
+    - [LLaDA-Image](./docs/llada_image.md)
+    - [Ming-Image Design](./docs/ming_image.md)
+    - [PixArt](./docs/pixart.md)
+    - [Iris-3B](./docs/iris.md)
+  - [Image Edit Models](./docs/edit.md)
    - [FLUX.1-Kontext-dev](./docs/kontext.md)
    - [Qwen Image Edit series](./docs/qwen_image_edit.md)
+    - [LongCat Image Edit](./docs/longcat_image.md)
+    - [Boogu Image Edit](./docs/boogu_image.md)
+    - [Mage-Flow-Edit](./docs/mage_flow.md#image-editing)
+    - [LLaDA-Image Edit](./docs/llada_image.md#image-editing)
+    - [Fibo 2.0 Dev / Edit Turbo](./docs/fibo2.md#editing-and-turbo)
  - Video Models
    - [Wan2.1/Wan2.2](./docs/wan.md)
-    - [LTX-2.3](./docs/ltx2.md)
-  - [PhotoMaker](https://github.com/TencentARC/PhotoMaker) support.
+    - [MiniMax-H3](./docs/minimax_h3.md)
+    - [LTX-2.3/LTX-2.5](./docs/ltx2.md)
+    - [HunyuanVideo 1.5](./docs/hunyuan_video.md)
+    - [LingBot-Video](./docs/lingbot_video.md)
+  - Video Upscale Models
+    - [RelayVSR](./docs/relay_vsr.md)
+    - [SeedVR2](./docs/seedvr2.md)
+  - [PhotoMaker](./docs/photo_maker.md) support.
+  - [IP-Adapter](./docs/ip_adapter.md) support (SD 1.5 and SDXL, including Plus)
  - 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)
  - Latent Consistency Models support (LCM/LCM-LoRA)
-  - Faster and memory efficient latent decoding with [TAESD](https://github.com/madebyollin/taesd)
-  - Upscale images generated with [ESRGAN](https://github.com/xinntao/Real-ESRGAN)
+  - Faster and memory efficient latent decoding with [TAESD](./docs/taesd.md)
+  - Upscale images generated with [ESRGAN](./docs/esrgan.md)
 - Supported backends
  - CPU (AVX, AVX2 and AVX512 support for x86 architectures)
  - CUDA
@@ -106,7 +150,7 @@ API and command-line option may change frequently.***
    - Stable Diffusion v1.5 from https://huggingface.co/stable-diffusion-v1-5/stable-diffusion-v1-5 

    ```sh
-    curl -L -O https://huggingface.co/runwayml/stable-diffusion-v1-5/resolve/main/v1-5-pruned-emaonly.safetensors
+    curl -L -O https://huggingface.co/stable-diffusion-v1-5/stable-diffusion-v1-5/resolve/main/v1-5-pruned-emaonly.safetensors
    ```

 ### Generate an image with just one command
@@ -124,29 +168,14 @@ For runtime and parameter backend placement, see the [backend selection guide](.

 ## More Guides

+- [Troubleshooting](./docs/troubleshooting.md)
 - [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)
- [FLUX.2-dev/FLUX.2-klein](./docs/flux2.md)
- [FLUX.1-Kontext-dev](./docs/kontext.md)
- [Chroma](./docs/chroma.md)
- [🔥Qwen Image](./docs/qwen_image.md)
- [🔥Qwen Image Edit series](./docs/qwen_image_edit.md)
- [🔥Wan2.1/Wan2.2](./docs/wan.md)
- [🔥LTX-2.3](./docs/ltx2.md)
- [🔥Z-Image](./docs/z_image.md)
- [Ovis-Image](./docs/ovis_image.md)
- [Anima](./docs/anima.md)
- [ERNIE-Image](./docs/ernie_image.md)
- [HiDream-O1-Image](./docs/hidream_o1_image.md)
+- [RPC](./docs/rpc.md)
 - [LoRA](./docs/lora.md)
 - [LCM/LCM-LoRA](./docs/lcm.md)
- [Using PhotoMaker to personalize image generation](./docs/photo_maker.md)
- [Using ESRGAN to upscale results](./docs/esrgan.md)
- [Using TAESD to faster decoding](./docs/taesd.md)
 - [Docker](./docs/docker.md)
 - [Quantization and GGUF](./docs/quantization_and_gguf.md)
+- [INT8 convrot safetensors](./docs/int8_convrot.md)
 - [Inference acceleration via caching](./docs/caching.md)

 ## Bindings
@@ -165,7 +194,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.

- [Jellybox](https://jellybox.com)
+- [GIMP Plugins](https://github.com/themanyone/gimp-plugins)
 - [Stable Diffusion GUI](https://github.com/fszontagh/sd.cpp.gui.wx)
 - [Stable Diffusion CLI-GUI](https://github.com/piallai/stable-diffusion.cpp)
 - [Local Diffusion](https://github.com/rmatif/Local-Diffusion)
@@ -173,29 +202,13 @@ These projects use `stable-diffusion.cpp` as a backend for their image generatio
 - [LocalAI](https://github.com/mudler/LocalAI)
 - [Neural-Pixel](https://github.com/Luiz-Alcantara/Neural-Pixel)
 - [KoboldCpp](https://github.com/LostRuins/koboldcpp)
+- [Kokon](https://github.com/GabrieleLozupone/kokon)
+- [NodeTool](https://github.com/nodetool-ai/nodetool)
+- [Quo](https://github.com/quo-ai/quo)
+- [Prisme](https://github.com/d3Ita/prisme)

 ## Contributors

 Thank you to all the people who have already contributed to stable-diffusion.cpp!

 [![Contributors](https://contrib.rocks/image?repo=leejet/stable-diffusion.cpp)](https://github.com/leejet/stable-diffusion.cpp/graphs/contributors)
-
-## Star History
-
-[![Star History Chart](https://api.star-history.com/svg?repos=leejet/stable-diffusion.cpp&type=Date)](https://star-history.com/#leejet/stable-diffusion.cpp&Date)
-
-## References
-
- [ggml](https://github.com/ggml-org/ggml)
- [diffusers](https://github.com/huggingface/diffusers)
- [stable-diffusion](https://github.com/CompVis/stable-diffusion)
- [sd3-ref](https://github.com/Stability-AI/sd3-ref)
- [stable-diffusion-stability-ai](https://github.com/Stability-AI/stablediffusion)
- [stable-diffusion-webui](https://github.com/AUTOMATIC1111/stable-diffusion-webui)
- [ComfyUI](https://github.com/comfyanonymous/ComfyUI)
- [k-diffusion](https://github.com/crowsonkb/k-diffusion)
- [latent-consistency-model](https://github.com/luosiallen/latent-consistency-model)
- [generative-models](https://github.com/Stability-AI/generative-models/)
- [PhotoMaker](https://github.com/TencentARC/PhotoMaker)
- [Wan2.1](https://github.com/Wan-Video/Wan2.1)
- [Wan2.2](https://github.com/Wan-Video/Wan2.2)
@@ -0,0 +1,17 @@
+#ifndef __SD_BACKEND_EXTENSIONS_H__
+#define __SD_BACKEND_EXTENSIONS_H__
+
+#include "core/cuda/registration.h"
+
+void sd_ggml_register_metal_i8(ggml_backend_dev_t device);
+
+void sd_ggml_register_vulkan_rope(ggml_backend_dev_t device);
+void sd_ggml_register_vulkan_i8(ggml_backend_dev_t device);
+void sd_ggml_register_vulkan_rms_affine(ggml_backend_dev_t device);
+void sd_ggml_register_vulkan_modulate(ggml_backend_dev_t device);
+
+static void (*const sd_backend_extensions[])(ggml_backend_dev_t) = {
+    @SD_BACKEND_REGISTRATIONS@
+};
+
+#endif
@@ -0,0 +1,20 @@
+set(sd_cutlass_default "${PROJECT_SOURCE_DIR}/thirdparty/cutlass")
+if(NOT DEFINED SD_CUDA_CUTLASS_PATH AND DEFINED GGML_CUDA_CUTLASS_PATH)
+    get_filename_component(sd_old_cutlass_default "${SD_GGML_SOURCE_DIR}/src/ggml-cuda/vendors/cutlass" ABSOLUTE)
+    if(NOT GGML_CUDA_CUTLASS_PATH STREQUAL sd_old_cutlass_default)
+        set(sd_cutlass_default "${GGML_CUDA_CUTLASS_PATH}")
+    endif()
+    message(STATUS "Migrating GGML_CUDA_CUTLASS_PATH to SD_CUDA_CUTLASS_PATH")
+endif()
+set(SD_CUDA_CUTLASS_PATH "${sd_cutlass_default}" CACHE PATH "sd: CUTLASS source directory for INT8 GEMM and CuTe attention, empty to disable")
+unset(GGML_CUDA_CUTLASS_PATH CACHE)
+if(SD_CUDA_CUTLASS_PATH)
+    if(NOT EXISTS "${SD_CUDA_CUTLASS_PATH}/include/cutlass/gemm/kernel/default_gemm_universal_with_visitor.h" OR
+       NOT EXISTS "${SD_CUDA_CUTLASS_PATH}/include/cute/tensor.hpp")
+        message(FATAL_ERROR
+            "CUTLASS headers not found at ${SD_CUDA_CUTLASS_PATH}. "
+            "Run git submodule update --init --recursive from the sd.cpp repository root, "
+            "set SD_CUDA_CUTLASS_PATH to a CUTLASS source tree (tested with 3.9.2), "
+            "or set SD_CUDA_CUTLASS_PATH to an empty string to disable CUTLASS.")
+    endif()
+endif()
@@ -0,0 +1,3 @@
+file(READ "${INPUT}" SPIRV HEX)
+string(REGEX REPLACE "(..)(..)(..)(..)" "0x\\4\\3\\2\\1," SPIRV_WORDS "${SPIRV}")
+file(WRITE "${OUTPUT}" "#include <cstdint>\nnamespace sd_vulkan_shaders {\nstatic const uint32_t ${NAME}[] = {${SPIRV_WORDS}};\n}\n")
@@ -0,0 +1,29 @@
+if(NOT TARGET ggml AND NOT TARGET ggml::ggml)
+    if(SD_USE_SYSTEM_GGML)
+        find_package(ggml REQUIRED)
+    else()
+        add_subdirectory("${SD_GGML_SOURCE_DIR}" "${CMAKE_CURRENT_BINARY_DIR}/ggml")
+    endif()
+endif()
+if(NOT TARGET ggml)
+    add_library(ggml ALIAS ggml::ggml)
+endif()
+
+if(SD_USE_UPSTREAM_GGML)
+    get_target_property(sd_ggml_imported ggml IMPORTED)
+    if(sd_ggml_imported)
+        set(sd_ggml_private_include "${SD_GGML_SOURCE_DIR}/src")
+    else()
+        get_target_property(sd_ggml_private_include ggml SOURCE_DIR)
+    endif()
+    if(NOT EXISTS "${sd_ggml_private_include}/ggml-impl.h")
+        message(FATAL_ERROR "Set SD_GGML_SOURCE_DIR to the source tree matching the selected ggml library (ggml-impl.h is required).")
+    endif()
+    target_include_directories(${SD_LIB} PRIVATE "${sd_ggml_private_include}")
+    set_property(TARGET ${SD_LIB} PROPERTY SD_GGML_PRIVATE_INCLUDE_DIR "${sd_ggml_private_include}")
+endif()
+
+if(SD_USE_UPSTREAM_GGML)
+    target_compile_definitions(${SD_LIB} PUBLIC SD_USE_UPSTREAM_GGML)
+    message(WARNING "Using upstream GGML: INT8 tensorwise/convrot is disabled and FP8 weights are converted to F16 at load time. Some operators may be unsupported and performance may be lower than with patched GGML.")
+endif()
@@ -0,0 +1,12 @@
+if(NOT WIN32 AND NOT CMAKE_CXX_COMPILER MATCHES "hipcc")
+    if(NOT CMAKE_HIP_ARCHITECTURES)
+        if(GPU_TARGETS)
+            set(CMAKE_HIP_ARCHITECTURES ${GPU_TARGETS})
+        elseif(AMDGPU_TARGETS)
+            set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
+        endif()
+    endif()
+    # HIP objects are linked by the parent target, so initialize its link rules here.
+    enable_language(HIP)
+    set_property(TARGET ${SD_LIB} PROPERTY HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}")
+endif()
@@ -0,0 +1,45 @@
+set(SD_VERSION      "@SD_INSTALL_VERSION@")
+set(SD_BUILD_COMMIT "@SD_INSTALL_COMMIT@")
+set(SD_SHARED_LIB    @SD_SHARED_LIB@)
+
+@PACKAGE_INIT@
+
+set_and_check(SD_INCLUDE_DIR "@PACKAGE_SD_INCLUDE_INSTALL_DIR@")
+set_and_check(SD_LIB_DIR     "@PACKAGE_SD_LIB_INSTALL_DIR@")
+set(SD_BIN_DIR "@PACKAGE_SD_BIN_INSTALL_DIR@")
+
+include(CMakeFindDependencyMacro)
+find_dependency(ggml REQUIRED HINTS "${SD_LIB_DIR}/cmake")
+find_dependency(Threads REQUIRED)
+if(@SD_CUDA@)
+    find_dependency(CUDAToolkit REQUIRED)
+endif()
+
+if(NOT TARGET stable-diffusion)
+    find_library(stable-diffusion_LIBRARY stable-diffusion
+        REQUIRED
+        HINTS "${SD_LIB_DIR}"
+        NO_CMAKE_FIND_ROOT_PATH
+    )
+
+    add_library(stable-diffusion UNKNOWN IMPORTED)
+    set_target_properties(stable-diffusion
+        PROPERTIES
+            INTERFACE_INCLUDE_DIRECTORIES "${SD_INCLUDE_DIR}"
+            INTERFACE_LINK_LIBRARIES "ggml::ggml;Threads::Threads"
+            IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
+            IMPORTED_LOCATION "${stable-diffusion_LIBRARY}"
+            INTERFACE_COMPILE_FEATURES "c_std_11;cxx_std_17"
+            POSITION_INDEPENDENT_CODE ON)
+
+    if(@SD_CUDA@)
+        set_property(TARGET stable-diffusion APPEND PROPERTY INTERFACE_LINK_LIBRARIES CUDA::cuda_driver)
+    endif()
+
+    if(SD_SHARED_LIB)
+        target_compile_definitions(stable-diffusion
+            INTERFACE SD_BUILD_SHARED_LIB)
+    endif()
+endif()
+
+check_required_components(stable-diffusion)
@@ -0,0 +1,11 @@
+prefix=@CMAKE_INSTALL_PREFIX@
+exec_prefix=${prefix}
+libdir=@CMAKE_INSTALL_FULL_LIBDIR@
+includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@
+
+Name: stable-diffusion
+Description: Diffusion model(SD,Flux,Wan,Qwen Image,Z-Image,...) inference in pure C/C++
+Version: @SDCPP_BUILD_VERSION@
+Libs: -L${libdir} -lstable-diffusion
+Libs.private: -lggml -lggml-base @CMAKE_THREAD_LIBS_INIT@
+Cflags: -I${includedir}
@@ -0,0 +1,45 @@
+ARG UBUNTU_VERSION=24.04
+
+FROM ubuntu:$UBUNTU_VERSION AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends build-essential git cmake ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && \
+    apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && \
+    rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+COPY . .
+
+RUN cmake . -B ./build \
+    -DSD_BUILD_SHARED_LIBS=ON \
+    -DGGML_NATIVE=OFF \
+    -DSD_BUILD_SHARED_GGML_LIB=ON \
+    -DGGML_BACKEND_DL=ON \
+    -DGGML_CPU_ALL_VARIANTS=ON \
+    -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+    -DCMAKE_INSTALL_RPATH='$ORIGIN'
+RUN cmake --build ./build --config Release --parallel
+
+FROM ubuntu:$UBUNTU_VERSION AS runtime
+
+RUN apt-get update && \
+    apt-get install --yes --no-install-recommends libgomp1 && \
+    apt-get clean
+
+COPY --from=build /sd.cpp/build/bin /sd.cpp/bin
+RUN printf '#!/bin/sh\nexec /sd.cpp/bin/sd-cli "$@"\n' > /sd-cli && \
+    printf '#!/bin/sh\nexec /sd.cpp/bin/sd-server "$@"\n' > /sd-server && \
+    chmod +x /sd-cli /sd-server
+
+WORKDIR /sd.cpp
+
+ENTRYPOINT [ "/sd-cli" ]
@@ -0,0 +1,76 @@
+# syntax=docker/dockerfile:1
+
+ARG CUDA_VERSION=12.6.3
+ARG UBUNTU_VERSION=24.04
+
+FROM nvidia/cuda:${CUDA_VERSION}-devel-ubuntu${UBUNTU_VERSION} AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends build-essential git ccache cmake ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && \
+    apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && \
+    rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+COPY . .
+
+ARG CUDACXX=/usr/local/cuda/bin/nvcc
+ARG CUDA_ARCHITECTURES=""
+ARG GGML_CUDA_FA_ALL_QUANTS=""
+ARG GGML_CUDA_ENABLE_DYNAMIC_CPU_BACKENDS=ON
+
+ENV CCACHE_DIR=/root/.cache/ccache \
+    CCACHE_MAXSIZE=2G \
+    CCACHE_COMPILERCHECK=content
+
+RUN set -- \
+        -DSD_CUDA=ON \
+        -DCMAKE_C_COMPILER_LAUNCHER=ccache \
+        -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
+        -DCMAKE_CUDA_COMPILER_LAUNCHER=ccache; \
+    if [ "${GGML_CUDA_ENABLE_DYNAMIC_CPU_BACKENDS}" = "ON" ]; then \
+        set -- "$@" \
+            -DSD_BUILD_SHARED_LIBS=ON \
+            -DGGML_NATIVE=OFF \
+            -DSD_BUILD_SHARED_GGML_LIB=ON \
+            -DGGML_BACKEND_DL=ON \
+            -DGGML_CPU_ALL_VARIANTS=ON \
+            -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+            "-DCMAKE_INSTALL_RPATH=\$ORIGIN"; \
+    fi; \
+    if [ -n "${CUDA_ARCHITECTURES}" ]; then \
+        set -- "$@" "-DCMAKE_CUDA_ARCHITECTURES=${CUDA_ARCHITECTURES}"; \
+    fi; \
+    if [ -n "${GGML_CUDA_FA_ALL_QUANTS}" ]; then \
+        set -- "$@" "-DGGML_CUDA_FA_ALL_QUANTS=${GGML_CUDA_FA_ALL_QUANTS}"; \
+    fi; \
+    cmake . -B ./build "$@"
+RUN --mount=type=cache,target=/root/.cache/ccache \
+    ccache --zero-stats && \
+    cmake --build ./build --config Release -j$(nproc) && \
+    ccache --cleanup && \
+    ccache --show-stats
+
+FROM nvidia/cuda:${CUDA_VERSION}-runtime-ubuntu${UBUNTU_VERSION} AS runtime
+
+RUN apt-get update && \
+    apt-get install --yes --no-install-recommends libgomp1 && \
+    apt-get clean
+
+COPY --from=build /sd.cpp/build/bin /sd.cpp/bin
+COPY --from=build /sd.cpp/thirdparty/cutlass/LICENSE.txt /sd.cpp/bin/cutlass.txt
+RUN printf '#!/bin/sh\nexec /sd.cpp/bin/sd-cli "$@"\n' > /sd-cli && \
+    printf '#!/bin/sh\nexec /sd.cpp/bin/sd-server "$@"\n' > /sd-server && \
+    chmod +x /sd-cli /sd-server
+
+WORKDIR /sd.cpp
+
+ENTRYPOINT [ "/sd-cli" ]
@@ -0,0 +1,46 @@
+ARG MUSA_IMAGE=registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64
+
+FROM ${MUSA_IMAGE} AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends ccache cmake git ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && \
+    apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && \
+    rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+COPY . .
+
+RUN mkdir build && cd build && \
+    cmake .. -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ \
+        -DSD_MUSA=ON \
+        -DMUSA_ARCHITECTURES=31 \
+        -DSD_BUILD_SHARED_LIBS=ON \
+        -DGGML_NATIVE=OFF \
+        -DSD_BUILD_SHARED_GGML_LIB=ON \
+        -DGGML_BACKEND_DL=ON \
+        -DGGML_CPU_ALL_VARIANTS=ON \
+        -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+        -DCMAKE_INSTALL_RPATH='$ORIGIN' \
+        -DCMAKE_BUILD_TYPE=Release && \
+    cmake --build . --config Release
+
+FROM ${MUSA_IMAGE} AS runtime
+
+COPY --from=build /sd.cpp/build/bin /sd.cpp/bin
+RUN printf '#!/bin/sh\nexec /sd.cpp/bin/sd-cli "$@"\n' > /sd-cli && \
+    printf '#!/bin/sh\nexec /sd.cpp/bin/sd-server "$@"\n' > /sd-server && \
+    chmod +x /sd-cli /sd-server
+
+WORKDIR /sd.cpp
+
+ENTRYPOINT [ "/sd-cli" ]
+CMD []
@@ -0,0 +1,50 @@
+FROM rocm/dev-ubuntu-24.04 AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends \
+        build-essential git cmake ninja-build ccache ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+# ROCm's clang is not on PATH in Docker RUN shells (profile scripts are not loaded).
+ENV PATH="/opt/rocm/llvm/bin:$PATH"
+
+COPY . .
+
+ARG AMDGPU_TARGETS=""
+
+RUN cmake . -B ./build -G Ninja \
+    -DCMAKE_C_COMPILER=clang \
+    -DCMAKE_CXX_COMPILER=clang++ \
+    -DSD_HIPBLAS=ON \
+    -DGGML_NATIVE=OFF \
+    -DAMDGPU_TARGETS="${AMDGPU_TARGETS}" \
+    -DCMAKE_BUILD_TYPE=Release \
+    -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+    -DCMAKE_POSITION_INDEPENDENT_CODE=ON
+RUN cmake --build ./build --config Release --parallel
+
+FROM rocm/dev-ubuntu-24.04 AS runtime
+
+# ROCm clang links against its own OpenMP runtime in /opt/rocm/llvm/lib,
+# and the ROCm runtime libs (hipblas, rocblas, amdhip64) live in /opt/rocm/lib.
+ENV LD_LIBRARY_PATH="/opt/rocm/llvm/lib:/opt/rocm/lib:$LD_LIBRARY_PATH"
+
+RUN apt-get update && \
+    apt-get install --yes --no-install-recommends libgomp1 && \
+    apt-get clean && rm -rf /var/lib/apt/lists/*
+
+COPY --from=build /sd.cpp/build/bin /sd.cpp/bin
+RUN printf '#!/bin/sh\nexec /sd.cpp/bin/sd-cli "$@"\n' > /sd-cli && \
+    printf '#!/bin/sh\nexec /sd.cpp/bin/sd-server "$@"\n' > /sd-server && \
+    chmod +x /sd-cli /sd-server
+
+ENTRYPOINT [ "/sd-cli" ]
@@ -0,0 +1,34 @@
+# ggml SYCL hardware detection uses BMG G31/WCL architecture enums added in oneAPI 2025.3.
+ARG SYCL_VERSION=2025.3.2-0
+
+FROM intel/oneapi-basekit:${SYCL_VERSION}-devel-ubuntu24.04 AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends cmake ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && \
+    apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && \
+    rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+COPY . .
+
+RUN mkdir build && cd build && \
+    cmake .. -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DSD_SYCL=ON -DCMAKE_BUILD_TYPE=Release && \
+    cmake --build . --config Release -j$(nproc)
+
+FROM intel/oneapi-basekit:${SYCL_VERSION}-devel-ubuntu24.04 AS runtime
+
+COPY --from=build /sd.cpp/build/bin/sd-cli /sd-cli
+COPY --from=build /sd.cpp/build/bin/sd-server /sd-server
+
+WORKDIR /sd.cpp
+
+ENTRYPOINT [ "/sd-cli" ]
@@ -0,0 +1,46 @@
+ARG UBUNTU_VERSION=24.04
+
+FROM ubuntu:$UBUNTU_VERSION AS build
+
+# sd-server embeds the web UI at build time, so the build image needs Node/pnpm.
+RUN apt-get update && apt-get install -y --no-install-recommends build-essential git cmake libvulkan-dev glslc spirv-headers ca-certificates curl gnupg && \
+    mkdir -p /etc/apt/keyrings && \
+    curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key -o /tmp/nodesource-repo.gpg.key && \
+    gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg /tmp/nodesource-repo.gpg.key && \
+    rm /tmp/nodesource-repo.gpg.key && \
+    echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_24.x nodistro main" > /etc/apt/sources.list.d/nodesource.list && \
+    apt-get update && \
+    apt-get install -y --no-install-recommends nodejs && \
+    npm install -g pnpm@10.15.1 && \
+    apt-get clean && \
+    rm -rf /var/lib/apt/lists/*
+
+WORKDIR /sd.cpp
+
+COPY . .
+
+RUN cmake . -B ./build \
+    -DSD_VULKAN=ON \
+    -DSD_BUILD_SHARED_LIBS=ON \
+    -DGGML_NATIVE=OFF \
+    -DSD_BUILD_SHARED_GGML_LIB=ON \
+    -DGGML_BACKEND_DL=ON \
+    -DGGML_CPU_ALL_VARIANTS=ON \
+    -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+    -DCMAKE_INSTALL_RPATH='$ORIGIN'
+RUN cmake --build ./build --config Release -j$(nproc)
+
+FROM ubuntu:$UBUNTU_VERSION AS runtime
+
+RUN apt-get update && \
+    apt-get install --yes --no-install-recommends libgomp1 libvulkan1 mesa-vulkan-drivers libglvnd0 libgl1 libglx0 libegl1 libgles2 && \
+    apt-get clean
+
+COPY --from=build /sd.cpp/build/bin /sd.cpp/bin
+RUN printf '#!/bin/sh\nexec /sd.cpp/bin/sd-cli "$@"\n' > /sd-cli && \
+    printf '#!/bin/sh\nexec /sd.cpp/bin/sd-server "$@"\n' > /sd-server && \
+    chmod +x /sd-cli /sd-server
+
+WORKDIR /sd.cpp
+
+ENTRYPOINT [ "/sd-cli" ]
@@ -0,0 +1,143 @@
+# ADetailer
+
+`sd-cli` can run a YOLOv8 object detector on an existing or newly generated
+image, and `sd-server` can run it after image generation. Both perform a cropped
+inpaint pass for every detected object. The implementation supports YOLOv8
+detection checkpoints. YOLOv8 segmentation and MediaPipe models are not
+supported yet.
+
+## Convert a detector
+
+Ultralytics checkpoints must be converted before use. The converter fuses
+BatchNorm into convolution layers and writes a safetensors file with the weight
+names expected by the native GGML implementation.
+
+```bash
+python scripts/convert_yolov8_to_safetensors.py face_yolov8n.pt face_yolov8n.safetensors
+```
+
+The converter requires Python packages `ultralytics`, `torch`, and
+`safetensors`.
+Only YOLOv8 detection checkpoints are accepted.
+PyTorch checkpoints use pickle internally, so only convert `.pt` files from a
+trusted source.
+
+## Repair an existing image
+
+Use the dedicated `adetailer` mode to detect and repair objects in an existing
+image:
+
+```bash
+./bin/sd-cli \
+  -M adetailer \
+  -m model.safetensors \
+  -i input.png \
+  -o repaired.png \
+  -p "detailed portrait photo" \
+  --negative-prompt "deformed face" \
+  --steps 24 \
+  --cfg-scale 6 \
+  --strength 0.4 \
+  --sampling-method dpm++2m \
+  --scheduler karras \
+  --ad-model face_yolov8n.safetensors \
+  --extra-ad-args "confidence=0.3,inpaint_padding=32,mask_blur=4"
+```
+
+This mode reuses the normal image-generation options for the detail pass:
+
+- `--init-img`, `--output`, `--prompt`, and `--negative-prompt`
+- `--steps`, `--cfg-scale`, `--sampling-method`, and `--scheduler`
+- `--strength`, `--seed`, LoRA settings, VAE tiling, and backend assignments
+- `--width` and `--height`, which also resize the input when specified
+
+`--ad-prompt` and `--ad-negative-prompt` optionally override the normal prompts.
+Values provided in `--extra-ad-args`, such as `steps`, `cfg_scale`,
+`denoising_strength`, or `inpaint_width`, take precedence over inherited values.
+
+## Repair generated images
+
+ADetailer can also run automatically after normal image generation:
+
+```bash
+./bin/sd-cli \
+  -m model.safetensors \
+  -p "portrait photo" \
+  --ad-model face_yolov8n.safetensors \
+  --ad-prompt "[PROMPT], detailed face" \
+  --ad-negative-prompt "" \
+  --extra-ad-args "confidence=0.3,denoising_strength=0.4,inpaint_width=512,inpaint_height=512"
+```
+
+The same options can be supplied when starting `sd-server`. When `--ad-model` is set, its detector runs after OpenAI-compatible and native asynchronous image generation. SDAPI requests enable ADetailer through `alwayson_scripts`, as described below. If no object is detected, the original image is returned unchanged. RGBA output keeps its original alpha channel, including when detected regions are repaired. Detector initialization or detail-pass failures are reported as request failures.
+
+Server image-generation requests limit base, hires, and detail-pass dimensions to 4096 pixels per side. ADetailer step overrides are limited to 1–100; `steps=0` in `extra_ad_args` keeps inheriting the main generation steps. These limits also apply to inherited dimensions and `extra_ad_args` overrides. Requests exceeding the limits return a parameter error before generation. CLI limits are unchanged.
+
+### SDAPI requests
+
+`/sdapi/v1/txt2img` and `/sdapi/v1/img2img` accept the WebUI ADetailer argument format:
+
+```json
+{
+  "prompt": "portrait photo",
+  "alwayson_scripts": {
+    "ADetailer": {
+      "args": [true, false, {
+        "ad_model": "face_yolov8n.pt",
+        "ad_prompt": "[PROMPT], detailed face",
+        "ad_use_steps": true,
+        "ad_steps": 20
+      }]
+    }
+  }
+}
+```
+
+The detector must be configured at server startup with `--ad-model`. `ad_model` must match that file's basename; the original `.pt` basename is also accepted for a converted detector. Requests cannot load arbitrary detector paths. Omitting the script, setting its first argument to `false`, selecting `"None"`, or setting `ad_tab_enable` to `false` disables detailing. Dictionary-only argument arrays are also accepted. Multiple enabled dictionaries run in order using the configured detector.
+
+Supported settings include prompts, confidence, mask filtering by area or confidence, mask count and area ratios, dilation/erosion, offsets, mask merging/inversion, blur, denoising strength, masked or full-image inpainting, padding, and optional overrides for inpaint dimensions, steps, CFG, sampler/scheduler, and CLIP skip. The corresponding `ad_use_*` flags control whether overrides apply. Otherwise dimensions, steps, CFG, sampler, scheduler, and CLIP skip inherit the main request. Denoising strength defaults to `0.4`. Pixel mask settings are limited to the larger output dimension, with a minimum limit of 32 and a maximum of 16384 pixels.
+
+Model/VAE switching, non-default noise multipliers, face restoration, ControlNet, and skipping the initial img2img pass are not supported; active requests for these features return an error. Unknown settings and invalid active values also return errors. During detailing, `/sdapi/v1/progress` remains active with an `adetailer` job; its percentage and preview stay fixed until detailing finishes.
+
+An empty ADetailer prompt inherits the main prompt. `[PROMPT]` inserts the main
+prompt, `[SEP]` assigns different prompts to consecutive masks, and `[SKIP]`
+skips the corresponding mask.
+
+All settings other than the detector path and prompts are passed through
+`--extra-ad-args` as a comma-separated `key=value` list:
+
+| Key | Default | Description |
+| --- | ---: | --- |
+| `input_size` | `640` | Square YOLO input size; must be a multiple of 32 |
+| `confidence` | `0.3` | Detection confidence threshold |
+| `nms` | `0.45` | NMS IoU threshold |
+| `max_detections` | `100` | Maximum detections retained after NMS |
+| `mask_k_largest` | `0` | Keep only the largest K masks; zero keeps all |
+| `mask_filter_method` | `area` | Rank retained masks by `area` or `confidence` |
+| `mask_min_ratio` | `0` | Minimum bbox area relative to the image |
+| `mask_max_ratio` | `1` | Maximum bbox area relative to the image |
+| `dilate_erode` | `4` | Positive values dilate; negative values erode |
+| `x_offset`, `y_offset` | `0` | Mask offset in pixels; positive Y moves upward |
+| `mask_mode` | `none` | `none`, `merge`, or `merge_invert` |
+| `merge_masks`, `invert_mask` | `false` | Boolean alternatives to `mask_mode` |
+| `mask_blur` | `4` | Final composite feather radius |
+| `inpaint_padding` | `32` | Padding around the detected region |
+| `inpaint_only_masked` | `true` | Crop to the detected region; `false` uses the full image with the detection mask |
+| `inpaint_width`, `inpaint_height` | mode-specific | `512x512` after generation; input/output size in `adetailer` mode |
+| `denoising_strength` | mode-specific | `0.4` after generation; inherits `--strength` in `adetailer` mode |
+| `steps` | `0` | Detail steps; zero inherits the main generation |
+| `cfg_scale` | `-1` | Detail CFG; a negative value inherits the main generation |
+| `sample_method` | inherited | Detail sampler name |
+| `scheduler` | inherited | Detail scheduler name |
+| `sort_by` | `none` | `none`, `left_to_right`, `center_to_edge`, or `area` |
+
+Multiple masks are processed serially. Each completed inpaint becomes the input
+for the next mask, and the seed is incremented by the mask index. Use
+`mask_mode=merge` to process all detections in one inpaint pass.
+
+The detector uses the `detector` backend module. For example, keep detection on
+the CPU while diffusion runs on CUDA:
+
+```bash
+--backend "diffusion=cuda0,detector=cpu"
+```
@@ -15,7 +15,7 @@
 ## Examples

 ```sh
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\anima-preview.safetensors --vae ..\..\ComfyUI\models\vae\qwen_image_vae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_06b_base.safetensors  -p "a lovely cat holding a sign says 'anima.cpp'" --cfg-scale 6.0 --sampling-method euler -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\anima-preview.safetensors --vae ..\models\vae\qwen_image_vae.safetensors  --llm ..\models\text_encoders\qwen_3_06b_base.safetensors  -p "a lovely cat holding a sign says 'anima.cpp'" --cfg-scale 6.0 --sampling-method euler -v --offload-to-cpu
 ```

 <img alt="anima image example" src="../assets/anima/example.png" />
@@ -0,0 +1,171 @@
+# AnimateDiff (SD 1.5)
+
+AnimateDiff adds motion to a frozen Stable Diffusion 1.5 checkpoint by
+injecting a temporal-attention module at 20 UNet slots. The base SD 1.5
+model, VAE, and text encoder are unchanged; only the motion module produces
+the temporal residual that turns a batch of independent frames into a
+coherent animation. Reference: Guo et al., "AnimateDiff: Animate Your
+Personalized Text-to-Image Diffusion Models without Specific Tuning"
+(https://arxiv.org/abs/2307.04725).
+
+## Download weights
+
+- Motion module (v3, recommended)
+    - fp16 safetensors: https://huggingface.co/conrevo/AnimateDiff-A1111/resolve/main/motion_module/mm_sd15_v3.safetensors
+    - original checkpoint: https://huggingface.co/guoyww/animatediff/resolve/main/v3_sd15_mm.ckpt
+- SD 1.5 base model
+    - any SD 1.5 checkpoint works. `realisticVisionV60B1` and `toonyou_beta3`
+      are the ones used in guoyww's reference configs.
+- Domain Adapter LoRA (optional, v3 only, sharpens the base UNet's output
+  toward the motion module's trained distribution)
+    - ckpt: https://huggingface.co/guoyww/animatediff/resolve/main/v3_sd15_adapter.ckpt
+    - place under your `--lora-model-dir` and reference in the prompt as
+      `<lora:v3_sd15_adapter:1.0>`.
+
+The motion module is `~836 MB` and loads alongside the SD 1.5 UNet via
+`--motion-module`.
+
+## Motion module versions
+
+Per [animatediff.net/models](https://animatediff.net/models):
+
+| Module              | Base | Native res | Character |
+|---------------------|------|------------|-----------|
+| `mm_sd_v14.ckpt`    | 1.5  | 256x256    | earliest, more jittery |
+| `mm_sd_v15.ckpt`    | 1.5  | 256x256    | improved stability over v1.4 |
+| `mm_sd_v15_v2.ckpt` | 1.5  | 384x384    | significantly better motion dynamics |
+| `v3_sd15_mm.ckpt`   | 1.5  | 512x512    | smoothest, highest quality; pairs with a Domain Adapter LoRA |
+| `mm_sdxl_v10_beta`  | SDXL | 512x512    | experimental, not yet supported here |
+
+Match your `-H -W` to the module's native resolution for best results. v3 is
+trained at 512x512 - going smaller (e.g. 384x384) still works but the motion
+character is closer to v2.
+
+## Examples
+
+Generate an 8-frame animation at 512x512, seed 42, 20 steps. The sampler /
+scheduler / CFG values below match what mm_sd15_v3 was trained with; using
+SD 1.5 defaults (euler_a, low CFG) produces noise-like output.
+
+```
+.\bin\Release\sd-cli.exe -M vid_gen \
+    --model ..\models\checkpoints\realisticVisionV60B1.safetensors \
+    --motion-module ..\models\animatediff\mm_sd15_v3.safetensors \
+    --offload-to-cpu \
+    -p "a red apple on a wooden table" \
+    --cfg-scale 8.0 --sampling-method euler --scheduler discrete \
+    -H 512 -W 512 --video-frames 8 --fps 8 --steps 20 -s 42 \
+    -o out.avi
+```
+
+Generate at the motion module's native 16-frame context (recommended for
+best temporal quality). Needs more VRAM at 512x512, so drop to 384x384 or
+use layer streaming:
+
+```
+.\bin\Release\sd-cli.exe -M vid_gen \
+    --model ..\models\checkpoints\realisticVisionV60B1.safetensors \
+    --motion-module ..\models\animatediff\mm_sd15_v3.safetensors \
+    --offload-to-cpu \
+    -p "photo of coastline, rocks, storm weather, wind, waves, lightning" \
+    --cfg-scale 8.0 --sampling-method euler --scheduler discrete \
+    -H 384 -W 384 --video-frames 16 --fps 8 --steps 20 -s 42 \
+    -o out.avi
+```
+
+Low-VRAM streaming (verified with a 2 GiB cap on RTX 3060):
+
+```
+.\bin\Release\sd-cli.exe -M vid_gen \
+    --model ..\models\checkpoints\realisticVisionV60B1.safetensors \
+    --motion-module ..\models\animatediff\mm_sd15_v3.safetensors \
+    --max-vram 2.0 \
+    -p "photo of coastline, rocks, storm weather, wind, waves, lightning" \
+    --cfg-scale 8.0 --sampling-method euler --scheduler discrete \
+    -H 384 -W 384 --video-frames 8 --fps 8 --steps 20 -s 42 \
+    -o out.avi
+```
+
+## Reference-quality reproduction
+
+Using guoyww's official reference configs on this impl (RealisticVision v6.0
+base + `mm_sd15_v3` or `mm_sd_v15_v2` + native resolution + 16 frames + euler
+ 25 steps + CFG 8 + linear beta schedule) reproduces the reference
+AnimateDiff output style.
+
+### v3 (512x512, `mm_sd15_v3`)
+
+| Prompt                                | Sample |
+|---------------------------------------|--------|
+| B&W man on stormy coastline           | <img src="../assets/animatediff/v3_man.gif" width="256"/>    |
+| Close-up rabbit macro shot            | <img src="../assets/animatediff/v3_rabbit.gif" width="256"/> |
+| Coastline, storm, waves, lightning    | <img src="../assets/animatediff/v3_coast.gif" width="256"/>  |
+| Old house, storm, forest, night       | <img src="../assets/animatediff/v3_house.gif" width="256"/>  |
+
+### v2 (384x384, `mm_sd_v15_v2.ckpt`)
+
+| Prompt                                | Sample |
+|---------------------------------------|--------|
+| B&W man on stormy coastline           | <img src="../assets/animatediff/v2_man.gif" width="256"/>    |
+| Close-up rabbit macro shot            | <img src="../assets/animatediff/v2_rabbit.gif" width="256"/> |
+| Coastline, storm, waves, lightning    | <img src="../assets/animatediff/v2_coast.gif" width="256"/>  |
+| Old house, storm, forest, night       | <img src="../assets/animatediff/v2_house.gif" width="256"/>  |
+
+Motion is strong for scenes with motion cues in the prompt (storm/waves/wind)
+and subtle for static subjects (close-up macro), matching reference behavior.
+v2 has an additional motion module at the UNet middle block that v3 dropped;
+this impl auto-detects the topology from the checkpoint.
+
+### v3 + Domain Adapter LoRA
+
+Attaching the `v3_sd15_adapter` LoRA sharpens the base UNet output toward
+the training distribution the motion module was fine-tuned against. Same
+prompt, seed, config as above:
+
+<img src="../assets/animatediff/v3_rabbit_domain_lora.gif" width="256"/>
+
+Individual fur strands, glowing inner-ear, and richer forest detail become
+visible compared to the no-LoRA rendering.
+
+```
+sd-cli -M vid_gen --model realisticVisionV60B1.safetensors \
+       --motion-module mm_sd15_v3.safetensors \
+       --lora-model-dir ./loras \
+       -p "close up photo of a rabbit ...<lora:v3_sd15_adapter:1.0>" ...
+```
+
+## img2video
+
+Pass a pre-rendered image via `-i / --init-img` to animate FROM it. All N output frames start from the encoded init latent, then per-frame noise is added at `--strength`. Character identity, composition, and quality are anchored by the init image; the motion module adds subtle motion on top.
+
+Left: init image rendered with `-M img_gen`. Right: 8-frame vid_gen output.
+
+<img src="../assets/animatediff/img2video_demo.gif" width="512"/>
+
+```
+sd-cli -M img_gen ... -o init.png                  # any high-quality still
+sd-cli -M vid_gen --motion-module mm_sd15_v3.safetensors \
+    -i init.png --strength 0.75 \
+    --cfg-scale 7.0 --sampling-method euler --scheduler karras \
+    -H 512 -W 512 --video-frames 8 --steps 25 -s 42 \
+    -p "..." -o out.avi
+```
+
+`--strength` controls how far the motion module is allowed to deviate from the init image (higher = more motion, lower = more static).
+
+## Notes
+
+- The motion module was trained at `video_length=16`. Running with
+  `--video-frames 16` gives the best coherence; F=8 works but shows a shorter
+  motion arc. Frame counts up to 32 are supported by the positional encoding
+  but exceed the trained regime and produce more static output.
+- At `--video-frames 1` the motion module is skipped entirely and the output
+  is bit-identical to `-M img_gen`. This avoids the single-token
+  temporal-attention degeneracy that would otherwise emit an untrained-magnitude
+  residual on a single-frame sample.
+- The base UNet is frozen, so character identity and style follow the SD 1.5
+  checkpoint you pass to `--model`. LoRAs and prompt weighting attach to the
+  base model in the usual way.
+- No mid_block motion module in v3. `mm_sdxl_v10_beta` (SDXL variant) is not
+  supported yet.
+- Output is written as MJPEG AVI. Use `--fps` to set playback speed.
@@ -3,9 +3,10 @@
 `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.
+- `--params-backend` selects where model parameters are kept.

-If `--params-backend` is not set, parameters use the same backend as their module runtime backend.
+If `--params-backend` is not set, auto-fit chooses parameter placement. With
+`--auto-fit off`, parameters use the same backend as their module runtime backend.

 ## Syntax

@@ -29,6 +30,24 @@ The same syntax is used for parameter placement:
 sd-cli -m model.safetensors -p "a cat" --backend cuda0 --params-backend te=cpu,vae=cpu
 ```

+`--params-backend` also accepts the special value `disk`:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --backend cuda0 --params-backend disk
+```
+
+`--max-vram` can target resolved backend/device names:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --backend diffusion=cuda0,vae=vulkan0 --max-vram cuda0=6,vulkan0=2
+```
+
+The value is a shared per-device budget for managed weights and registered
+runner compute/cache buffers. Live free memory can lower the effective limit
+for each graph run. Driver contexts and allocations made outside the managed
+model runners are not part of this accounting, so it is not a hard physical
+VRAM cap.
+
 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:
@@ -37,6 +56,300 @@ Module names are case-insensitive. Hyphens and underscores in module names are i
 sd-cli -m model.safetensors -p "a cat" --backend all=cuda0,te=cpu
 ```

+## Multiple devices per module (layer split)
+
+A `--backend` module assignment can list several devices separated by `&`:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --backend "diffusion=cuda0&cuda1"
+```
+
+The module's transformer blocks are then distributed across the listed devices
+in contiguous ranges sized proportionally to each device's free memory (minus a
+compute-buffer headroom of about 2 GiB per device), and the
+module's graphs are executed with a `ggml_backend_sched` that runs each block
+on the device holding its weights, copying the residual stream at the range
+boundaries. The first device in the list is the module's main device: it also
+holds the non-block tensors (embeddings, final norms, small sub-runners such as
+CLIP models or projectors) and the graph inputs/outputs.
+
+Layer split is supported for the `diffusion` and `te` modules. For `te` it
+applies to the dominant text encoder (`t5xxl` or the LLM); other modules accept
+only a single device. If the module has no recognizable transformer blocks, the
+assignment falls back to the first listed device.
+
+`--params-backend` accepts no device lists. If the module has no explicit
+params assignment, each block range's parameters are loaded directly to (and,
+with `--params-backend diffusion=disk`, released directly from) its own device;
+an explicit assignment such as `te=cpu` keeps the parameters on that backend
+and stages each range to its device on demand.
+
+Layer split uses the fixed graph-cut plan to assign blocks across devices, but
+single-device segmented execution and next-segment prefetch are disabled for
+the split module. `--max-vram` can still provide the per-device limits used by
+layer split and auto-fit.
+
+Use `--list-devices` to see the device names available on the system.
+
+### Row split (`--split-mode row`)
+
+`--split-mode` selects how a multi-device module distributes its weights:
+`layer` (the default, described above) or `row`. It accepts a single mode or
+per-module assignments:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --backend "diffusion=cuda0&cuda1" --split-mode row
+sd-cli -m model.safetensors -p "a cat" --backend "diffusion=cuda0&cuda1,te=cuda0&cuda1" --split-mode diffusion=row,te=layer
+```
+
+In row mode the module keeps executing on its main (first listed) device, but
+its transformer-block matmul weights are allocated in the backend's row-split
+buffer type, which slices each weight's rows across the listed devices in
+proportion to free memory and runs those matmuls on all devices in parallel.
+Compared to a layer split this uses all GPUs within every layer (instead of
+sequentially device by device) at the cost of a cross-device reduction per
+matmul - usually the faster option when the devices have fast interconnect.
+
+Row split requires a compatible split-buffer export from the linked GGML
+backend. If it is unavailable (or the listed devices belong to different backend
+registries), the module falls back to a layer split.
+Embeddings, normalization weights, biases and other non-block tensors stay in
+regular buffers on the main device.
+
+Row-split execution can use graph segments, but split weights are loaded
+synchronously instead of using the normal single-device prefetch path. Because
+GGML does not expose exact shard allocation sizes, the managed budget currently
+counts a split buffer's full size on each participating device. This is a
+conservative bound and can reject otherwise feasible layouts.
+
+Direct ("immediately") LoRA application cannot patch row-split tensors; with
+`--split-mode row` the automatic LoRA mode selects runtime application, and an
+explicit `--lora-apply-mode immediately` skips the split tensors with a
+warning.
+
+## Automatic placement (`--auto-fit on|off`)
+
+`--auto-fit` requires `on` or `off` and defaults to `on` when omitted.
+Explicit `--params-backend` assignments disable auto-fit,
+regardless of argument order, even with `--auto-fit on`.
+
+Auto-fit preserves explicit `--backend` assignments, including per-module
+assignments and device lists. For modules without a runtime assignment, it chooses
+the discrete GPU with the largest positive memory budget (the first device on a
+tie), falling back to integrated GPUs when no discrete GPU has a positive budget.
+An integrated GPU's shared system memory does not give it priority over a discrete
+GPU. Explicit assignments such as `--backend Vulkan1` still select that device.
+It then derives parameter placements from the model metadata, each module's
+compute devices, and the remaining memory budgets. The chosen backend
+specifications are printed.
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --auto-fit on
+sd-cli -m model.safetensors -p "a cat" --auto-fit on --max-vram cuda0=8,cuda1=14
+sd-cli -m model.safetensors -p "a cat" --backend cuda0
+sd-cli -m model.safetensors -p "a cat" --backend diffusion=cuda0,te=cpu,vae=cuda1
+sd-cli -m model.safetensors -p "a cat" --auto-fit off
+```
+
+Budgets reuse `--max-vram`: a positive per-device value caps what auto-fit
+plans with on that device, a negative value means "free memory minus that many
+GiB", and with no budget set each device's free memory minus a 512 MiB margin
+is used. These resolved GPU budgets, including the safety margin, also drive
+the runner's graph-cut capacity checks.
+
+Runtime capacity checks also leave 512 MiB of currently free device memory for
+backend scratch buffers and pipelines, including with explicit backend assignments.
+They cap free-memory reports by the device's total memory minus tracked
+resident allocations. Vulkan reports exceeding total memory are rejected on
+discrete GPUs because its heap-budget subtraction can underflow. On integrated
+GPUs, a wrapped report triggers reclamation before the tracked estimate can be
+accepted, subject to the byte checks and configured budgets. Other backends use
+the cap instead of treating such reports as zero free memory. Failed checks log
+the reported free and total memory alongside tracked weight and runtime allocations.
+With `--mmap`, device-backed mappings count toward these budgets at their full
+mapped-file size, once per device buffer even when multiple parameter blocks
+share it. Mappings retained in the loader cache continue to count.
+
+Components are considered in `diffusion`, `te`, `vae` order so that repeatedly
+used diffusion weights have priority. Each component's weights use the first
+storage location with enough remaining budget:
+
+1. The component's compute GPU, leaving estimated space for computation and weight staging.
+2. CPU RAM, reserving the larger of 2 GiB or 10% of available RAM for other work.
+3. Another GPU, choosing the one with the largest remaining budget that fits.
+4. Disk, reloading weights on demand.
+
+GPU cache space follows the same component priority. Before a lower-priority
+component can become permanently resident, the planner leaves room for the full
+weights and estimated compute space of higher-priority offloaded components.
+If offloaded diffusion already needs the entire main GPU budget, TE and VAE also
+use offloaded parameters. GPU copies remain cached across sampling steps while
+a component's execution is active and are released under memory pressure.
+Ending a component's execution releases its staged GPU copies, disk-backed
+weight buffers, temporary compute buffers and execution caches. Parameters
+assigned permanently to a backend remain resident. Hires model upscaling shares
+the runtime budget and can
+reclaim idle weight copies from other components.
+Reclamation releases idle compute buffers first, then cached weight allocations
+until the current work fits. Weights used by active computation or pending
+transfers stay protected.
+For CPU computation with disk-backed parameters, cached weights are released
+when execution ends so they do not accumulate in RAM across components.
+CPU parameter residency allows GPU weight caching; it does not force every
+weight to be copied again at every step.
+
+RAM and GPU budgets are shared across components. Each component uses a single
+parameter backend; several other GPUs' capacities are not combined to store
+one component. If available RAM cannot be queried, RAM residency is skipped.
+Weights stored on another GPU are copied to the component's compute devices for
+execution. CPU modules use RAM or disk. Compute reserves and cache priority are
+accounted for separately on each device, so a CPU module does not reserve GPU
+space. Storage on another module's GPU also leaves room for that module's work.
+
+Auto-fit does not select multi-GPU layer/row computation itself. Explicit device
+lists and `--split-mode` still control that computation. Before the runners have
+built their split plans, auto-fit conservatively counts the full component size
+on each listed GPU when checking residency and cache space. This can offload
+parameters even when a split layout would fit; use `--auto-fit off` to keep the
+default split-device parameter placement.
+
+For example, a diffusion model whose full weights exceed the main GPU's budget
+can use `--backend diffusion=cuda0 --params-backend diffusion=cpu` when RAM is
+sufficient. Automatic graph segmentation can then load the required weights
+for each segment and reclaim idle GPU copies. `--disable-segmented-compute`
+still disables segmentation.
+
+Initial compute reserves are estimates (2 GiB for diffusion and text encoders,
+1 GiB for VAE); higher-priority placements also leave staging space for the
+largest weight tensor of each lower-priority offloaded component. Actual segment
+weights, compute buffers and caches must
+still fit the runner's capacity checks. Offloading weights does not guarantee
+that every resolution or frame count will fit, and auto-fit does not change a
+component to CPU computation solely because its full weights exceed VRAM.
+If a VAE decode fails, decoding retries with spatial tiling even when `--auto-fit`
+is off; supported video decoders try temporal tiling first and can then add
+spatial tiling. Spatial retries use half-size tiles along each latent dimension.
+
+## Direct convolution (`--conv-direct on|off|auto`)
+
+`--conv-direct` defaults to `auto` and accepts module assignments for
+`diffusion`, `vae` (including preview VAE), and `controlnet`, using the same
+aliases and precedence rules as `--flash-attn`:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --conv-direct auto,vae=off
+sd-cli -m model.safetensors -p "a cat" --conv-direct off,diffusion=on
+```
+
+- `on` requests direct 2D and 3D convolution.
+- `off` uses the existing im2col-based convolution paths.
+- `auto` first checks support for a representative F16-weight/F32-input direct
+  convolution graph. Vulkan then benchmarks direct convolution against im2col
+  using two small shapes per dimension, with warmup, output validation, and five
+  interleaved measurements per shape. Direct convolution must improve the sum of
+  median times by at least 5%; a failed probe defaults to `off`.
+  Other backends retain their implicit GEMM checks without benchmarking:
+  CUDA 2D requires a compiled Turing-or-newer MMA architecture; HIP 2D requires
+  the supported CDNA MFMA or RDNA WMMA architectures described below for flash
+  attention. MUSA 2D has a SIMT implementation. CUDA, HIP, and MUSA 3D also have
+  a SIMT implementation, so they do not require MMA. Other backends default to
+  `off`. Results and decision logs are cached per compute backend and dimension
+  in the context, as with flash attention. Each dimension is selected independently,
+  and all compute devices in a module must qualify. Parameter-storage devices do
+  not affect this choice.
+
+These are backend-level defaults. GGML still chooses the actual kernel based on
+weight type, shape, and hardware; enabling direct convolution does not guarantee
+that every operation uses implicit GEMM. The CUDA-family implicit GEMM paths
+require F16 weights and can use ordinary direct kernels for other cases.
+For 3D convolution, a supported direct kernel is also used as a fallback when
+the im2col decomposition is unsupported for the actual tensor sizes, including
+when an intermediate reshape or matrix multiplication exceeds backend limits.
+This option does not change the separately created image upscaler context.
+
+Repeating `--conv-direct` replaces the previous assignment string. Existing
+`--diffusion-conv-direct` (diffusion and ControlNet) and `--vae-conv-direct`
+flags are deprecated and emit a warning, but remain supported when
+`--conv-direct` is omitted. Use `--conv-direct auto,diffusion=on,controlnet=on`
+and `--conv-direct auto,vae=on`, respectively. An explicit
+`--conv-direct` takes precedence over both legacy flags regardless of order.
+Library callers can set the appended `sd_ctx_params_t::conv_direct_mode` field
+to the same string; null or empty selects `auto` while honoring legacy true
+flags. The context copies the string.
+
+## Flash attention (`--flash-attn on|off|auto`)
+
+`--flash-attn` defaults to `auto`. Selection happens while creating the model
+context, before inference, and is cached once per compute backend in that context.
+Parameter-storage backends are not benchmarked.
+
+The option also accepts module assignments, using the same module aliases as
+`--backend`:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --flash-attn auto,te=off,vae=on
+sd-cli -m model.safetensors -p "a cat" --flash-attn off,diffusion=auto
+```
+
+Supported modules are `diffusion`, `te`, `clip_vision`, `vae`, `controlnet`, and
+`audio_encoder`. Modules without an assignment use the default (`auto` if no
+default is supplied). A bare mode or `all=`, `default=`, or `*=` sets the default;
+module-specific settings take precedence regardless of their position. Within
+one assignment string, the last setting for the same module or default wins.
+Repeating `--flash-attn` replaces the entire previous string.
+
+Only modules set to `auto` need backend probing. An explicit `on` or `off` for
+one module does not change another module's decision, even when they share a
+backend. Modules without configurable attention are rejected instead of silently
+ignoring their settings.
+
+- `on` requests flash attention for all models. Unsupported attention operations
+  still fall back to ordinary attention.
+- `off` disables flash attention, including the flash-attention fallback from
+  SageAttention. SageAttention itself remains controlled by `--sage-attn`.
+- `auto` enables CPU and supported CUDA MMA configurations without timing a
+  benchmark. CUDA uses both the device's capability and the compiled GGML
+  architectures. HIP also skips timing on known CDNA MFMA architectures
+  (`gfx908`, `gfx90a`, `gfx940`/`gfx941`/`gfx942`, `gfx950`) and RDNA 3/3.5/4
+  WMMA architectures (`gfx11xx`/`gfx12xx`). The CDNA shortcut requires
+  `GGML_HIP_MMQ_MFMA` to be enabled at build time. All shortcuts still check
+  support for the FA probe graph. Vulkan devices whose description identifies
+  NVIDIA RTX hardware (GeForce RTX, TITAN RTX, Quadro RTX and professional RTX)
+  also skip timing. This is a device-name heuristic and needs no CUDA runtime.
+  Unsupported graphs are disabled; other configurations, including older or
+  unknown HIP architectures and other Vulkan devices, are benchmarked.
+
+The probe compares ordinary and flash attention through the same attention
+builder used by models, including layout conversion and F16 conversion. It uses
+two small self-attention graphs (512 tokens with head size 64, and 1024 tokens
+with head size 128, both with four heads), checks their outputs, warms up each
+path, then alternates five synchronized measurements. Flash attention is enabled
+when the sum of its median times is at least 5% lower. These representative
+shapes are a startup heuristic, not a guarantee for every model or resolution.
+Allocation/execution failures and inconclusive probes select `off`. Timing stops
+starting new measurement pairs after one second of measured computation;
+backend initialization, warmup and an in-flight pair can take longer.
+
+Different modules can select different results. A multi-device model enables
+flash attention only if all of its compute devices select `on`, preserving the
+model's existing single boolean setting. This conservative rule also covers
+row-split configurations, which can fall back to layer split during initialization.
+
+`--fa` and `--diffusion-fa` are deprecated and emit a warning when used.
+Use `--flash-attn on` and `--flash-attn off,diffusion=on`, respectively.
+Both aliases write the same mode string as `--flash-attn`; the last occurrence
+of any of these options wins. Library callers
+can set `sd_ctx_params_t::flash_attn_mode` to the same string, including
+`"on"`, `"off"`, `"auto"`, or `"auto,te=off,vae=on"`.
+A null or empty string uses automatic selection unless a
+legacy `flash_attn` or `diffusion_flash_attn` boolean is set. The context copies
+the string; no new enum is required.
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --flash-attn auto
+sd-cli -m model.safetensors -p "a cat" --flash-attn off
+```
+
 ## Modules

 | Module | Purpose | Accepted names |
@@ -48,6 +361,7 @@ sd-cli -m model.safetensors -p "a cat" --backend all=cuda0,te=cpu
 | `controlnet` | ControlNet | `controlnet`, `control` |
 | `photomaker` | PhotoMaker ID encoder and PhotoMaker LoRA | `photomaker`, `photomakerid`, `pmid`, `photo` |
 | `upscaler` | ESRGAN upscaler | `upscaler`, `esrgan`, `hires` |
+| `detector` | ADetailer YOLOv8 detector | `detector`, `adetailer`, `yolo` |

 `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.

@@ -64,9 +378,11 @@ The special values `auto`, `default`, and an empty backend name select the defau

 The special value `gpu` selects the first GPU backend, falling back to the first integrated GPU backend.

+The special value `disk` is accepted only by `--params-backend`. `--backend disk` is invalid because `disk` is a parameter residency mode, not a runtime compute backend.
+
 ## Runtime backend vs. parameter backend

-The runtime backend controls where graph execution runs. The parameter backend controls where model weights are allocated.
+The runtime backend controls where graph execution runs. The parameter backend controls where model weights are allocated or whether they are reloaded from disk on demand.

 For example:

@@ -76,6 +392,16 @@ 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.

+For example:
+
+```shell
+sd-cli -m model.safetensors -p "a cat" --backend cuda0 --params-backend disk
+```
+
+This runs all modules on `cuda0`, reloads parameters from the model file as needed, and releases those parameter buffers after use.
+
+Outside `--auto-fit`, `disk` is never selected implicitly. If `--params-backend` is not set, parameters use the runtime backend.
+
 Per-module assignments can be mixed:

 ```shell
@@ -100,23 +426,31 @@ 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.

+`--params-backend disk` does not create a separate backend instance. Parameters are loaded lazily using the module runtime backend.
+
 `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:
+The example CLI/server still accepts these older CPU placement flags as compatibility aliases:

 - `--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`.
+`--clip-on-cpu`, `--vae-on-cpu`, and `--control-net-cpu` are deprecated. The example argument layer prepends `te=cpu`, `vae=cpu`, and `controlnet=cpu` to `--backend` before creating the context.

-`--offload-to-cpu` affects parameter backend assignment only when `--params-backend` is not set. It is equivalent to:
+`--offload-to-cpu` prepends a CPU default to the parameter assignment in the caller before creating the context:

 ```shell
--params-backend cpu
+--params-backend '*=cpu'
 ```

-Explicit `--backend` and `--params-backend` assignments are preferred for new commands.
+Because this default is inserted first, later explicit `--params-backend` entries can still override it, for example `--offload-to-cpu --params-backend te=disk` keeps non-TE parameters on CPU and reloads TE parameters from disk.
+
+Library callers should set `backend` and `params_backend` directly. `sd_ctx_params_init()`
+enables `auto_fit` by default; a nonempty `params_backend` assignment disables it.
+The `backend` assignment constrains auto-fit's compute placement.
+The old CPU/offload fields are no longer part of the C API. Explicit `--backend` and
+`--params-backend` assignments are preferred for new commands.
@@ -0,0 +1,50 @@
+# Adding backend performance probes
+
+`src/core/probe/backend_probe.h` provides the shared measurement machinery. Operator
+policy stays separate: `src/core/probe/flash_attn.cpp`, for example, defines attention
+shapes and its CPU/CUDA/HIP/Vulkan shortcuts, then calls `sd_run_backend_probe` for unknown
+configurations. `src/core/probe/conv_direct.cpp` uses metadata-only probe graphs
+to check direct 2D/3D convolution support, then enables known implicit GEMM
+paths on non-Vulkan backends. Vulkan compares direct convolution with im2col
+using the shared measurement machinery.
+
+For each representative shape, create a `BackendProbeCase`:
+
+- `build(graph, candidate)` creates inputs and either the reference or candidate
+  graph in `graph.ctx`, registers input tensors in `graph.inputs`, and assigns
+  `graph.output`. Leave the output null if the requested implementation is
+  unavailable; do not silently benchmark its fallback as the candidate.
+- By default, inputs receive deterministic F32 data, with the same seed for both
+  variants. Supply `initialize` for F16 convolution weights, integer inputs, or
+  operator-specific values. It runs after allocation, outside timed execution.
+  Both variants must receive equivalent logical inputs.
+- Default output validation requires matching shapes and contiguous F32 outputs,
+  checks finiteness, and applies absolute and relative tolerances. Supply
+  `compare` for other output types or validation rules.
+
+Graphs must be safe to execute repeatedly without modifying their inputs.
+`BackendProbeGraph` owns its GGML context and buffer, so all return paths release
+probe resources. Constructors allocate metadata only; `supported()` can also be
+used before a known-architecture shortcut, without allocating device tensors.
+
+`BackendProbeOptions` controls graph metadata capacity, warmup count, measurement
+count, measured-time budget, minimum free device memory, comparison tolerances,
+and minimum improvement. Warmup and allocation are excluded from timing. The
+runner alternates the two variants, synchronizes the backend around execution,
+and compares sums of per-case medians. The time budget is checked before each
+measurement pair; it cannot interrupt an in-flight backend operation.
+
+Only use `candidate_faster` when `completed` is true. On incomplete probes,
+`error` reports why comparison failed and the timing totals may be partial.
+The caller handles logging, known architecture rules, fallback policy, and
+context-lifetime caching. Keep caches separate by operation (or key a shared
+cache by both operation and backend); an attention result must not select a
+convolution implementation. If shapes or settings vary within one context,
+include them in the corresponding cache key as well.
+
+Convolution results are cached separately for 2D and 3D, once per compute backend
+in the context, just like attention results. Each probe logs its backend, decision,
+and reason; mixed-device modules log when their compute devices disagree.
+
+Module assignment parsing remains in `src/core/ggml_extend_backend.cpp`, next
+to backend assignment parsing. The probe directory does not parse CLI/API options.
@@ -0,0 +1,31 @@
+# How to Use
+
+Boogu Image uses a Boogu diffusion transformer, the FLUX VAE, and Qwen3-VL as the LLM text and vision encoder.
+
+## Download weights
+
+- Download Boogu Image
+    - safetensors: https://huggingface.co/Comfy-Org/Boogu-Image/tree/main/diffusion_models
+- Download vae
+    - safetensors: https://huggingface.co/black-forest-labs/FLUX.1-dev/blob/main/ae.safetensors
+- Download Qwen3-VL 8B
+    - gguf: https://huggingface.co/unsloth/Qwen3-VL-8B-Instruct-GGUF/tree/main
+        - For image editing with GGUF text encoders, also download the matching mmproj file and pass it with `--llm_vision`.
+
+## Examples
+
+### Boogu Image Base
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\boogu_image_base_bf16.safetensors --llm ..\models\text_encoders\Qwen3VL-8B-Instruct-Q4_K_M.gguf --vae ..\models\vae\ae.sft -p "a lovely cat" -v --offload-to-cpu
+```
+
+<img width="256" alt="Boogu Image Base example" src="../assets/boogu/example.png" />
+
+### Boogu Image Edit
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\boogu_image_edit_bf16.safetensors --llm ..\models\text_encoders\Qwen3VL-8B-Instruct-Q4_K_M.gguf --llm_vision ..\models\text_encoders\mmproj-Qwen3VL-8B-Instruct-F16.gguf --vae ..\models\vae\ae.sft -v --offload-to-cpu -r ..\assets\flux\flux1-dev-q8_0.png -p "change 'flux.cpp' to 'boogu.cpp'"
+```
+
+<img width="256" alt="Boogu Image Edit example" src="../assets/boogu/edit_example.png" />
@@ -2,20 +2,82 @@

 ## Get the Code

-```
+```shell
 git clone --recursive https://github.com/leejet/stable-diffusion.cpp
 cd stable-diffusion.cpp
 ```

- If you have already cloned the repository, you can use the following command to update the repository to the latest code.
+If you already cloned the repository without submodules, initialize them from the repository root:

+```shell
+git submodule update --init --recursive
 ```
+
+To update an existing checkout and its submodules:
+
+```shell
 cd stable-diffusion.cpp
 git pull origin master
-git submodule init
-git submodule update
+git submodule update --init --recursive
 ```

+## MSVC compilation parallelism
+
+MSVC builds limit parallel C/C++ compilation within each project by default.
+For `n` logical processors, the limit is `n` when `n <= 2`, `n - 1` when
+`3 <= n <= 8`, and `n - 2` when `n > 8`. The default is calculated when the
+build directory is first configured with this setting. Override it with
+`-DSD_MSVC_MP_COUNT=4`, or pass `-USD_MSVC_MP_COUNT` to recalculate it.
+
+For Visual Studio builds, keep project-level parallelism at one to avoid
+multiplying the per-project limit:
+
+```shell
+cmake --build build --config Release --parallel 1
+```
+
+This setting controls MSVC's `/MP` option, not CUDA compiler parallelism or a
+hard CPU usage limit. Other build tools, such as Ninja, schedule compilation
+jobs separately; use `cmake --build build --parallel N` to limit those jobs.
+
+## Selecting a GGML source tree
+
+By default, sd.cpp builds the patched GGML submodule in `ggml/`. To build with
+an upstream GGML checkout instead, enable `SD_USE_UPSTREAM_GGML` and set
+`SD_GGML_SOURCE_DIR`:
+
+```shell
+cmake -S . -B build-upstream -DSD_USE_UPSTREAM_GGML=ON -DSD_GGML_SOURCE_DIR=../ggml-upstream
+cmake --build build-upstream --config Release
+```
+
+Upstream compatibility mode uses the selected library and matching private
+headers. The patched GGML build uses public headers, including `ggml-ext.h` and
+`ggml-graph.h`. Backend options such as `-DSD_CUDA=ON` apply to the selected tree
+as usual.
+
+`SD_USE_UPSTREAM_GGML` defaults to `OFF`, which enables the patched GGML
+extensions. Set it to `ON` when using upstream GGML; it selects the compatibility
+mode and does not download or replace the GGML source tree. Upstream mode keeps
+the original FP8 safetensors handling: FP8 tensors are converted to F16 at load
+time (one byte per element in the file, two in RAM and VRAM). INT8
+tensorwise/convrot is disabled and its model files are rejected with an explicit
+error. FP8 GGUF files, FP8 weight type requests and tensor type rules are also
+rejected; no automatic conversion is performed.
+
+Upstream GGML may lack some operators and performance optimizations provided by
+the patched version. A warning is emitted during CMake configuration and when
+creating an inference context. Ordinary floating-point and shared GGML
+quantization types remain available, subject to backend operator support.
+
+`SD_USE_SYSTEM_GGML=ON` instead links an installed GGML CMake package, located
+with `ggml_DIR` or `CMAKE_PREFIX_PATH`. An installed patched package must include
+the extension and graph APIs; its source tree is not required for private
+headers. The installed library must use the same ABI settings as sd.cpp,
+including `GGML_MAX_NAME`. For an upstream package, also set
+`SD_USE_UPSTREAM_GGML=ON` and point `SD_GGML_SOURCE_DIR` to its matching source
+tree for the compatibility layer's private headers.
+
 ## WebP and WebM Support in Examples

 The example applications (`examples/cli` and `examples/server`) use `libwebp` to support WebP image I/O, and `examples/cli` can also use `libwebm` for `.webm` video output. Both are enabled by default. WebM output currently reuses `libwebp` to encode each frame as VP8 before muxing with `libwebm`.
@@ -56,6 +118,10 @@ cmake --build . --config Release

 ## Build with CUDA

+The sd.cpp SageAttention extension is included when using CUDA with patched
+GGML (`SD_USE_UPSTREAM_GGML=OFF`).
+See [SageAttention](sage_attention.md) for GPU requirements and `--sage-attn` usage.
+
 This provides GPU acceleration using NVIDIA GPU. Make sure to have the CUDA toolkit installed. You can download it from your Linux distro's package manager (e.g. `apt install nvidia-cuda-toolkit`) or from here: [CUDA Toolkit](https://developer.nvidia.com/cuda-downloads). Recommended to have at least 4 GB of VRAM.

 ```shell
@@ -64,6 +130,28 @@ cmake .. -DSD_CUDA=ON
 cmake --build . --config Release
 ```

+CUDA extensions use sd.cpp's `thirdparty/cutlass` submodule by default
+for eligible INT8 GEMM and CuTe attention operations. If its headers are missing,
+CMake stops with an initialization hint; it does not download them automatically.
+Run `git submodule update --init --recursive` from the repository root before
+configuring.
+
+To build with the existing CUDA implementations without CUTLASS, configure from
+the repository root with:
+
+```shell
+cmake -S . -B build -DSD_CUDA=ON -DSD_CUDA_CUTLASS_PATH:PATH=
+cmake --build build --config Release
+```
+
+Alternatively, set `SD_CUDA_CUTLASS_PATH` to another local CUTLASS source
+directory (tested with 3.9.2). Existing CMake caches retain their configured path;
+pass `-USD_CUDA_CUTLASS_PATH` when reconfiguring to restore the bundled default.
+Existing `GGML_CUDA_CUTLASS_PATH` caches are migrated on first configuration:
+the old bundled path becomes `thirdparty/cutlass`, while custom paths and an
+empty disabled value are preserved. An explicitly set `SD_CUDA_CUTLASS_PATH`
+takes precedence. Non-CUDA builds do not require CUTLASS.
+
 ## Build with HipBLAS

 This provides GPU acceleration using AMD GPU. Make sure to have the ROCm toolkit installed.
@@ -81,14 +169,16 @@ cmake --build . --config Release

 ## Build with MUSA

-This provides GPU acceleration using Moore Threads GPU. Make sure to have the MUSA toolkit installed.
+This provides GPU acceleration using Moore Threads GPUs. The configuration below follows upstream validation with MUSA SDK 5.2.0 and MTT S5000 (PH1, `mp_31`). Install that toolkit before building. Older SDKs such as 4.2 may lack the CUB headers required by the current backend.

 ```shell
 mkdir build && cd build
-cmake .. -DCMAKE_C_COMPILER=/usr/local/musa/bin/clang -DCMAKE_CXX_COMPILER=/usr/local/musa/bin/clang++ -DSD_MUSA=ON -DCMAKE_BUILD_TYPE=Release
+cmake .. -DCMAKE_C_COMPILER=/usr/local/musa/bin/clang -DCMAKE_CXX_COMPILER=/usr/local/musa/bin/clang++ -DSD_MUSA=ON -DMUSA_ARCHITECTURES=31 -DCMAKE_BUILD_TYPE=Release
 cmake --build . --config Release
 ```

+The MUSA Docker image uses the same SDK 5.2.0 development image as upstream for both build and runtime stages, and targets `mp_31` only. It includes the development tools and PyTorch from that base image, so it is larger than a minimal runtime image. Other GPU architectures require separate validation.
+
 ## Build with Metal

 Using Metal makes the computation run on the GPU. Currently, there are some issues with Metal when performing operations on very large matrices, making it highly inefficient at the moment. Performance improvements are expected in the near future.
@@ -2,6 +2,16 @@

 Caching methods accelerate diffusion inference by reusing intermediate computations when changes between steps are small.

+### Conditioning Cache
+
+Conditioning results are cached per model context using an LRU cache. The default
+capacity is **0 (disabled) for `sd-cli`** and **4 entries for `sd-server` and the C
+API**. Set `--conditioning-cache-size N` to change the limit; `0` disables caching.
+For example, `sd-cli -m model.safetensors -p "a cat" --conditioning-cache-size 4`
+enables the cache in the CLI. The C API option is
+`sd_ctx_params_t::conditioning_cache_size`, initialized by `sd_ctx_params_init()`.
+This cache is independent of the diffusion-step `--cache-mode` options below.
+
 ### Cache Modes

 | Mode | Target | Description |
@@ -15,7 +15,7 @@ You can run Chroma using stable-diffusion.cpp with a GPU that has 6GB or even 4G
 You can download the preconverted gguf weights from [silveroxides/Chroma-GGUF](https://huggingface.co/silveroxides/Chroma-GGUF), this way you don't have to do the conversion yourself.

 ```
-.\bin\Release\sd-cli.exe -M convert -m ..\..\ComfyUI\models\unet\chroma-unlocked-v40.safetensors -o ..\models\chroma-unlocked-v40-q8_0.gguf -v --type q8_0
+.\bin\Release\sd-cli.exe -M convert -m ..\models\unet\chroma-unlocked-v40.safetensors -o ..\models\chroma-unlocked-v40-q8_0.gguf -v --type q8_0
 ```

 ## Run
@@ -24,10 +24,9 @@ You can download the preconverted gguf weights from [silveroxides/Chroma-GGUF](h
 For example:

 ```
- .\bin\Release\sd-cli.exe --diffusion-model  ..\models\chroma-unlocked-v40-q8_0.gguf --vae ..\models\ae.sft --t5xxl ..\models\t5xxl_fp16.safetensors  -p "a lovely cat holding a sign says 'chroma.cpp'" --cfg-scale 4.0 --sampling-method euler -v --chroma-disable-dit-mask --clip-on-cpu
+ .\bin\Release\sd-cli.exe --diffusion-model  ..\models\chroma-unlocked-v40-q8_0.gguf --vae ..\models\ae.sft --t5xxl ..\models\t5xxl_fp16.safetensors  -p "a lovely cat holding a sign says 'chroma.cpp'" --cfg-scale 4.0 --sampling-method euler -v --model-args chroma_use_dit_mask=false --clip-on-cpu
 ```

 ![](../assets/flux/chroma_v40.png)


-
@@ -12,7 +12,7 @@
 ## Examples

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\Chroma1-Radiance-v0.4-Q8_0.gguf --t5xxl ..\..\ComfyUI\models\clip\t5xxl_fp16.safetensors  -p "a lovely cat holding a sign says 'chroma  radiance cpp'" --cfg-scale 4.0 --sampling-method euler -v
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\Chroma1-Radiance-v0.4-Q8_0.gguf --t5xxl ..\models\clip\t5xxl_fp16.safetensors  -p "a lovely cat holding a sign says 'chroma  radiance cpp'" --cfg-scale 4.0 --sampling-method euler -v
 ```

 <img alt="Chroma1-Radiance" src="../assets/flux/chroma1-radiance.png" />
@@ -18,16 +18,43 @@ docker run --rm --init -v /path/to/models:/models -v /path/to/output/:/output -p

 ## Building using Docker

+Initialize the submodules from the repository root before building. Docker copies
+the local checkout, so its build context must include the submodule contents:
+
 ```shell
-docker build -t sd .
+git submodule update --init --recursive
+docker build -f docker/Dockerfile -t sd .
 ```

 ## Building variants using Docker

-Vulkan:
+### CUDA (NVIDIA GPU):

 ```shell
-docker build -f Dockerfile.vulkan -t sd .
+docker build -f docker/Dockerfile.cuda -t sd-cuda .
+```
+
+To select GPU architectures, pass `--build-arg "CUDA_ARCHITECTURES=89"` (for
+example, RTX 4090). Run the image on a host configured with NVIDIA Container
+Toolkit, using `docker run --rm --gpus all ... sd-cuda [args...]`.
+
+### Vulkan:
+
+```shell
+docker build -f docker/Dockerfile.vulkan -t sd .
+```
+
+### ROCm (AMD GPU):
+
+Defining one or multiple targets with `--build-arg AMDGPU_TARGETS` is strongly recommended.
+Building the image will otherwise likely fail or build for a target that does not fit your
+hardware.
+
+```shell
+# Build for one target
+docker build -f docker/Dockerfile.rocm --build-arg "AMDGPU_TARGETS=gfx1100" -t sd .
+# Build for multiple targets
+docker build -f docker/Dockerfile.rocm --build-arg "AMDGPU_TARGETS=gfx1030;gfx1100" -t sd .
 ```

 ## Run locally built image's CLI
@@ -37,3 +64,21 @@ docker run --rm -v /path/to/models:/models -v /path/to/output/:/output sd [args.
 # For example
 # docker run --rm -v ./models:/models -v ./build:/output sd -m /models/sd-v1-4.ckpt -p "a lovely cat" -v -o /output/output.png
 ```
+
+### ROCm (AMD GPU)
+
+The user that executes this command needs to be in the `video` group and have read and write
+access to `/dev/dri` and `/dev/kfd`.
+
+`--device` and `--group-add` allow the container access to the GPU.
+
+```shell
+docker run --rm --device=/dev/dri --device=/dev/kfd --group-add video \
+  -v /path/to/models:/models -v /path/to/output/:/output \
+  sd [args...]
+# For example
+# docker run --rm --device=/dev/dri --device=/dev/kfd --group-add video \
+#   -v ./models:/models -v ./build:/output \
+#   sd -m /models/sd-v1-4.ckpt -v -o /output/output.png \
+#   -p "a lovely cat"
+```
@@ -0,0 +1,114 @@
+# Image Editing
+
+Image editing in `stable-diffusion.cpp` allows you to use reference images to guide the generation process, enabling tasks like identity preservation, style transfer, or layout modification.
+
+
+## Supported Models
+
+Depending on the architecture, different models handle reference images differently.
+
+| Model | Default Preset |
+| :--- | :--- |
+| [**FLUX.1-Kontext-dev**](./kontext.md) | `flux_kontext` |
+| [**LongCat Image Edit**](./longcat_image.md) | `longcat` |
+| [**Qwen Image Edit**](./qwen_image_edit.md) | `qwen` |
+| **Qwen Image LAYERED** | `qwen_layered` |
+| [**Flux.2 [Dev] / Flux.2 [Klein]**](./flux2.md) | `flux2` |
+| [**FLUX.3 Image Klein / Klein Base**](./flux3.md) | `flux2` |
+| [**Boogu Image Edit**](./boogu_image.md) | `z_image_omni` |
+| **Krea2 (Community Edit LoRAs)** | `krea2_ostris_edit` |
+| [**Mage-Flow-Edit**](./mage_flow.md#image-editing) | `mage_flow` |
+| [**LLaDA-Image**](./llada_image.md#image-editing) | `llada_image` |
+| [**Fibo 2.0 Dev / Edit Turbo**](./fibo2.md#examples) | `fibo2` |
+| **Anima (Community Edit LoRAs)** | `cosmos_reference` |
+
+Stable-diffusion.spp also supports basic Unet-based editing models like instruct-pix2pix or CosXL-Edit. This document is not about those.
+
+---
+
+## Configuring Reference Modes (`--ref-image-args`)
+
+For a one-time input transform before reference presets and model processing,
+including cropping, padding, and resizing algorithms, see [Image preprocessing](./image_preprocessing.md).
+
+Different DiT-based editing models require different configurations to process reference images correctly (e.g., whether to use a Vision Language Model (VLM) encoder or pass VAE-encoded images directly to the DiT).
+
+To simplify this, we provide **Presets**. By default, the system automatically selects the best preset based on the model architecture. However, you can override this using the `--ref-image-args` argument.
+
+### Usage
+The `--ref-image-args` argument accepts a comma-separated list of key-value pairs:
+
+**Using a preset:**
+`--ref-image-args "preset=qwen_layered"`
+
+**Using a preset with a specific override:**
+`--ref-image-args "preset=krea2_edit,force_ref_timestep_zero=true"`
+
+### Available Presets
+
+| Preset | Primary Use Case |
+| :--- | :--- |
+| `flux_kontext` | FLUX.1 Kontext |
+| `longcat` | LongCat Image Edit |
+| `flux2` | FLUX.2 models |
+| `fibo2` | Fibo 2.0 Dev / Edit Turbo |
+| `qwen` | Qwen Image Edit |
+| `qwen_layered` | Qwen Image Layered |
+| `z_image_omni` | Boogu, Z-Image Omni |
+| `krea2_ostris_edit` | Most Krea2 Community edit LoRAs (trained with Ostris script) |
+| `mage_flow` | Mage-Flow-Edit |
+| `krea2_edit` | Specifically for [lbouaraba/krea2edit](https://huggingface.co/conradlocke/krea2-identity-edit). (or similar) |
+| `cosmos_reference` | For Anima |
+| `default` | Uses the automatic detection based on model architecture. |
+
+---
+
+## Advanced Parameter Reference
+
+If presets are insufficient, you can manually configure the following parameters via `--ref-image-args`:
+
+| Key | Type | Description | Allowed Values |
+| :--- | :--- | :--- | :--- |
+| `preset` | string | Overrides the automatic preset. | (See the Presets table above) |
+| `pass_to_vlm` | bool | Whether reference images are passed to the VLM encoder. | `true`, `false` |
+| `pass_to_dit` | bool | Whether VAE-encoded references are passed directly to the DiT. | `true`, `false` |
+| `ref_index_mode` | string | Behavior of the RoPE index. | `fixed`, `increase`, `decrease` |
+| `force_ref_timestep_zero` | bool | Forces timestep=0 for reference tokens. | `true`, `false` (Krea2 only) |
+| `resize_before_vae` | bool | Whether reference images are resized before VAE encoding. | `true`, `false` |
+| `resize_vae_to_target` | bool | Use output dimensions for every reference, overriding area-based sizing. | `true`, `false` |
+| `resize_single_ref_to_target` | bool | Use output dimensions when exactly one reference is supplied. | `true`, `false` |
+| `vae_size_buckets` | bool | Choose aligned dimensions within ±10% of the requested area, minimizing aspect-ratio cropping. Ignored when using output dimensions. | `true`, `false` |
+| `vae_center_crop` | bool | Scale to cover the selected dimensions with Lanczos, then center-crop. Otherwise resize directly. | `true`, `false` |
+| `vae_input_max_pixels` | int | VAE area budget, also capped by output area in ordinary resizing. With size buckets, this is the target area and output area does not cap it. | Integer |
+| `vlm_use_vae_input` | bool | Reuse the resized/cropped VAE input for the VLM, before the VLM's own resize. Has no effect when `resize_before_vae=false`. | `true`, `false` |
+| `vlm_resize_mode` | string | How to resize VLM reference inputs. | `longest_side`, `area`, `none` |
+| `vlm_max_size` | int | Maximum VLM input size; interpreted according to `vlm_resize_mode`. | Integer |
+| `vlm_min_size` | int | Minimum VLM input size; interpreted according to `vlm_resize_mode`. | Integer |
+| `vlm_size` | int | Shortcut to set both VLM min and max size to the same value. | Integer |
+| `ref_image_size` | string | MiniMax-H3 reference image sizing before VAE encoding. `match` scales each reference down to the generation's pixel area; `max` uses a 2048 px short edge. Only references above the limit are downscaled; each axis is then snapped to the nearest 32 px multiple (an axis of at least 16 px grows by at most 16 px, and references are never narrower than 32 px). | `match`, `max` |
+
+`ref_image_size` applies to `vid_gen` with MiniMax-H3 Ref2VA. Reference latents are attended to at every sampling step, so `max` can be several times slower.
+
+### Preset Default Values
+
+For a technical overview of how each preset is configured, see the table below.
+
+| Preset | VLM | RoPE Index | Cond Resize | Special Notes |
+| :--- | :---: | :---: | :---: | :--- |
+| `flux_kontext` | No | `fixed` | `none` | |
+| `longcat` | Yes | `fixed` | `area` | |
+| `flux2` | No | `increase` | `none` | |
+| `fibo2` | Yes | `increase` | `longest` | Shared reference crops; clean reference timesteps; see [Fibo 2.0](./fibo2.md) |
+| `qwen` | Yes | `increase` | `area` | |
+| `qwen_layered` | Yes | `decrease` | `area` | |
+| `mage_flow` | Yes | `increase` | `longest` | `vlm_max_size = 384`, VAE input resized to target |
+| `z_image_omni` | Yes | `fixed` | `area` | |
+| `krea2_ostris_edit`| Yes | `increase` | `area` | `force_ref_timestep_zero = true` |
+| `krea2_edit` | Yes | `increase` | `longest` | `vlm_size = 768` |
+| `cosmos_reference` | No | `fixed` | `none` | `resize_before_vae = false` |
+
+**Additional Default Notes:**
+- **VLM Input Sizes:** For most presets, `vlm_max_size` and `vlm_min_size` are set to `-1`, meaning the values are model-dependent and handled automatically. In `area` mode they represent pixel area; in `longest_side` mode they represent a side length in pixels.
+- **VAE Input Size:** `vae_input_max_pixels` defaults to $1024 \times 1024$ pixels (`1048576`).
+- **VAE Geometry:** `resize_vae_to_target` defaults to `true` for Mage-Flow and LLaDA-Image. Fibo enables `resize_single_ref_to_target`, `vae_size_buckets`, and `vae_center_crop`. These options are otherwise `false`. Bucket dimensions use the model's alignment (normally 16, Qwen Image 32), with a minimum side of four alignment units; if no bucket fits, output dimensions are used.
+- **Shared Encoder Input:** `vlm_use_vae_input` defaults to `true` for Fibo and Qwen Image 2.1, otherwise `false`. All these options can be overridden individually after selecting a preset.
@@ -21,7 +21,7 @@ You can run ERNIE-Image with stable-diffusion.cpp on GPUs with 4GB of VRAM — o
 ### ERNIE-Image-Turbo

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\ernie-image-turbo.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\ministral-3-3b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 8 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\ernie-image-turbo.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\ministral-3-3b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 8 -v --offload-to-cpu
 ```

 <img width="256" alt="ERNIE-Image Turbo example" src="../assets/ernie_image/turbo_example.png" />
@@ -29,7 +29,7 @@ You can run ERNIE-Image with stable-diffusion.cpp on GPUs with 4GB of VRAM — o
 ### ERNIE-Image

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\ernie-image-UD-Q4_K_M.gguf --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\ministral-3-3b.safetensors -p "a lovely cat" --cfg-scale 5.0 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\ernie-image-UD-Q4_K_M.gguf --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\ministral-3-3b.safetensors -p "a lovely cat" --cfg-scale 5.0 -v --offload-to-cpu
 ```

 <img width="256" alt="ERNIE-Image example" src="../assets/ernie_image/example.png" />
@@ -2,6 +2,10 @@

 You can use ESRGAN—such as the model [RealESRGAN_x4plus_anime_6B.pth](https://github.com/xinntao/Real-ESRGAN/releases/download/v0.2.2.4/RealESRGAN_x4plus_anime_6B.pth)—to upscale the generated images and improve their overall resolution and clarity.

+MoSR models with the PixelShuffle upsampler, such as [4xNomos2_hq_mosr.pth](https://github.com/Phhofm/models/releases/download/4xNomos2_hq_mosr/4xNomos2_hq_mosr.pth), are also supported. The architecture is detected from the weights, so they are used the same way as ESRGAN models.
+
+RGBA images, including Qwen Image 2.1 output, keep their alpha channel during model upscaling and hires fix. The upscaler processes the RGB channels; the alpha channel is resized with bilinear interpolation and recombined with the upscaled image.
+
 - Specify the model path using the `--upscale-model PATH` parameter. example:

 ```bash
@@ -0,0 +1,53 @@
+# How to Use
+
+Fibo 2.0 uses a Fibo diffusion transformer, the standard FLUX.2 VAE, and Qwen3-VL 4B as the LLM text encoder. The Qwen tokenizer is built in.
+
+## Download weights
+
+- Download Fibo 2.0 Dev
+    - safetensors: https://huggingface.co/Comfy-Org/Fibo-2.0/tree/main/diffusion_models
+- Download Fibo 2.0 Edit Turbo
+    - safetensors: https://huggingface.co/Comfy-Org/Fibo-2.0/tree/main/diffusion_models
+- Download vae
+    - safetensors: https://huggingface.co/black-forest-labs/FLUX.2-dev/tree/main
+- Download Qwen3-VL 4B
+    - safetensors: https://huggingface.co/Comfy-Org/Fibo-2.0/tree/main/text_encoders
+    - gguf: https://huggingface.co/Qwen/Qwen3-VL-4B-Instruct-GGUF/tree/main
+
+The original Diffusers weights from [briaai/Fibo-2.0-dev](https://huggingface.co/briaai/Fibo-2.0-dev) and [briaai/Fibo-2.0-edit-turbo](https://huggingface.co/briaai/Fibo-2.0-edit-turbo) are also supported. Pass the transformer's `diffusion_pytorch_model.safetensors.index.json` to `--diffusion-model` and keep its shards alongside it.
+
+INT8 convrot weights load directly; see [INT8 convrot](./int8_convrot.md) for backend requirements.
+
+## Examples
+
+### Fibo 2.0 Dev
+
+Dev supports both text-to-image generation and reference-image editing.
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\fibo2_dev_int8_convrot.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --llm_vision ..\models\text_encoders\Qwen3-VL-4B-Instruct-mmproj-BF16.gguf --vae ..\models\vae\flux2_ae.safetensors -p "a lovely cat holding a sign says 'fibo2.cpp'" --cfg-scale 5.0 --sampling-method euler -v  --offload-to-cpu
+```
+
+<img src="../assets/fibo2/t2i.png" />
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\fibo2_dev_int8_convrot.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --llm_vision ..\models\text_encoders\Qwen3-VL-4B-Instruct-mmproj-BF16.gguf --vae ..\models\vae\flux2_ae.safetensors -r ..\assets\flux\flux1-dev-q8_0.png -p "change 'flux.cpp' to 'fibo2.cpp'" --cfg-scale 5.0 --sampling-method euler -v  --offload-to-cpu
+```
+
+<img src="../assets/fibo2/dev_edit.png" />
+
+### Fibo 2.0 Edit Turbo
+
+Edit Turbo is distilled for reference-image editing in 4 steps with CFG 1.
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\fibo2_edit_turbo_int8_convrot.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --llm_vision ..\models\text_encoders\Qwen3-VL-4B-Instruct-mmproj-BF16.gguf --vae ..\models\vae\flux2_ae.safetensors -r ..\assets\flux\flux1-dev-q8_0.png -p "change 'flux.cpp' to 'fibo2.cpp'" --cfg-scale 1 --steps 4 --sampling-method euler -v  --offload-to-cpu
+```
+
+The BF16 files accept the same options. Use dimensions divisible by 16. Repeat `-r` to use multiple reference images; instructions can refer to `<image_1>`, `<image_2>`, etc. See [Image Editing](./edit.md) for reference-image preprocessing options. Separate GGUF text encoders need their vision projector passed with `--llm_vision` for editing.
+
+The automatic `fibo2` reference preset combines `resize_single_ref_to_target=true`, `vae_size_buckets=true`, `vae_center_crop=true`, and `vlm_use_vae_input=true`. A single reference is scaled to cover the output dimensions and center-cropped. Multiple references use dimensions divisible by 16, chosen within ±10% of one megapixel to minimize cropping for each image's aspect ratio; `vae_input_max_pixels` overrides that target area. The VAE and vision encoder receive the same crops, with smaller copies used by the vision encoder. These are generic reference-image options and can be overridden individually: for example, `--ref-image-args vae_center_crop=false` uses direct resizing, while `resize_before_vae=false` skips this automatic crop and resize step entirely.
+
+The vision encoder defaults to `vlm_resize_mode=longest_side,vlm_max_size=768`. Override these with `--ref-image-args`, for example `vlm_max_size=512`, or use `vlm_resize_mode=none` to skip size-based scaling. Vision inputs are still aligned to the encoder's patch grid.
+
+Plain prompts and structured JSON captions are accepted. This integration uses the caption-only text-rendering path (`render_text_images=False`); it does not generate auxiliary glyph images or expand prompts with Bria VGL. Native Fibo mask preprocessing is not connected to `--mask`.
@@ -17,7 +17,7 @@ You can download the preconverted gguf weights from [FLUX.1-dev-gguf](https://hu

 For example:
 ```
-.\bin\Release\sd-cli.exe -M convert -m ..\..\ComfyUI\models\unet\flux1-dev.sft -o ..\models\flux1-dev-q8_0.gguf -v --type q8_0
+.\bin\Release\sd-cli.exe -M convert -m ..\models\unet\flux1-dev.sft -o ..\models\flux1-dev-q8_0.gguf -v --type q8_0
 ```

 ## Run
@@ -16,7 +16,7 @@
 ### Examples

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux2-dev-Q4_K_S.gguf --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\Mistral-Small-3.2-24B-Instruct-2506-Q4_K_M.gguf -r .\kontext_input.png -p "change 'flux.cpp' to 'flux2-dev.cpp'" --cfg-scale 1.0 --sampling-method euler -v --diffusion-fa --offload-to-cpu
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux2-dev-Q4_K_S.gguf --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\Mistral-Small-3.2-24B-Instruct-2506-Q4_K_M.gguf -r .\kontext_input.png -p "change 'flux.cpp' to 'flux2-dev.cpp'" --cfg-scale 1.0 --sampling-method euler -v --offload-to-cpu
 ```

 <img alt="flux2 example" src="../assets/flux2/example.png" />
@@ -42,19 +42,19 @@
 ### Examples

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-4b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_4b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 4 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-4b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_4b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 4 -v --offload-to-cpu
 ```

 <img alt="flux2-klein-4b" src="../assets/flux2/flux2-klein-4b.png" />

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-4b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_4b.safetensors -r .\kontext_input.png -p "change 'flux.cpp' to 'klein.cpp'" --cfg-scale 1.0 --sampling-method euler -v --diffusion-fa --offload-to-cpu --steps 4
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-4b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_4b.safetensors -r .\kontext_input.png -p "change 'flux.cpp' to 'klein.cpp'" --cfg-scale 1.0 --sampling-method euler -v --offload-to-cpu --steps 4
 ```

 <img alt="flux2-klein-4b-edit" src="../assets/flux2/flux2-klein-4b-edit.png" />

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-base-4b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_4b.safetensors -p "a lovely cat" --cfg-scale 4.0 --steps 20 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-base-4b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_4b.safetensors -p "a lovely cat" --cfg-scale 4.0 --steps 20 -v --offload-to-cpu
 ```

 <img alt="flux2-klein-base-4b" src="../assets/flux2/flux2-klein-base-4b.png" />
@@ -78,19 +78,19 @@
 ### Examples

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-9b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_8b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 4 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-9b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_8b.safetensors -p "a lovely cat" --cfg-scale 1.0 --steps 4 -v --offload-to-cpu
 ```

 <img alt="flux2-klein-9b" src="../assets/flux2/flux2-klein-9b.png" />

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-9b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_8b.safetensors -r .\kontext_input.png -p "change 'flux.cpp' to 'klein.cpp'" --cfg-scale 1.0 --sampling-method euler -v --diffusion-fa --offload-to-cpu --steps 4
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-9b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_8b.safetensors -r .\kontext_input.png -p "change 'flux.cpp' to 'klein.cpp'" --cfg-scale 1.0 --sampling-method euler -v --offload-to-cpu --steps 4
 ```

 <img alt="flux2-klein-9b-edit" src="../assets/flux2/flux2-klein-9b-edit.png" />

 ```
-.\bin\Release\sd-cli.exe --diffusion-model  ..\..\ComfyUI\models\diffusion_models\flux-2-klein-base-9b.safetensors --vae ..\..\ComfyUI\models\vae\flux2_ae.safetensors  --llm ..\..\ComfyUI\models\text_encoders\qwen_3_8b.safetensors -p "a lovely cat" --cfg-scale 4.0 --steps 20 -v --offload-to-cpu --diffusion-fa
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-2-klein-base-9b.safetensors --vae ..\models\vae\flux2_ae.safetensors  --llm ..\models\text_encoders\qwen_3_8b.safetensors -p "a lovely cat" --cfg-scale 4.0 --steps 20 -v --offload-to-cpu
 ```

 <img alt="flux2-klein-base-9b" src="../assets/flux2/flux2-klein-base-9b.png" />
@@ -0,0 +1,50 @@
+# How to Use
+
+FLUX.3 Image Klein uses a FLUX.3 diffusion transformer, the FLUX.2 VAE, and Qwen3-VL 4B as the LLM text encoder.
+
+## Download weights
+
+- Download FLUX.3 Image Klein
+    - safetensors: https://huggingface.co/black-forest-labs/flux-3-image-klein/tree/main
+    - gguf: https://huggingface.co/leejet/flux-3-image-klein-GGUF/tree/main
+- Download FLUX.3 Image Klein Base
+    - safetensors: https://huggingface.co/black-forest-labs/flux-3-image-klein-base/tree/main
+    - gguf: https://huggingface.co/leejet/flux-3-image-klein-base-GGUF/tree/main
+- Download vae
+    - safetensors: https://huggingface.co/black-forest-labs/FLUX.2-dev/tree/main
+    - The `ae.safetensors` included with FLUX.3 can also be used.
+- Download Qwen3-VL 4B
+    - safetensors: https://huggingface.co/Comfy-Org/Krea-2/tree/main/text_encoders
+    - gguf: https://huggingface.co/Qwen/Qwen3-VL-4B-Instruct-GGUF/tree/main
+
+Use the root `model.safetensors` for the diffusion model.
+
+## Examples
+
+### FLUX.3 Image Klein
+
+The distilled model uses 16 steps and CFG 1.0.
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\flux-3-image-klein.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --vae ..\models\vae\flux2_ae.safetensors -p "a lovely cat holding a sign says 'flux3.cpp'" --cfg-scale 1.0 --steps 16 --sampling-method euler -v --offload-to-cpu
+```
+
+<img src="../assets/flux3/flux3-image-klein.png" />
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model  ..\models\diffusion_models\flux-3-image-klein.safetensors --vae ..\models\vae\flux2_ae.safetensors --llm  ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf -r ..\assets\flux\flux1-dev-q8_0.png -p "change 'flux.cpp' to 'flux3-image-klein.cpp'" --cfg-scale 1.0 --sampling-method euler -v  --offload-to-cpu --steps 16
+```
+
+<img src="../assets/flux3/flux3-image-klein-edit.png" />
+
+Repeat `-r` to use multiple reference images, up to 10. See [Image Editing](./edit.md) for reference-image preprocessing options.
+
+### FLUX.3 Image Klein Base
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\flux-3-image-klein-base.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --vae ..\models\vae\flux2_ae.safetensors -p "a lovely cat holding a sign says 'flux3.cpp'" --cfg-scale 3.5 --steps 20 --sampling-method euler -v --offload-to-cpu
+```
+
+<img src="../assets/flux3/flux3-image-klein-base.png" />
+
+The default discrete scheduler uses the upstream time shift of 6.93 for both models. Use `--flow-shift` to override it.
@@ -0,0 +1,165 @@
+# Generation results and metadata
+
+`generate_image()`, `generate_video()`, `adetail_image()` and `upscale()` return
+owned results containing the media and its generation record. This is a breaking
+C API change: applications and bindings must update their calls and rebuild.
+
+```c
+sd_img_gen_params_t params;
+sd_img_gen_params_init(&params);
+params.prompt = "a small house";
+
+sd_image_generation_result_t* result = NULL;
+if (generate_image(ctx, &params, &result)) {
+    for (int i = 0; i < result->count; ++i) {
+        const sd_image_result_t* item = &result->items[i];
+        /* Consume item->image and item->generation_info together. */
+    }
+    free_sd_image_generation_result(result);
+}
+```
+
+`result_out` is required and is set to `NULL` on failure. A successful result
+owns its pixels and UTF-8 JSON strings. The video result additionally owns its
+audio and contains the effective `fps` and `frame_count`. Results remain valid
+after freeing the generation context. Use the matching library result destructor;
+do not separately free fields unless transferring ownership and clearing them.
+Both result destructors accept `NULL`.
+
+## Reproduction parameters
+
+The `sdcpp.generation/v1` JSON is a snapshot of generation parameters for the
+same loaded context and the same input images, masks, reference media and other
+assets. It is created after `GenerationRequest` resolves the request, before
+preprocessing or sampling. A failed call returns no result or metadata.
+
+Fields follow the C generation parameter structs: `prompt`, `negative_prompt`,
+`width`, `height`, `seed`, `sample_params`, `hires`, `cache`,
+`vae_tiling_params`, `image_preprocess`, and the image/video-specific options.
+`kind` distinguishes image and video parameters. There is no inferred `task`.
+Enums use their API string names. Seeds are decimal strings to preserve their
+full integer range. Omitted optional fields retain the parameter initializer's
+defaults, including automatic values represented by infinity in the C API.
+
+Disabled optional features are omitted entirely: `cache`, `hires`, VAE tiling,
+zero/nonpositive `shifted_timestep`, inactive SLG, false circular axes, and empty
+preprocessing/reference arguments. SLG layers remain present when `slg_uncond`
+is enabled even with zero SLG scale. Enabled caches store their mode's parameters;
+TaylorSeer retains the shared DBCache thresholds and limits used by the runtime.
+Image CFG at its neutral value of 1 is omitted and resolves back to 1 by default.
+
+Image strength, ControlNet/IP-Adapter strength and PhotoMaker settings are stored
+when their corresponding input images are supplied. PuLID requires an embedding
+path. Layer count is stored for layered generation; video VACE and high-noise
+settings require those models. High-noise settings omit scheduler, custom sigmas,
+flow shift and timestep shifting because that pass shares the base configuration.
+Explicit nonempty argument strings are retained without rewriting their contents.
+Core sampling values such as seed, sampler, scheduler, steps and guidance remain
+explicit even when equal to their defaults. ADetailer always retains its inpaint
+strength because its input crops are created later.
+Model-required VAE tiling is resolved before serialization, so forced tiling
+(such as MiniMax-H3) remains present even if the caller did not enable it.
+
+The request fixes the seed, dimensions, clip skip, guidance, sampler, scheduler
+and eta before taking the snapshot. `custom_sigmas` is stored only when supplied
+as an input. Automatically computed schedules, preprocessing transforms,
+conditioning/cache decisions, detector regions and output statistics are not
+stored: they are recomputed from the same parameters and inputs.
+
+Each image records the seed used to sample its latent. `batch_count` is omitted;
+replay image parameters with `batch_count = 1`. Layers or frames decoded from
+the same latent share its seed, so layered generation reproduces that group of
+layers. Video frames share the video request. Partial batch cancellation returns
+metadata only for completed outputs. Full cancellation returns no result.
+
+The current image pipeline prepares input latents once using the starting seed
+and shares them across a batch. With a stochastic VAE encoder, replaying a later
+image alone using its sampling seed can change those input latents. Per-image
+seeds therefore do not guarantee independent reproduction for these paths.
+
+To replay a record, initialize the matching C parameter struct, populate its
+fields from the JSON, restore the same input buffers, and call generation with
+the same context. Keep custom-sigma/SLG arrays and string storage alive for the
+call. Paths supplied as generation parameters (LoRAs, hires models and identity
+embeddings) are retained. `models` identifies successfully loaded context model
+files by component (`model`, `diffusion_model`, `clip_l`, `vae`, etc.), using
+filenames with extensions and omitting directories. Optional components disabled
+for the request and preview-only TAE are omitted. WebUI text includes `Model`
+(the diffusion model when supplied separately, otherwise the checkpoint, without
+the extension) and `VAE` for a separate active VAE. These names are identifiers,
+not hashes or sufficient configuration to rebuild a context; replay still requires
+the same model files and context settings. The JSON does not embed media buffers
+or model weights.
+
+When `--lora-model-dir` is configured, CLI/server output records express LoRA
+paths relative to that directory, preserving subdirectories and file extensions
+and using `/` separators. This applies to file metadata, API responses and
+ADetailer records. Resolve these paths against the same LoRA directory before
+passing them back to the C API. Core results retain the supplied loading paths.
+Files outside the directory retain `..` components; paths on another volume,
+or records without a configured LoRA directory, retain their original paths.
+
+## Postprocessing
+
+Upscale records `upscale.upscale_factor`. ADetailer records `adetailer` and
+`inpaint_params`, with a concrete starting seed chosen before detection. These
+operations require the same upscaler/detector context and source image on replay.
+They do not store detection results or per-region generation traces.
+
+The CLI/server append operation parameter snapshots to `postprocess` in call
+order. The source generation parameters remain unchanged. Frame postprocessing
+uses the same operation parameters for each video frame. Standalone
+postprocessing does not invent source diffusion parameters for an input image.
+
+## File metadata
+
+Image encoders accept the native record and derive WebUI infotext from it.
+Images use the existing metadata locations and combined text layout:
+`<WebUI infotext>, SDCPP: <native JSON>`. Standalone postprocessing without
+source diffusion parameters stores only native JSON. Video metadata stores
+only native JSON.
+
+| Container | Metadata location |
+| --- | --- |
+| PNG | `tEXt`, keyword `parameters` |
+| JPEG | COM segment, `parameters` label followed by a NUL byte |
+| WebP image / animated WebP | XMP property `sdcpp:parameters` |
+| AVI | RIFF INFO `ICMT` |
+| WebM | `COMMENT` tag |
+
+Inspect images or videos with `sd-cli -M metadata --image <file>`; the same
+command supports all containers above, with `--metadata-format json` for a
+structured report. Video metadata inspection does not require a model or video
+decoder.
+Matroska containers are also supported, including `.webm` outputs whose audio
+codec causes the muxer to select the Matroska document type.
+
+The XMP namespace is `https://github.com/leejet/stable-diffusion.cpp/ns/1.0/`.
+PNG and JPEG retain the existing UTF-8 text payload convention. Memory encoding
+and file saving use the same implementation. The core result and API
+`generation_info` fields remain pure JSON; the combined layout applies only
+to embedded image metadata.
+
+WebUI conversion is lossy: known samplers and schedules have explicit name
+mappings; unsupported ones use `SDCPP Sampler`/`SDCPP Schedule` fields. Hires
+keeps the base request size and separate configured hires size/steps. The native
+record retains the full generation parameters. Standalone postprocessing has no WebUI
+infotext because its original diffusion settings are unknown.
+
+WebUI prompt text appends ordinary LoRAs as `<lora:filename_without_extension:weight>`.
+Directory components are omitted to match WebUI's filename lookup; the native
+`loras` array retains relative paths and weights. High-noise-only LoRAs and names
+that cannot be represented by this tag syntax remain in the native record only.
+The native prompt is not modified, and postprocessing LoRAs are not added to the
+source generation prompt.
+
+JPEG requires the COM payload, including the label and terminating NUL bytes,
+to fit within 65,533 bytes. Larger records fail encoding with an error rather
+than being truncated. Use PNG or WebP for larger records.
+
+Disabling file metadata omits both image representations, but does not disable
+the core generation record or its API response. Native async image results and
+OpenAI image response items contain `generation_info`; video results contain a
+sequence-level `generation_info`. The SD API `info` object contains per-output
+`generation_info` alongside derived `infotexts` and each image's seed. Its outer
+`parameters` field continues to echo the request.
@@ -0,0 +1,552 @@
+# ggml operator extensions
+
+The patched ggml exposes `ggml-ext.h` for application-owned operators, kernel
+overrides, and graph fusion rules. Registrations are in-process and scoped to a
+backend device. Register them before capability queries or graph allocation;
+callback code and user data must remain alive until all dependent graphs and
+backends have been destroyed. Registration names are unique, and duplicate
+registrations fail rather than replacing live implementations.
+
+## Operator definitions
+
+`ggml_ext_register_op` registers a versioned name, attribute size, and an output
+metadata inference callback. `ggml_ext_op_create` creates an ordinary ggml tensor
+with explicit input dependencies. The initial interface supports one contiguous,
+out-of-place output with an existing `ggml_type`.
+
+Attributes are copied into the unused portion of `op_params`. They are small,
+plain values; use `memcpy` when reading attributes that need stricter alignment.
+Auxiliary tensors must appear in `src`, not in an attribute pointer or user data.
+Registered operators use `GGML_OP_CUSTOM` internally, with an identifiable
+callback marker. Unsupported devices do not mistake them for CPU callbacks.
+These process-local descriptors are not a graph serialization format.
+
+## Kernel selection
+
+`ggml_ext_register_kernel(device, &desc)` associates a kernel with a built-in
+`ggml_op`, or with a registered custom operator. `supports` describes correctness
+constraints; `preferred` optionally decides whether to override an available
+native implementation. Higher priorities are considered first. When native
+execution is unsupported, a supported extension can run regardless of its
+preference callback.
+
+Capability queries include registered single-node implementations, so the
+scheduler can select a device even when its native implementation rejects the
+operator. Callbacks inspect metadata only and must not recursively call the
+public capability query for the same device and operator. A fusion-only
+registration does not claim single-node support.
+
+`ggml_ext_dev_supports_native_op` queries the original device implementation
+without extension dispatch. Fusion policies can use it to select a supported
+extension even when their performance preference normally favors native kernels.
+
+`prepare` creates optional host-side plan state and reports workspace bytes.
+It can run during memory measurement, before tensor addresses are assigned, and
+must not launch work or allocate device buffers. Workspace is allocated by ggml
+before execution and included in scheduler memory estimates and buffer-size
+reports. `execute` receives the backend, workspace buffer, and current tensor
+addresses. `destroy` releases prepared state. A preparation or execution error
+is propagated; execution never retries a native implementation after launch.
+
+Workspace is scratch storage shared by sequential extension steps in a plan;
+its capacity is the maximum request, not the sum across nodes. Contents are not
+preserved between calls. Queue work in backend execution order, and finish any
+private-stream use before subsequent backend work can access the workspace.
+Persistent state belongs in `prepare`/`destroy` state or explicit tensor inputs.
+
+Overriding a built-in operator preserves its numerical and aliasing contract,
+including allocator reuse of input storage. Explicit views/in-place nodes are
+not intercepted by this first interface. Rebuild the scheduler allocation after
+changing graph metadata, selection policy, or registrations.
+
+## Fusion
+
+A `match` callback returns the number of consecutive nodes it can replace.
+The scheduler freezes this selection before allocating tensors. ggml checks
+compute flags, external uses, output flags, and views; the initial implementation
+accepts a conservative set of pure operators with only the final node exposed
+as the fused output. The rule remains responsible for mathematical equivalence,
+dtype, shape, layout, and hardware constraints. Internal views are accepted when
+their storage belongs to the fused subgraph. In-place intermediates are accepted
+only when their storage owner is an earlier node in the same match. External
+consumers and output flags still prevent eliminating those results. The final
+output must own storage, use a terminal `CPY`, or write back to an earlier owning
+node within the match. Internal writeback must cover exactly the owner's full
+shape, type, strides and storage offset, without external consumers of that
+owner. Output flags propagated to that owner are retained. Terminal `CPY` keeps
+its original destination and write semantics. The extension must handle input/output storage overlap, including
+allocator reuse, with scratch space when necessary. Other in-place writes are
+not eligible for extension fusion when they target storage outside the match.
+
+An optional `optimize` callback runs before matching and allocation. Use
+`ggml_ext_graph_delay` to move a producer range after independent nodes. The
+helper rejects dependencies, unknown side effects, and writes that alias moved
+inputs or outputs. It can cross independent copies and known in-place operators
+when their storage does not alias the moved inputs or outputs. Owning dtype
+conversions can also be moved or crossed.
+Reordering must preserve edges and tensor metadata. Fusion legality is checked
+again after reordering.
+
+External inputs stay allocated through the final node. Native graph optimization
+and execution operate on the remaining contiguous ranges, so native fusion cannot
+consume nodes selected by an extension. Fused intermediates are excluded from
+allocation through the backend's `skip_alloc` callback and
+`ggml_gallocr_set_skipped_tensors`. Their metadata and graph edges remain available,
+but execution must not access their data or buffers. Storage backing the final
+output is retained, including an internal owner used by an in-place final write.
+The skipped set is replaced for each allocation graph; changing which tensors are skipped
+invalidates cached allocation assignments even when tensor shapes are unchanged.
+
+Automatic extension fusion is enabled in scheduler execution with one pipeline
+copy and without evaluation callbacks. Multiple pipeline copies and evaluation
+callbacks retain individual operations. Direct backend graph computation supports
+single-node extensions; it does not fuse an already allocated graph. Direct calls
+use temporary plans and synchronize before releasing them; scheduler plans and
+workspaces are reused across executions.
+
+Use the backend graph APIs or the scheduler. The legacy CPU `ggml_graph_compute`
+and backend `graph_plan_*` interfaces do not dispatch these extensions.
+
+## Backend services
+
+Packed ComfyUI INT4 linear layers use `sd.linear_w4_convrot.v1`, built by
+`ggml_ext_mul_mat_w4_convrot`. The operator consumes byte-packed I8 weights and the
+explicit output of `sd.quantize_i8_convrot.v1`; F32 callers get a quantization
+node during graph construction. Logical output dimensions are stored in the
+operator attributes, so backend copies of the packed activation do not need to
+retain the producer's op marker. W4A4 uses row scales; W4A8 additionally exposes
+the codebook, channel scales, and FP8 relative scales as tensor dependencies.
+CPU, CUDA/HIP/MUSA, and Vulkan register kernels through the same extension API
+as INT8. No ggml submodule changes are required.
+
+Services are discovered through `ggml_backend_reg_get_proc_address`:
+
+| Service | Purpose |
+| --- | --- |
+| `ggml_ext_cpu_parallel_v1` | Run an application callback on the CPU backend's configured thread pool, with its abort callback. |
+| `ggml_ext_cuda_context_v1` | Borrow the current CUDA device and stream, ordered with native work. |
+| `ggml_ext_cuda_device_v1` | Get the physical CUDA ordinal for capability checks before backend creation, including virtual ggml devices. |
+| `ggml_ext_cuda_blas_v1` | Borrow the current stream's cuBLAS handle during execution, with host scalar pointers. |
+| `ggml_ext_gpu_device_v1` | Query CUDA/HIP/MUSA runtime kind, physical ordinal, shared-memory limit, warp size, and INT8 BLAS capability. |
+| `ggml_ext_gpu_context_v1` | Borrow the matching runtime's device and execution stream. |
+| `ggml_ext_gpu_blas_v1` | Borrow the matching cuBLAS, hipBLAS, or muBLAS handle on that stream. |
+| `ggml_ext_metal_device_v1` | Borrow the Metal device for application pipeline compilation. |
+| `ggml_ext_metal_buffer_v1` | Resolve a backend buffer range to its Metal buffer and offset, including private storage and split allocations. |
+| `ggml_ext_metal_submit_v1` | Encode application commands on a backend-owned command buffer, ordered with native Metal work and tracked by backend synchronization. |
+
+The borrowed cuBLAS handle remains owned by the backend. Extensions must not retain
+or destroy it, or change its stream, pointer mode, math mode, or workspace. This
+CUDA-specific BLAS service is retained for existing callers; the runtime-aware
+services support CUDA, HIP, and MUSA with the same ownership rules. Extensions
+call vendor runtime and BLAS APIs directly. These services do not execute GEMM,
+quantization, or model fusion on the application's behalf.
+
+Workspace is a backend buffer. Its base is a CUDA device pointer on CUDA; callers
+must use the representation appropriate to their backend. Vulkan extensions use
+the command recording services described below.
+
+CUDA INT8 Linear, RoPE, ConvRot, RMS affine and INT8 gated residual kernels omit output
+scratch when allocation guarantees that a fresh owning output cannot overlap
+live inputs. Existing outputs and views retain overlap protection. Internal
+GEMM and quantization workspace is still reserved when required; kernels also
+handle a null workspace when the entire step needs no scratch.
+
+Compute workspace measurement preserves the storage identities and offsets of
+external tensors, including views. Separate allocations must not share a synthetic
+buffer identity: that can invent overlap, overestimate scratch, and repeatedly
+rebuild a workspace whose actual allocation is smaller than the estimate.
+
+## sd.cpp integration
+
+`src/core/ggml_extensions.cpp` and `.h` contain the shared operator schemas,
+validation, and backend registration entry point. Graph construction helpers live
+in `src/core/ggml_extend.cpp` and `.h`; they take `GGMLRunnerContext*` so backend
+selection, compatibility workarounds, and shared ConvRot nodes use the same
+context as the calling model. Per-call precision, scaling, attention, and padding
+options remain explicit. Model blocks retain weight loading and adapter
+orchestration, while the shared helpers own FP8 fallback, embedding lookup,
+full/partial RoPE, grouped convolution, and backend-dependent repeat handling.
+LoRA/LoKr construction, n-mode products, and Kronecker products stay in
+`src/model/adapter/lora_ops.cpp` and `.h`, with the same runner-context interface.
+CPU kernels and fusion rules live in `src/core/cpu/`, with registration collected
+in `cpu/registration.cpp`. CUDA-style and Vulkan implementations live in their
+respective backend directories.
+
+`src/core/metal/int8.mm` registers Metal INT8 linear and ConvRot kernels. Their
+Metal source is embedded at build time and compiled for the selected device
+during registration. The public `ggml-ext-metal.h` services expose borrowed
+objects without private backend headers. The submit callback only encodes work;
+GGML owns submission, cancellation checks, command-buffer lifetime, and
+asynchronous error tracking. Rotation scratch and any overlapping output use
+the scheduler workspace. No full floating-point weight copy is allocated.
+When runtime compilation supports Metal 4 TensorOps, INT8 multiplication uses
+`matmul2d` with INT32 cooperative results and fused scale/bias restoration.
+The SIMD-group implementation remains the fallback for older systems and
+devices where the TensorOps pipeline is unavailable.
+TensorOps multiplies at most 1024 pairs per partial result and combines those
+results with wrapping integer addition. This preserves exact INT32 semantics
+even when TensorOps uses floating-point arithmetic internally.
+
+`ggml_ext_modulate(ctx, x, shift, scale)` emits `sd.modulate.v1` for F32 inputs
+when the runner context uses CPU, CUDA/HIP/MUSA, or Vulkan.
+Its explicit inputs are `x`, `scale`, and optional `shift`; it computes
+`(x + x * scale) + shift` with separate FP32 multiply/add rounding. The output
+is contiguous and out of place. Inputs may be strided and use ggml's
+dimension-wise repetition rules. Callers prepare token, batch, or frame layouts
+before constructing the operation; normalization and residual gates remain
+separate operations. CPU, CUDA/HIP/MUSA, and Vulkan have standalone kernels.
+Other backends (including Metal, SYCL, and OpenCL), upstream ggml builds, and
+other input types retain the original MUL/ADD graph.
+
+ConvRot fusion matches the Modulate node directly, including scale-only
+modulation, instead of recognizing its former MUL/ADD/ADD decomposition.
+CUDA accepts contiguous per-channel parameters for `Modulate -> ConvRot` and
+`RMSNorm -> Mul(weight) -> Modulate -> ConvRot`. Vulkan limits the first fusion
+to widths 4096/6144/5376 and the RMSNorm fusion to width 5376.
+Other layouts execute the standalone operation followed by ConvRot. The
+gated-residual MUL/ADD fusion remains available for residual branches.
+
+`src/core/cpu/rope.cpp` registers the CPU implementation of
+`sd.rope_apply.v1` and a
+`RMSNorm -> Mul -> RoPE` fusion. The fusion accepts a one-dimensional weight
+vector. Both interleaved and split-half RoPE, partial rotation, and strided input
+rows are supported. Model construction uses the custom operator on CPU, CUDA,
+HIP, MUSA, and Vulkan. Other backends can use the scheduler's CPU fallback.
+The former `ggml_rope_apply` API, builtin opcode, and CPU implementation have
+been removed from ggml.
+
+`src/core/cuda/rope.cu` supplies standalone CUDA RoPE,
+`RMSNorm -> RoPE` fusion for non-affine normalization, and
+`RMSNorm -> Mul -> RoPE [-> Scale -> FP16 Copy]` fusion. CuTe attention shares
+its RMSNorm/RoPE kernel. Partial rotation and arbitrary float-aligned input and
+position-encoding strides are supported by the standalone kernel; RMSNorm fusion
+requires contiguous position encoding and, when present, weights. The non-affine
+fusion uses 64-channel F32 inputs and split-half RoPE. Copy fusion uses workspace
+when the destination overlaps an input. These registrations do not require
+CUTLASS. The old CUDA/HIP/MUSA RoPE dispatch is removed.
+
+Shared CUDA-style kernels live in `src/core/cuda/` and are compiled separately
+for each enabled runtime. `SD_CUDA_NAMESPACE` expands to `sd_cuda`, `sd_hip`, or
+`sd_musa`, so implementations can coexist. `runtime.h` handles the small runtime
+differences, including logical 32-lane shuffles on HIP. The application defines
+and launches its kernels directly; it includes only public ggml headers.
+CMake generates the registration table from enabled backend options. The shared
+entry point uses one `SD_BACKEND_EXTENSIONS` macro; backend registrars verify the
+device kind, and disabled backends add no symbol references.
+
+`SD_USE_UPSTREAM_GGML` excludes these registrations. The stable-diffusion public C
+API is unchanged. INT8 ConvRot retains its existing packed activation layout and
+separate weight scale tensor.
+
+`src/core/cuda/cute_attention.cu` supplies CuTe Flash Attention on CUDA when
+CUTLASS headers are configured. It registers an override of `FLASH_ATTN_EXT`,
+with the existing sequence-length preference and native fallback. It also
+registers Q preparation (`RMSNorm -> Mul -> RoPE`), output scaling, and packed
+ConvRot output fusion. These kernels use only the public extension API, borrowed
+CUDA stream, and scheduler-owned workspace; they do not use ggml's CUDA pool or
+private conversion helpers. The existing CUDA fusion and CuTe disable environment
+variables are honored.
+Aligned contiguous FP16 K/V inputs are read directly without reserving conversion
+buffers. Fresh F32 outputs are written by the attention epilogue, retaining the
+original FP16 rounding before output scaling. Packed ConvRot and preallocated or
+view outputs retain the intermediate half buffer.
+
+`src/core/cuda/i8_projection.cu` supplies the INT8 GEMM projection epilogues:
+`GEMM -> Reshape -> RMSNorm -> Mul -> RoPE`, optional scaling and FP16 copy,
+and K/V prefix concatenation with FP16 output. CuTe can consume a directly
+produced FP16 Q without materializing the FP32 GEMM or RoPE intermediates.
+Independent affine operands, including broadcasted constants, are computed before
+their normalization/projection
+chain when graph dependencies allow it. The attention optimizer can then move
+the complete five-node projection when safe.
+These rules accept custom RoPE nodes. They preserve the original
+large-matrix selection thresholds, full interleaved head dimension 128, packed
+activation layout, and separate weight scales. External consumers prevent fusion;
+aliasing terminal copies use scheduler workspace. The CUDA projection kernels
+and matchers now live in sd.cpp. `GGML_CUDA_FORCE_CUBLAS` disables these epilogues.
+
+`src/core/ggml_extensions.cpp` defines `sd.linear_i8.v1` and
+`sd.quantize_i8_convrot.v1`. `src/core/cpu/int8.cpp` supplies their CPU
+implementations. The CPU kernels use the
+configured thread pool and scheduler workspace; the old CPU INT8 GEMM and ConvRot
+implementations have been removed from ggml.
+Model code keeps calling `ggml_ext_linear_i8_tensorwise`,
+`ggml_ext_mul_mat_i8_tensorwise`, and `ggml_ext_quantize_i8_convrot`. The latter two
+constructors are declared and implemented in sd.cpp and create custom nodes.
+Their former ggml declarations, implementations, ConvRot opcode, and native
+backend paths have been removed. CPU, CUDA/HIP/MUSA, and Vulkan register kernels
+and fusion rules through the extension interfaces.
+
+For already quantized activations, `ggml_ext_mul_mat_i8_tensorwise` accepts ordinary
+I8 weight and activation tensors. Its optional final `input_scale` parameter is a
+separate F32 tensor with one value per activation row. Weight scales have one
+value per output channel; weight scales and bias can be null. Existing calls
+omit `input_scale` and retain their original argument order. It computes `float(int32_dot(W, X)) * X_scale * W_scale + bias`, in that
+order. All scales are explicit graph dependencies; there is no quantized-tensor
+wrapper. Shapes use ggml's low-to-high dimension order. Call
+`sd_ggml_register_extensions` before scheduling these custom nodes.
+
+`ggml_ext_mul_mat_i8_tensorwise` also accepts F32 activations (quantized into temporary
+workspace) or the output of `ggml_ext_quantize_i8_convrot`. ConvRot keeps the existing
+packed I8 activation rows, padded to a multiple of four rows, followed by F32 row
+scales. Weight scales remain separate. CPU ConvRot groups are 1, 4, 16, 64, or 256.
+Unsupported layouts and malformed scale shapes are rejected. Separate CPU quantize
+and GEMM phases propagate cancellation, and overlapping outputs use workspace.
+
+`src/core/cuda/int8.cu` implements the same custom schemas on CUDA. It uses the
+borrowed cuBLAS handle for I8-by-I8-to-I32 GEMM and application-owned quantization,
+ConvRot, and rescaling kernels. F32 inputs are quantized into workspace; ordinary
+I8 inputs retain separate scale tensors and use a padded temporary only when
+needed by cuBLAS. Packed ConvRot inputs are consumed directly. Large ConvRot rows
+use multiple kernels when they exceed shared memory. Aliasing outputs use scratch.
+CUDA GEMM requires Turing or newer, contiguous tensors, K/N divisible by four,
+float-aligned inputs, and row counts that fit a padded signed 32-bit dimension;
+CUDA ConvRot uses groups of 256.
+Unsupported custom nodes can fall back to CPU through the scheduler.
+
+When CUTLASS is configured, `src/core/cuda/i8_gemm.cu` supplies dense and SwiGLU
+INT8 GEMM kernels on Ampere or newer. It preserves the native tile selection,
+small-row policy, Stream-K scheduling, Ada column bands, and separate scaling
+order. The same extension matches `GEMM -> SwiGLU [-> ConvRot]` and row
+concatenation for custom GEMM nodes. Weight scales remain explicit
+inputs, including the gate/up scale permutation recorded on every SwiGLU launch.
+Two independent INT8 GEMMs sharing a quantized input can also fuse with split
+SwiGLU. A paired-column iterator reads the separate gate/up weight allocations,
+and the GEMM epilogue writes only the activated output. Both intermediate
+activations are omitted through the generic fusion allocation interface; weights
+are neither copied nor concatenated. This path requires bias-free branches,
+compatible layouts and alignments, and no external intermediate consumers.
+Only interleaved weight scales and any Stream-K workspace remain temporary.
+Plans reserve bounded Stream-K workspace, cache occupancy on first execution,
+and initialize barriers on every launch, so CUDA Graph replay observes updated
+inputs and scales. External consumers retain ordinary execution; overlapping
+outputs use workspace. For fresh outputs, the CUTLASS SwiGLU epilogue writes
+an allocator-managed tensor that the subsequent ConvRot launch consumes. This
+keeps the same GPU launches while allowing its storage to be reused by other
+operations. The combined SwiGLU/ConvRot path retains its workspace intermediate
+for other cases. Shapes or alignments unsuitable for CUTLASS use cuBLAS.
+Custom I8 nodes also retain cuBLAS when the reduction could overflow I32, to
+preserve their previous accumulation behavior. Fusion disable and force-cuBLAS
+settings are honored. No CUTLASS headers are needed for the cuBLAS path.
+
+The native CUDA/HIP INT8 GEMM, standalone ConvRot, and associated fusion paths
+have been removed from ggml. The shared application implementation calls cuBLAS
+or hipBLAS directly. HIP retains the existing architecture eligibility policy;
+MUSA does not advertise INT8 BLAS support and uses scheduler fallback for GEMM.
+Standalone ConvRot kernels are compiled for all three runtimes. CUTLASS
+implementations remain CUDA-specific.
+
+`src/core/cuda/i8_gated_residual.cu` implements
+`INT8 GEMM -> [Scale] -> Mul(gate) -> Add(residual) -> [Concat]` as a registered
+CUDA fusion, for packed ConvRot inputs and ordinary I8 nodes with separate activation
+scales. The corresponding native CUDA fusion rules and CUTLASS residual epilogue
+have been removed from ggml. Its optimizer moves independent gate producers ahead
+of GEMM and optional scaling through `ggml_ext_graph_delay`, before fusion matching.
+The former native INT8 affine reorder rule has also moved into this extension.
+It preserves the original tile selection, small-row
+threshold, Stream-K scheduling, Ada column bands, and separate FP32 rounding of
+gate multiplication and residual addition. The plan bounds Stream-K workspace
+over possible occupancy values using host metadata; execution verifies the
+actual requirement against that bound. Device occupancy is cached on first
+execution; each launch rebuilds parameters for the current tensor addresses and
+initializes its Stream-K barriers. Overlapping outputs use scratch, except
+for an exact in-place residual without concatenation. External consumers and
+unsupported layouts retain ordinary operator execution. Fusion and force-cuBLAS
+disable settings are honored.
+
+`src/core/cuda/convrot.cu` registers CUDA fusion for `Scale -> ConvRot`,
+`Norm -> Mul -> ConvRot`, `Modulate -> ConvRot`,
+`RMSNorm -> Mul(weight) -> Modulate -> ConvRot`, `SwiGLU -> ConvRot`, and
+`SiLU/Sigmoid -> Mul -> ConvRot`.
+ConvRot outputs are custom nodes. SwiGLU accepts packed,
+swapped, or separate gate/up inputs. Independent modulation and up-branch
+producers are moved ahead of their consumers through the safe graph-delay API.
+Reordering keeps an owning producer together with in-place consumers, such as
+Krea2's RMSNorm weight multiplication, so the alias checks can preserve their
+storage relationship.
+The application kernel preserves the original reduction order, thread selection,
+ConvRot row padding, and separate scale tail. Prepared plans cache operation
+metadata, use the borrowed CUDA stream, and reserve scheduler workspace for
+overlapping outputs. Long rows or external intermediate consumers retain
+ordinary execution. These fusions do not require CUTLASS and honor
+`GGML_CUDA_DISABLE_FUSION`.
+
+The corresponding scale, normalization, and unary-activation fusion rules have
+been removed from ggml for CUDA and HIP. GEMM/SwiGLU/ConvRot chains are selected by the
+application INT8 extension before the standalone activation fusion.
+
+`src/core/cuda/affine.cu` registers FP32 `Mul(gate) -> Add(residual)`,
+`Concat -> [Permute] -> FP16 Copy`, and `Contiguous -> Scale -> Permute -> FP16
+Copy` fusions. Gate multiplication and residual addition retain separate FP32
+rounding. Copy layouts and alignment requirements match the original CUDA
+kernels; unsupported layouts, external intermediate consumers, self-residual
+Add inputs, and fusion disabling retain ordinary operators. Internal in-place
+Add outputs are accepted when they overwrite the complete fused Mul allocation. Exact in-place residual
+or input addresses are safe for the pointwise kernel; partial overlap, broadcast
+gate overlap, and FP16 copy overlap use scheduler workspace. These registrations
+require no CUTLASS, compile for CUDA/HIP/MUSA, and support reusable CUDA Graph
+replay. Their former native fusion implementations have been removed from ggml.
+
+CUDA kernels can opt into device replay with `ggml_ext_register_kernel_ex` and
+`GGML_EXT_KERNEL_CAPTURE_SAFE`. The original registration function retains its
+per-execution callback behavior, and the descriptor ABI is unchanged. Opting in
+means the callback enqueues replayable device work: it has no required
+per-execution host side effects or decisions based on tensor contents, and any
+private streams rejoin the borrowed stream. Lazy setup must complete during the
+first ordinary execution. Tensor metadata, attributes, and host plan parameters
+remain frozen until scheduler reallocation; tensor contents may change normally.
+
+When every selected extension opts in, the CUDA backend can capture the whole
+scheduled plan, including its native fragments, into one CUDA Graph. The first
+execution warms up kernels and library state while retaining native fragment
+caches. The next stable execution captures
+and launches the graph; later calls replay it. Changing a tensor address or the
+execution stream invalidates the capture and repeats warmup, even when the graph
+UID is unchanged. Capture failures propagate without retrying partially executed
+work. An enclosing stream capture records the plan directly instead of starting
+a nested capture.
+
+`GGML_EXT_KERNEL_REPLAY_REUSABLE` additionally permits reuse across scheduler
+plan rebuilds and requires `CAPTURE_SAFE`. Recorded work must depend only on
+complete tensor metadata, addresses, the device and immutable registration data.
+Captured resources must belong to the backend or the supplied workspace, never
+per-plan host state. Workspace contents must be initialized by recorded device
+work on each execution. Prepare/destroy still run for every rebuilt host plan.
+The backend cache retains only the executor and workspace; it compares graph
+structure, tensor metadata/addresses, selected kernels and workspace requirements
+exactly, without dereferencing old graph objects. A new metadata arena can reuse
+an equivalent capture. Each backend retains at most 16 retired plans and 256 MiB
+of their workspaces, evicting the oldest entries first. Larger workspaces retain
+plan-local capture only. Backend destruction synchronizes and releases the cache.
+This allows sd.cpp's per-step scheduler rebuilds to retain device replay.
+
+Mixed plans containing an extension without the flag retain ordinary dispatch.
+Evaluation callbacks, multiple pipeline copies, and direct backend graph calls
+also retain their previous behavior. `GGML_CUDA_DISABLE_GRAPHS` disables whole-plan
+capture as well as native CUDA Graph caching. Native operations that cannot be
+captured retain the existing fragment execution path. Whole-plan replay requires
+CUDA VMM so native scratch addresses remain valid for the backend lifetime;
+legacy scratch pools retain the existing fragment path. HIP and MUSA do not yet
+provide this plan executor. sd.cpp's CUDA registrations opt in to reusable replay.
+
+## Vulkan execution services
+
+`ggml-ext-vulkan.h` exposes `ggml_ext_vulkan_v1` through the Vulkan backend
+registry. The versioned function table reports device limits and subgroup
+capabilities, resolves tensor/workspace byte ranges, dispatches application
+SPIR-V, and copies nonoverlapping buffer ranges. Extensions do not include
+private ggml Vulkan headers or own backend command buffers and descriptors.
+
+Shaders use the `main` entry point, consecutive storage-buffer bindings starting
+at zero, and consecutive 32-bit specialization constants. Dispatch accepts actual
+workgroup counts. The backend validates ranges and dispatch limits, caches the
+complete shader/configuration per device, and releases pipelines with that
+device. Push constants and specialization values are copied; the backend does
+not retain application execution state. Pipeline creation occurs on execution,
+never during host-only plan preparation.
+
+Bindings identify either a tensor or a backend buffer, with a byte offset and
+optional byte count. `buffer_info` returns an opaque storage identity and byte
+range for alias checks. Vulkan descriptor offsets are rounded down to the
+device's required alignment; shaders must add the returned `shader_offset`.
+Workspace bindings follow the same rules. Invalid or foreign-device ranges are
+rejected before recording commands. Kernel bounds and shader capability
+requirements remain the extension's responsibility.
+
+Dispatch and copy use the backend's existing compute command context, descriptor
+pool, transfer waits, and memory barriers. They preserve ordering with native
+fragments and invalidate cached converted operands after writes. Normal backend
+synchronization completes queued work before host access or resource release.
+Vulkan copy ranges must be equal in size, nonoverlapping, and aligned to four
+bytes. There is no per-dispatch host wait or extension-owned queue.
+
+Scheduled Vulkan plans share submission counters and FLOP thresholds across
+native fragments and extension callbacks. Native scratch is initialized once
+per plan invocation, with its offsets retained across fragments. Fragment
+boundaries do not force a
+queue submission, and the final pending commands are submitted when the plan
+returns. Extension dispatches and copies count toward the device's node limit.
+Every callback still executes on every invocation; this batching neither
+captures device replay nor retains plan state across scheduler rebuilds.
+Callback errors propagate without retry. Evaluation callbacks, direct backend
+calls, synchronous devices, serialized diagnostics, and timestamp profiling
+retain fragment execution.
+
+`src/core/vulkan/rope.cpp` registers the custom RoPE implementation,
+plus `RMSNorm -> Mul -> RoPE` and `RMSNorm -> Mul -> RoPE -> Scale -> FP16 Copy`
+fusion. These kernels retain the original strides, partial-rotation behavior,
+128-channel subgroup reduction, and separate arithmetic rounding. Their SPIR-V
+is generated from application shaders with the native RoPE compilation policy;
+SPIR-V optimization is disabled because it changes the reduction results.
+Safe graph-delay rules bring independent producers ahead of fusion candidates.
+External intermediate consumers retain ordinary operators, and overlapping
+outputs use plan workspace. `GGML_VK_DISABLE_FUSION` is honored.
+
+The corresponding Vulkan shader sources and native fusion rules have been
+removed from ggml. Unsupported custom nodes retain scheduler fallback to
+another supported backend.
+
+RoPE Apply, INT8/ConvRot, CuTe attention, SageAttention, and Sol-Attn use
+application-owned implementations. This interface does not add dynamic dtypes.
+
+`src/core/vulkan/int8.cpp` owns Vulkan INT8 GEMM, ConvRot, and their fusion rules.
+The Vulkan compute metadata service exposes shared-memory limits, subgroup
+arithmetic, and cooperative-matrix properties; shape selection stays in sd.cpp.
+Shader variants preserve the former tiling and compilation policy, including
+SwiGLU, residual, row-concatenation, and modulation paths. Kernels use public
+Vulkan dispatch, buffer, copy, and workspace services without private backend
+headers. Vulkan currently handles packed group-256 activations with separate
+weight scales; unsupported custom nodes fall back to CPU.
+
+## SageAttention
+
+The `ggml_ext_sage_attn` entry point is declared and implemented in sd.cpp's
+graph construction helpers. It constructs `sd.sage_attn.v1` with explicit Q/K/V inputs
+and scale/mode attributes. ggml has no Sage-specific opcode, allocation rule, or
+backend dispatch branch.
+
+`src/core/cuda/sage_attention.cu` and its `sageattention/` device headers own the
+CUDA implementation. Quantization and temporary attention output use scheduler
+workspace, rather than an oversized output tensor. The implementation obtains
+the backend stream through the public GPU service and converts the final FP16
+output with a local CUDA kernel. It retains the SM80 compatibility path and the
+SageAttention2/2++ toolkit and architecture requirements. Capability checks use
+the application's compiled CUDA target as well as the physical device.
+
+CPU, HIP, MUSA, and Vulkan do not implement this NVIDIA-specific kernel. Existing
+model-level capability checks retain flash/default attention for unsupported
+layers; a custom Sage node itself does not silently change algorithms.
+
+## Sol-Attn
+
+The `ggml_ext_sol_attn` entry point lives in the common sd.cpp graph construction
+files and constructs `sd.sol_attn.v1`. It accepts contiguous F32 Q/K/V tensors
+with shape `[128, tokens, heads, batch]`, a positive finite scale, and finite tau.
+The output remains F32 `[128, heads, tokens, batch]`. No current sd.cpp model or
+CLI path calls this entry point; the migration preserves its low-level behavior.
+
+`src/core/cuda/sol_attention.cu` owns packing, block summaries, threshold
+selection, and the BF16 tensor-core kernel. Temporary storage is separate
+scheduler workspace, so the output allocation contains only the output tensor.
+The kernel uses the public backend stream service and direct CUDA/PTX operations.
+Registration requires CUDA 12 or newer, an SM80-or-newer device, and an application
+kernel compiled for SM80 or newer. CPU, HIP, MUSA, and Vulkan do not advertise
+support. ggml no longer contains a Sol-Attn opcode, constructor, allocation rule,
+or backend kernel.
+
+## Public graph access
+
+`ggml-graph.h` exposes leaf enumeration, append-only leaf construction, recorded
+use counts, mapped inference-graph copying, use-count copying for segments, and
+backend execution of a node range. Mapped copies retain node order and consumer
+counts from the original graph, including consumers outside a graph view. The
+caller owns tensor copies and remaps their storage and edges; invalid mappings
+leave the destination graph unchanged. Range execution uses already allocated
+tensors and retains the original graph's consumer counts.
+
+Workspace measurement, graph cuts, and evaluation callbacks use these APIs.
+Normal sd.cpp builds require only ggml's public headers. Compatibility with an
+upstream ggml lacking this header is isolated in `src/core/ggml_compat.cpp`; that
+optional build still requires matching upstream private headers. The compatibility
+layer also retains the workaround for old Metal backends wrapping failed host
+mappings. The patched Metal backend returns null directly when mapping fails.
@@ -12,7 +12,7 @@
 ### HiDream-O1-Image-Dev

 ```
-.\bin\Release\sd-cli.exe -m  ..\..\ComfyUI\models\diffusion_models\hidream_o1_image_dev_bf16.safetensors -p "a lovely cat holding a sign says 
+.\bin\Release\sd-cli.exe -m  ..\models\diffusion_models\hidream_o1_image_dev_bf16.safetensors -p "a lovely cat holding a sign says 
 'hidream o1 cpp'" --cfg-scale 1.0  -v -H 1024 -W 1024
 ```

@@ -0,0 +1,24 @@
+# HunyuanVideo 1.5
+
+HunyuanVideo 1.5 uses a HunyuanVideo diffusion transformer, a causal video VAE, Qwen2.5-VL 7B for the main text conditioning, 
+and ByT5 Small GlyphXL for glyph-aware text conditioning.
+
+## Download weights
+
+- Download HunyuanVideo 1.5
+    - safetensors: https://huggingface.co/Comfy-Org/HunyuanVideo_1.5_repackaged/tree/main/split_files/diffusion_models
+- Download vae
+    - safetensors: https://huggingface.co/Comfy-Org/HunyuanVideo_1.5_repackaged/tree/main/split_files/vae
+- Download qwen_2.5_vl 7b
+    - safetensors: https://huggingface.co/Comfy-Org/Qwen-Image_ComfyUI/tree/main/split_files/text_encoders
+    - gguf: https://huggingface.co/mradermacher/Qwen2.5-VL-7B-Instruct-GGUF/tree/main
+- Download byt5 small glyphxl
+    - safetensros: https://huggingface.co/Comfy-Org/HunyuanVideo_1.5_repackaged/tree/main/split_files/text_encoders
+
+## Text-to-video example
+
+```shell
+.\bin\Release\sd-cli.exe -M vid_gen --diffusion-model  ..\models\diffusion_models\hunyuanvideo1.5_720p_t2v_fp16.safetensors --vae ..\models\vae\hunyuanvideo15_vae_fp16.safetensors --llm ..\models\text_encoders\qwen_2.5_vl_7b.safetensors --t5xxl ..\models\text_encoders\byt5_small_glyphxl_fp16.safetensors  -p "a lovely cat" --cfg-scale 6.0 --sampling-method euler -v -W 1280 -H 720 --offload-to-cpu --video-frames 33 --vae-tiling
+```
+
+<video src=../assets/hunyuan_video/hy1.5_t2v.mp4 controls="controls" muted="muted" type="video/mp4"></video>
@@ -0,0 +1,55 @@
+# How to Use
+
+## Download weights
+
+- Download Ideogram4
+    - safetensors: https://huggingface.co/ideogram-ai/ideogram-4-fp8/tree/main/transformer
+- Download Ideogram4 uncond
+    - safetensors: https://huggingface.co/ideogram-ai/ideogram-4-fp8/tree/main/unconditional_transformer
+- Download vae
+    - safetensors: https://huggingface.co/black-forest-labs/FLUX.2-dev/tree/main
+- Download Qwen3-VL-8B-Instruct
+    - gguf: https://huggingface.co/unsloth/Qwen3-VL-8B-Instruct-GGUF/tree/main
+
+## Use original FP8 weights
+
+The original Ideogram4 FP8 safetensors can be loaded directly. FP8 tensors stay
+at one byte per element in RAM and VRAM. Backends that cannot multiply FP8
+weights directly cast only the active layer to a temporary BF16 tensor during
+execution; the loader does not expand the entire checkpoint to BF16.
+
+With `SD_USE_UPSTREAM_GGML=ON`, FP8 tensors are converted to F16 at load time
+instead (two bytes per element in RAM and VRAM).
+
+Use `ideogram4_fp8.safetensors` and `ideogram4_uncond_fp8.safetensors` directly
+with `--diffusion-model` and `--uncond-diffusion-model`, respectively.
+
+## Optional conversion for quantization
+
+The following conversion is only needed when creating a quantized GGUF model.
+
+fp8 scale -> bf16
+
+```
+python .\convert_fp8_scale_to_bf16.py --input .\ideogram4_fp8.safetensors --output ideogram4_bf16.safetensors
+python .\convert_fp8_scale_to_bf16.py --input .\ideogram4_uncond_fp8.safetensors --output ideogram4_uncond_bf16.safetensors
+```
+
+bf16 -> q8
+
+```
+.\bin\Release\sd-cli.exe -M convert -m ideogram4_bf16.safetensors -o ideogram4-Q8_0.gguf --tensor-type-rules "^layers.*adaln_modulation.*weight=q8_0,layers.*attention.o.*weight=q8_0,layers.*attention.qkv.*weight=q8_0,layers.*feed_forward.*weight=q8_0" -v
+
+.\bin\Release\sd-cli.exe -M convert -m ideogram4_uncond_bf16.safetensors -o ideogram4_uncond-Q8_0.gguf --tensor-type-rules "^layers.*adaln_modulation.*weight=q8_0,layers.*attention.o.*weight=q8_0,layers.*attention.qkv.*weight=q8_0,layers.*feed_forward.*weight=q8_0" -v
+```
+
+If you want lower VRAM usage, you can change the quantization from q8_0 to a lower-level quantization, such as q4_0.
+
+
+## Examples
+
+```sh
+.\bin\Release\sd-cli.exe --diffusion-model ideogram4-Q8_0.gguf --uncond-diffusion-model ideogram4_uncond-Q8_0.gguf --llm ..\models\text_encoders\Qwen3VL-8B-Instruct-Q4_K_M.gguf --vae ..\models\vae\flux2_ae.safetensors -p '{"high_level_description":"A square 1024 x 1024 luxury fashion magazine cover featuring exactly one short chubby fluffy cat as the main model. The cat sits on a soft ivory studio floor, facing the viewer with a stylish calm expression, wearing tiny black sunglasses, a red silk scarf, and a small gold collar charm. In front of the cat on the floor is a wide horizontal luxury nameplate that clearly reads ideogram4.cpp. The whole design feels premium, fashionable, clean, and editorial.","style_description":{"aesthetics":"luxury fashion magazine cover, high-end pet couture campaign, minimalist editorial design, elegant studio photography, soft paper texture, refined typography, fashionable and polished","lighting":"Soft diffused studio lighting, gentle spotlight on the cat, subtle floor shadow, warm ivory highlights, clean separation between subject and background","photo":"high-resolution fashion editorial photography look, front-facing cat portrait, crisp fur details, glossy sunglasses, clear readable nameplate text, shallow depth of field","medium":"mixed media fashion photography and premium editorial graphic design","color_palette":["#F4EFE7","#111111","#D8B56D","#B73A3A","#FFFFFF","#8A7A6A"]},"compositional_deconstruction":{"canvas":"Square 1024 x 1024 canvas with a normal upright orientation. Do not rotate the poster or any text. Use a clean fashion magazine cover layout.","background":"Warm ivory studio backdrop with subtle paper grain, a soft spotlight gradient, faint floor shadow, and a few minimal gold editorial lines. The background is spacious, premium, and uncluttered.","layout":"Top center has a small elegant headline. Center area features one cat as the main fashion model. Lower foreground has a wide horizontal luxury nameplate placed on the floor in front of the cat. Bottom center has a small footer. All text is horizontal, upright, and readable left to right.","elements":[{"type":"text","desc":"Top center headline reading LOOK WHAT I FOUND in a refined high-fashion serif font. The headline is horizontal, centered, elegant, and secondary to the nameplate text."},{"type":"obj","desc":"Exactly one short chubby fluffy cat sitting in the center like a luxury fashion model. The cat has a large round head, compact body, short legs, soft detailed fur, expressive eyes, and a calm confident pose. The cat is cute and rounded, not tall, not stretched, not duplicated."},{"type":"obj","desc":"Tiny glossy black sunglasses worn naturally by the cat, slightly oversized but still showing the cat face clearly. The sunglasses add a chic fashion-editorial attitude."},{"type":"obj","desc":"A red silk scarf tied neatly around the cat neck, with soft folds and a couture feeling. The scarf must not cover the cat face or the nameplate."},{"type":"obj","desc":"A small gold collar charm or fashion accessory under the scarf, subtle and premium, adding a luxury campaign detail."},{"type":"obj","desc":"In the lower foreground, place a wide horizontal luxury nameplate on the floor in front of the cat. The nameplate is low, flat, landscape-oriented, much wider than tall, like a fashion show seat card or premium display plaque. It is centered, front-facing, level, and fully visible. It must not become vertical, tall, standing, rotated, or side-facing."},{"type":"text","desc":"Print the exact text ideogram4.cpp only on the wide horizontal nameplate. Use clean bold black lettering, perfectly spelled, lowercase, with the number 4 and .cpp extension. The text must fit completely inside the nameplate, stay horizontal, and be readable from left to right."},{"type":"obj","desc":"Add sparse premium editorial accents around the edges: thin gold lines, small code brackets, tiny cursor marks, subtle dots, and minimal geometric details. No extra cats, no stickers, no animal faces, no busy decorations."},{"type":"text","desc":"Bottom center footer reading tiny paws, big compile energy in a small refined monospace or editorial font. The footer is horizontal, centered, understated, and much smaller than the nameplate text."}]}}' -v --offload-to-cpu -H 1024 -W 1024
+```
+
+<img alt="ideogram4 image example" src="../assets/ideogram4/example.png" />
@@ -0,0 +1,180 @@
+# Image preprocessing
+
+Use `--image-preprocess` to transform each image input once, before generation:
+
+```sh
+sd-cli ... \
+  --image-preprocess "target=init,mode=crop-resize,filter=lanczos,antialias=true" \
+  --image-preprocess "target=mask,filter=nearest-exact" \
+  --image-preprocess "target=ref,index=0,mode=fit-pad,width=768,height=768,filter=bicubic"
+```
+
+CLI and server image loaders decode at the original resolution. The generation
+entry point merges input defaults with user rules and prepares one transformed
+image per input. The original pipeline then consumes those images, including
+its mandatory canvas adaptation, reference resizing, and encoder preprocessing.
+
+```text
+native-resolution image
+  -> input defaults + user overrides
+  -> one input transform
+  -> original generation pipeline and model-specific processing
+```
+
+These rules do not override internal VAE, CLIP/VLM, ControlNet, or pixel-patch preprocessing.
+`--ref-image-args` retains its existing meaning and runs after this input transform.
+
+## Inputs and defaults
+
+| `target` | Input | Default geometry | Indexed? |
+| --- | --- | --- | --- |
+| `init` | img2img image or video first frame | Center crop to the generation aspect ratio, then resize | No |
+| `end` | Video last frame | Center crop, then resize | No |
+| `mask` | Inpainting mask | Inherit init geometry; otherwise center crop, then resize | No |
+| `control` | Control image | Center crop, then resize | No |
+| `ref` | Reference images | Preserve source dimensions | Yes |
+| `ip-adapter` | IP-Adapter image | Preserve source dimensions | No |
+| `id` | PhotoMaker identity images | Preserve source dimensions | Yes |
+| `control-frame` | Control video frames | Center crop, then resize | Yes |
+
+Canvas defaults use the aligned generation dimensions. Reference, IP-Adapter,
+and identity inputs use their original dimensions unless overridden. Default
+resampling is nearest for images and nearest-exact for masks.
+
+These defaults are shared by CLI, server, and C API. Moving geometry out of
+the loaders replaces the previous CLI/server BOX/sRGB resizing, so default
+pixels are not guaranteed to match earlier builds.
+
+Reference video and audio preprocessing are outside these image rules.
+Preprocessing options apply to `img_gen` and `vid_gen`, not standalone upscale
+or ADetailer mode. ADetailer clears the user's rules for its internal crops.
+
+## Rules
+
+Rules are comma-separated `key=value` lists. Repeat the CLI option or separate
+rules with semicolons. Every rule requires a `target` and at least one option.
+Rule syntax and input compatibility are checked when image/video generation
+starts. Unknown keys, invalid values, duplicate keys in a rule, missing images,
+and out-of-range indices cause generation to fail with an error log.
+
+Omit `index` to configure every image of that type; otherwise use a zero-based
+index. CLI directory inputs follow filename order. Indexed rules override
+type-wide rules field by field, regardless of order. At equal specificity,
+the last value for a field wins. `auto` selects the input preset.
+
+| `mode` | Input transform |
+| --- | --- |
+| `auto` | Use the input's default geometry |
+| `none` | Keep source dimensions without resizing, cropping, or padding |
+| `stretch` | Resize to the target dimensions |
+| `crop` | Crop a target-sized rectangle without resizing; fail if the source is too small |
+| `crop-resize` | Crop to the target aspect ratio, then resize |
+| `fit-pad` | Fit the entire image inside the target dimensions, preserving aspect ratio, then pad |
+
+`width` and `height` must be specified together as positive integers. They
+override the input transform's dimensions, not the generation or encoder size.
+For a native-size preset, specifying dimensions without a mode selects stretch.
+`mode=none` with explicit dimensions different from the source is contradictory
+and is rejected.
+
+`anchor=center|top|bottom|left|right` selects crop/padding placement.
+`pad_color=#RRGGBB` or `#RRGGBBAA` selects padding, defaulting to opaque black.
+A grayscale mask uses the first color component.
+
+`filter=auto|nearest|nearest-exact|bilinear|bicubic|lanczos` selects resampling.
+`antialias=auto|true|false` enables antialiasing automatically for filtered
+downscaling; explicit true requires bilinear, bicubic, or Lanczos.
+Filtered RGBA resizing uses premultiplied alpha.
+
+`canny=true|false` enables edge detection for any supported image target,
+defaulting to `false`. It runs once after geometry, before the original
+generation pipeline, including with `mode=none`. Grayscale, grayscale-alpha,
+RGB, and RGBA inputs are supported; alpha is preserved.
+
+Each input has its own Canny setting. Indexed rules can enable or disable it
+for individual references, identity images, or video control frames.
+
+```sh
+--image-preprocess "target=init,mode=fit-pad,canny=true"
+--image-preprocess "target=ref,index=0,mode=none,canny=true"
+--image-preprocess "target=control-frame,index=2,canny=true"
+```
+
+If init and mask source dimensions differ, the mask is first center-cropped to
+the init source aspect ratio and resized to the init source dimensions using
+nearest-exact interpolation. The mask then inherits the init crop, resize, and
+padding coordinates, while retaining its own filter, padding value, and Canny
+setting for that transform. Explicit mask geometry is evaluated after this
+alignment; conflicting geometry is rejected. An omitted mask remains absent
+until the original pipeline creates its default mask.
+
+## Downstream behavior
+
+`mode=none` only skips the input geometry transform. For example:
+
+```sh
+--image-preprocess "target=init,mode=none" \
+--image-preprocess "target=ref,mode=none"
+```
+
+The init image is still adapted to the generation canvas by the original
+pipeline. Reference images still follow `--ref-image-args` and model-specific
+resizing. CLIP retains its fixed input dimensions and normalization. HiDream-O1
+retains its original pixel-reference and visual preprocessing.
+
+Existing sharing between consumers is preserved: for example, Wan img2video
+uses the same adapted first frame for VAE conditioning and CLIP. High-resolution
+passes reuse the prepared images and apply their original size adaptation;
+they do not apply the user's crop a second time.
+
+To disable reference resizing before VAE encoding, use
+`--ref-image-args "resize_before_vae=false"` or the server field
+`"ref_image_args": "resize_before_vae=false"`. This is separate from
+`target=ref,mode=none`, which only skips input geometry. Model constraints
+still apply; MiniMax-H3 requires 32-pixel-aligned references and warns instead
+of honoring the switch.
+
+## Server requests
+
+Native image/video requests and SDAPI accept `image_preprocess` as a string or
+an array of rule strings:
+
+```json
+{
+  "image_preprocess": [
+    "target=init,mode=fit-pad,filter=bicubic",
+    "target=mask,filter=nearest-exact",
+    "target=ref,index=0,mode=none"
+  ]
+}
+```
+
+OpenAI-compatible requests accept it through
+`<sd_cpp_extra_args>{...}</sd_cpp_extra_args>` in the prompt.
+Request rules replace server-default rules. Generation metadata records the
+user rules for replay with the same input images; image encodings and channel conventions are unchanged.
+
+## C API
+
+Set `image_preprocess` on the existing image/video generation parameters.
+Generation returns media and actual parameters together; see [generation metadata](generation_metadata.md):
+
+```c
+sd_img_gen_params_t params;
+sd_img_gen_params_init(&params);
+/* Set prompt, original-resolution input images, and generation options. */
+params.image_preprocess.rules = "target=init,mode=crop-resize,filter=lanczos;"
+                                "target=mask,filter=nearest-exact";
+sd_image_generation_result_t* result = NULL;
+bool ok = generate_image(ctx, &params, &result);
+/* Consume result->items and their generation_info on success. */
+free_sd_image_generation_result(result);
+```
+
+Both generation parameter initializers set `image_preprocess.rules` to `NULL`,
+selecting input presets. Rule strings are borrowed for the synchronous call.
+The library owns temporary transformed pixels; caller images and arrays are
+not modified. Add `canny=true` to the desired target's rule in
+`image_preprocess.rules` to enable Canny.
+
+The parameter structs have grown; applications and bindings must be rebuilt.
@@ -0,0 +1,59 @@
+# Importance Matrix (imatrix) Quantization
+
+## What is an Importance Matrix?
+
+Quantization reduces the precision of a model's weights, decreasing its size and computational requirements. However, this can lead to a loss of quality. An importance matrix helps mitigate this by identifying which weights are *most* important for the model's performance. During quantization, these important weights are preserved with higher precision, while less important weights are quantized more aggressively.  This allows for better overall quality at a given quantization level.
+
+This originates from work done with language models in [llama.cpp](https://github.com/ggml-org/llama.cpp/blob/master/tools/imatrix/README.md).
+
+## Usage
+
+The imatrix feature involves two main steps: *training* the matrix and *using* it during quantization.
+
+### Training the Importance Matrix
+
+To generate an imatrix, run stable-diffusion.cpp with the `--imat-out` flag, specifying the output filename.  This process runs alongside normal image generation.
+
+```bash
+sd.exe [same exact parameters as normal generation] --imat-out imatrix.dat
+```
+
+*   **`[same exact parameters as normal generation]`**:  Use the same command-line arguments you would normally use for image generation (e.g., prompt, dimensions, sampling method, etc.).
+*   **`--imat-out imatrix.dat`**: Specifies the output file for the generated imatrix.
+
+You can generate multiple images at once using the `-b` flag to speed up the training process.
+
+### Continuing Training an Existing Matrix
+
+If you want to refine an existing imatrix, use the `--imat-in` flag *in addition* to `--imat-out`. This will load the existing matrix and continue training it.
+
+```bash
+sd.exe [same exact parameters as normal generation] --imat-out imatrix.dat --imat-in imatrix.dat
+```
+With that, you can train and refine the imatrix while generating images like you'd normally do. 
+
+### Using Multiple Matrices
+
+You can load and merge multiple imatrices together:
+
+```bash
+sd.exe [same exact parameters as normal generation] --imat-out imatrix.dat --imat-in imatrix.dat --imat-in imatrix2.dat
+```
+
+### Quantizing with an Importance Matrix
+
+To quantize a model using a trained imatrix, use the `-M convert` option (or equivalent quantization command) and the `--imat-in` flag, specifying the imatrix file.
+
+```bash
+sd.exe -M convert [same exact parameters as normal quantization] --imat-in imatrix.dat
+```
+
+*   **`[same exact parameters as normal quantization]`**: Use the same command-line arguments you would normally use for quantization (e.g., target quantization method, input/output filenames).
+*   **`--imat-in imatrix.dat`**: Specifies the imatrix file to use during quantization.  You can specify multiple `--imat-in` flags to combine multiple matrices.
+
+## Important Considerations
+
+*   The quality of the imatrix depends on the prompts and settings used during training. Use prompts and settings representative of the types of images you intend to generate for the best results.
+*   Experiment with different training parameters (e.g., number of images, prompt variations) to optimize the imatrix for your specific use case.
+*   The performance impact of training an imatrix during image generation or using an imatrix for quantization is negligible.
+*   Using already quantized models to train the imatrix seems to be working fine.
@@ -0,0 +1,173 @@
+# INT8 Convrot Safetensors
+
+sd.cpp can load and execute ComfyUI `int8_tensorwise` safetensors with `convrot` metadata directly. The stored INT8 weights are not converted to another weight type at load time.
+
+This uses sd.cpp's INT8 tensorwise/convrot operators through the patched GGML
+extension API. Schemas and graph builders live in sd.cpp, alongside the
+backend kernels and fusion rules; GGML supplies generic dispatch and workspace
+services. See [operator extensions](ggml_extensions.md) for the implementation
+and the preserved INT8 API entry points.
+Builds with `SD_USE_UPSTREAM_GGML=ON` reject these files during loading.
+
+## Checkpoint format
+
+Each quantized linear module contains the following tensors:
+
+- `<module>.weight`: an I8 weight matrix.
+- `<module>.weight_scale`: one floating-point scale for each output row. ComfyUI's two-dimensional `[out_features, 1]` representation is normalized to a one-dimensional tensor while loading.
+- `<module>.comfy_quant`: a U8 tensor containing the JSON quantization configuration.
+
+A supported configuration has this form:
+
+```json
+{
+  "format": "int8_tensorwise",
+  "convrot": true,
+  "convrot_groupsize": 256
+}
+```
+
+The convrot group size must be a power of four and must divide the input feature dimension. The commonly used configuration is H256, with `convrot_groupsize` set to `256`.
+
+## How INT8 convrot works
+
+Convrot combines an offline rotation of the weights with the same rotation of the activations at runtime. The rotation uses a normalized regular Hadamard matrix constructed recursively from
+
+```text
+     [ 1  1  1 -1 ]
+     [ 1  1 -1  1 ]
+H4 = [ 1 -1  1  1 ] / 2
+     [-1  1  1  1 ]
+```
+
+For a group size `G = 4^n`, the transform is the normalized Kronecker power of `H4`. It is applied independently to every contiguous group of `G` input features. The resulting block-diagonal rotation matrix `R` is orthogonal and symmetric, so `R R^T = I`.
+
+For an original floating-point linear layer
+
+```text
+Y = X W^T + b
+```
+
+the checkpoint stores a rotated weight matrix `W_rot = W R`, quantized per output row. At runtime sd.cpp computes `X_rot = X R`. Ignoring quantization error,
+
+```text
+X_rot W_rot^T = X R (W R)^T = X R R^T W^T = X W^T
+```
+
+The rotation therefore preserves the linear operation. Its purpose is to spread isolated large values across each feature group, reducing the effect of outliers on tensorwise INT8 quantization.
+
+### Weight quantization
+
+The rotated weights are quantized offline with one scale per output row:
+
+```text
+s_w[o]      = max_i(abs(W_rot[o, i])) / 127
+Q_w[o, i]   = clamp(round(W_rot[o, i] / s_w[o]), -127, 127)
+```
+
+`Q_w` is stored in `<module>.weight`, and `s_w` is stored in `<module>.weight_scale`.
+
+### Runtime activation quantization
+
+For every activation row, sd.cpp applies the group-wise Hadamard rotation and then calculates one dynamic scale across the entire rotated row:
+
+```text
+s_x[r]      = max_i(abs(X_rot[r, i])) / 127
+Q_x[r, i]   = clamp(round(X_rot[r, i] / s_x[r]), -127, 127)
+```
+
+The matrix multiplication accumulates into signed 32-bit integers:
+
+```text
+A[r, o] = sum_i(Q_x[r, i] * Q_w[o, i])
+```
+
+The floating-point output is reconstructed as
+
+```text
+Y[r, o] ~= A[r, o] * s_x[r] * s_w[o] + b[o]
+```
+
+The packed runtime activation tensor contains the I8 activation rows and their floating-point row scales. Linear layers with an input scale of `1` that share the same input and convrot group size reuse this packed tensor, avoiding repeated rotation and activation quantization within the graph. For other input scales, activations are scaled before explicit convrot quantization, and the linear output is unscaled before adding bias.
+
+## Backend support
+
+- CPU provides the portable regular Hadamard, activation quantization, INT8 matrix multiplication, and scale restoration implementations.
+- Metal supports INT8 tensorwise and convrot on Apple silicon (Apple GPU family 7 or newer). It keeps weights in INT8 and rotates and quantizes activations on the GPU. On macOS/iOS 26 or newer, it probes Metal TensorOps `matmul2d` with INT8 inputs and INT32 cooperative results. It combines 1024-product partial sums using wrapping integer addition to preserve exact accumulation on older GPUs, then applies scales and bias in registers. Older systems or unavailable TensorOps pipelines use SIMD-group matrix multiplication: INT8 tiles are converted to FP16, and FP32 partial sums cover at most 256 products before being combined as integers, preserving the I32 accumulation contract. Convrot groups 1, 4, 16, 64, and 256 are supported. These kernels require the patched GGML Metal extension services and are enabled by `SD_METAL=ON`.
+- NVIDIA CUDA devices with compute capability 7.5 or newer use the native accelerated path. For H256, CUDA fuses the rotation, row-wise maximum reduction, and activation quantization. Eligible shapes use the CUTLASS INT8 implementation when enabled. For linear layers, it restores scales and adds any aligned bias before writing the result; other supported shapes use cuBLAS for I8 x I8 to I32 GEMM and a CUDA kernel for scale restoration and bias addition.
+- AMD HIP devices in the CDNA, RDNA3 (including RDNA3.5), and RDNA4 families use the same INT8/H256 kernels with hipBLAS for I8 x I8 to I32 GEMM. Other AMD architectures fall back to CPU for INT8 matrix multiplication.
+- MUSA shares the CUDA-style H256 activation kernels. INT8 GEMM falls back to CPU because the MUSA backend does not advertise INT8 BLAS support.
+- Vulkan provides H256 activation quantization and INT8 matrix multiplication when the build and device support accelerated packed INT8 dot products. Unsupported configurations and GPU backends without these kernels use the backend scheduler to fall back to CPU.
+
+LoRA adapters are applied at runtime without modifying the INT8 weights. The INT8 convrot path computes the base linear output, while LoRA, LoHa, LoKr, and raw weight-difference adapters compute their output corrections from the original, unrotated activation and add them to the base output. `--lora-apply-mode auto` selects this path for models containing INT8 tensorwise weights. If `immediately` is requested, sd.cpp falls back to runtime application because merging an adapter would require dequantizing and rotating its weight update, then recalculating the per-row scales and requantizing the result.
+
+The shared CUDA/HIP/MUSA convrot activation paths support H64 and H256. Vulkan also supports H64 activation quantization; its INT8 matrix multiplication path currently requires H256.
+
+## Packed INT4 convrot
+
+ComfyUI `convrot_w4a4` and `asym_w4a8_int8` safetensors load directly, including
+checkpoints that mix packed INT4, INT8 tensorwise, and floating-point layers.
+Packed weights retain their checkpoint byte layout; they are not expanded or
+converted to GGUF Q4 blocks at load time.
+
+- `convrot_w4a4`: signed 4-bit weights, low nibble first, with F32 per-output
+  `weight_scale`. Activations use the same INT8 convrot quantization as above;
+  the format name does not select a native INT4 activation kernel.
+- `asym_w4a8_int8`: 4-bit indices into `weight_codebook` (16 F32 values), F32
+  `weight_s_channel`, and F8_E4M3 `weight_s_rel` per 16 input features. The
+  metadata must specify `group_size: 16`. CPU evaluates the floating-point
+  codebook; GPU kernels requantize it to INT8 with `amax / 127`, introducing a
+  small additional rounding error.
+
+Both formats imply convrot rotation even without a `convrot` flag; missing
+`convrot_groupsize` defaults to 256.
+
+Both formats reuse the convrot activation cache when the input scale is `1` and
+produce F32 outputs with optional bias. CPU provides a portable implementation.
+CUDA/HIP/MUSA use shared packed-weight kernels without requiring INT8 BLAS.
+With CUTLASS enabled, NVIDIA
+SM80 and newer devices use mixed-input S8 x S4 GEMM for eligible W4A4 shapes,
+reusing the INT8 scale epilogue without expanding the full weight tensor. W4A8
+uses CUTLASS integer MMA per 16-value group with FP32 scale restoration.
+Other CUDA shapes and builds retain the packed byte-dot path. Set
+`GGML_CUDA_W4_FORCE_DP4A=1` to compare against that path. Supported AMD targets
+use packed byte-dot instructions, with portable byte arithmetic for other targets.
+Vulkan requires integer dot product support and at least 13312 bytes of shared
+memory for W4A4 or 16452 bytes for W4A8. GPU W4 kernels require the logical input
+width to be divisible by 64; H64 and H256 are supported. When supported by the
+compiler and device, Vulkan uses NV cooperative matrix integer GEMM. W4A4
+decodes packed weights directly into cooperative matrices for aligned input
+and output widths. Complete activation tiles and the final partial tile use
+separate shader specializations. Other W4A4 shapes use double-buffered weight
+and activation tiles (20480 bytes of shared memory, or 30720 for larger tiles).
+W4A8 stages both weights and activations in shared memory in K=64 tiles,
+reusing them for each 16-value group's integer
+MMA and FP32 scale restoration. FP8 scales use a shared lookup table, and
+expanding projections visit activation tiles in groups of weight panels for
+cache reuse. Adjacent groups share a K=32 activation fragment, with
+complementary zero padding in the weights to preserve each group's scale.
+This path uses 17476 bytes of shared memory, or 22596 for larger tiles.
+Both paths use 256-thread workgroups and keep the model's
+weights packed; neither expands a full layer in device memory. Other Vulkan
+devices retain the packed dot-product path; set `GGML_VK_W4_FORCE_DP4A=1` to compare
+against it. FP32 accumulation order can differ between these paths.
+Other shapes use CPU fallback. HIP W4A8 selects wave64 WMMA by the active
+compilation target's intrinsic support, preferring the GFX12 fragment layout
+over the GFX11 layout.
+There is no GPU architecture allowlist: newer targets supporting either
+intrinsic automatically use its matching layout. Selection happens per device
+image, including multi-architecture builds; targets supporting neither intrinsic
+retain the integer-dot fallback. W4A4 retains integer dot products on HIP.
+
+LoRA application remains at runtime, using the original unrotated activations
+and logical weight dimensions. Packed weights and their codebook/scale tensors
+are excluded from weight type conversion. `SD_USE_UPSTREAM_GGML` builds reject
+these formats because they require the application extension API.
+
+## Example
+
+ComfyUI INT8 convrot safetensors can be passed to `--diffusion-model` without conversion:
+
+```powershell
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\krea2_turbo_int8_convrot.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf --vae ..\models\vae\wan_2.1_vae.safetensors -p "a lovely cat holding a sign says 'krea2.cpp'" --steps 8 --cfg-scale 1 -v --offload-to-cpu
+```
@@ -0,0 +1,87 @@
+# 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.
+
+Both the classic adapters and the higher-fidelity **Plus** adapters are
+supported; see [Plus variants](#plus-variants) below. The variant is
+detected from the weight file, so the same options work for both.
+
+## 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`
+   - SD 1.5 Plus: `models/ip-adapter-plus_sd15.safetensors`
+   - SDXL Plus: `sdxl_models/ip-adapter-plus_sdxl_vit-h.safetensors`
+
+   The Plus files (`ip-adapter-plus_*`) are used exactly like the classic
+   ones; see [Plus variants](#plus-variants).
+
+## 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 --vae-tiling
+```
+
+The SDXL VAE decode at 1024x1024 is memory heavy; add `--vae-tiling` (and
+`--offload-to-cpu`) on GPUs with limited VRAM.
+
+## Plus variants
+
+The Plus adapters (`ip-adapter-plus_sd15`, `ip-adapter-plus_sdxl_vit-h`)
+replace the small linear image projection with a Resampler (a
+Perceiver-style module with learned latent queries). Instead of pooling the
+CLIP-Vision output into one vector, the Resampler attends over the full grid
+of penultimate CLIP-Vision hidden states and emits more image tokens (16
+instead of 4). The result transfers finer detail and layout from the
+reference, at a small extra cost in the image-projection step.
+
+No extra flags are needed. The variant is detected from the weight file (the
+Resampler's `image_proj.latents` tensor), and every Resampler dimension is
+read from the tensor shapes, so the same `--ip-adapter`,
+`--ip-adapter-image`, and `--ip-adapter-strength` options apply. Plus
+composes with Control Net in the same way as the classic adapters.
+
+```
+sd-cli -m ..\models\sd_v1.5.safetensors --clip_vision ..\models\clip_vision_h.safetensors --ip-adapter ..\models\ip-adapter-plus_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
+```
+
+The startup log line `IP-Adapter: 16 image tokens` (versus `4` for the
+classic adapters) confirms a Plus file was loaded.
+
+## 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 --vae-tiling
+```
@@ -0,0 +1,24 @@
+# How to Use
+
+Iris-3B generates images directly in pixel space and uses Qwen3-VL-4B-Instruct as the text encoder. No VAE is required.
+
+## Download weights
+
+- Download Iris-3B
+    - safetensors: https://huggingface.co/speridlabs/iris-3b/tree/main (`model.safetensors` in the root directory)
+- Download Qwen3-VL-4B-Instruct
+    - safetensors: https://huggingface.co/Comfy-Org/Krea-2/tree/main/text_encoders
+    - gguf: https://huggingface.co/Qwen/Qwen3-VL-4B-Instruct-GGUF/tree/main
+
+## Examples
+
+```
+.\bin\Release\sd-cli.exe --diffusion-model ..\models\diffusion_models\iris-3b.safetensors --llm ..\models\text_encoders\Qwen3-VL-4B-Instruct-Q4_K_M.gguf -p "a lovely cat" --cfg-scale 3 -H 1024 -W 1024
+```
+
+<img width="256" alt="iris-3b example" src="../assets/iris/example.png" />
+
+## Notes
+
+- Width and height must be multiples of 16. Do not pass `--vae`.
+- Captions are limited to 300 tokens including the assistant-turn suffix. Positive and negative prompts use the same template.