fix: reclaim before accepting wrapped Vulkan iGPU memory estimates (#2124)

docker/Dockerfile.rocm

@@ -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_20.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" ]

docs/docker.md

@@ -24,12 +24,25 @@ docker build -f docker/Dockerfile -t sd .
 
 ## Building variants using Docker
 
-Vulkan:
+### 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
 
 ```shell
@@ -37,3 +50,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"
+```

docs/qwen_image_2.1.md

@@ -32,6 +32,16 @@ Run the following commands from the build directory. Use image dimensions divisi
 
 To use GGUF diffusion weights, set `--diffusion-model` to the path of a file such as `qwen_image_2.1-Q4_K.gguf`.
 
+### Turbo
+
+For [Qwen Image 2.1 Turbo](https://huggingface.co/Qwen/Qwen-Image-2.1-Turbo), use the Turbo diffusion weights and select its fixed 8-step sigma schedule with `--sigmas qwen-image-2.1-turbo`. For example, add these sampling options to the command above:
+
+```text
+--sigmas qwen-image-2.1-turbo --steps 8 --cfg-scale 1.0 --sampling-method euler
+```
+
+The preset is equivalent to `--sigmas 1.0,0.978453,0.95418,0.926626,0.89508,0.845148,0.704534,0.414568,0.0`. It only supplies the sigma values; CFG scale and sampling method remain independently configurable. As with an explicit sigma list, the sigma count determines the actual number of sampling steps. The same preset name is also accepted by `--hires-sigmas` for the highres fix second pass.
+
 ### Image editing
 
 Pass the reference image with `-r` and describe the edit in `-p`. Vision weights are required; the example below loads them separately with `--llm_vision`.

examples/common/common.cpp

@@ -1475,6 +1475,9 @@ ArgOptions SDGenerationParams::get_options() {
             return -1;
         }
         std::string sigmas_str = value;
+        if (sigmas_str == "qwen-image-2.1-turbo") {
+            sigmas_str = "1.0,0.978453,0.95418,0.926626,0.89508,0.845148,0.704534,0.414568,0.0";
+        }
         if (!sigmas_str.empty() && sigmas_str.front() == '[') {
             sigmas_str.erase(0, 1);
         }
@@ -1721,11 +1724,11 @@ ArgOptions SDGenerationParams::get_options() {
          on_scheduler_arg},
         {"",
          "--sigmas",
-         "custom sigma values for the sampler, comma-separated (e.g., \"14.61,7.8,3.5,0.0\").",
+         "custom sigma values for the sampler, comma-separated (e.g., \"14.61,7.8,3.5,0.0\") or preset: qwen-image-2.1-turbo.",
          on_sigmas_arg},
         {"",
          "--hires-sigmas",
-         "custom sigma values for the highres fix second pass, comma-separated (e.g., \"0.85,0.725,0.421875,0.0\").",
+         "custom sigma values for the highres fix second pass, comma-separated (e.g., \"0.85,0.725,0.421875,0.0\") or preset: qwen-image-2.1-turbo.",
          on_hires_sigmas_arg},
         {"",
          "--skip-layers",

src/model_manager.cpp

@@ -1757,7 +1757,14 @@ ModelManager::CapacityCheck ModelManager::check_capacity(
         }
         // Vulkan's heap budget subtraction can underflow when usage exceeds the budget.
         if (total_bytes > 0 && free_bytes > total_bytes && sd_backend_is(backend, "Vulkan")) {
-            return size_t{0};
+            if (ggml_backend_dev_type(device) != GGML_BACKEND_DEVICE_TYPE_IGPU) {
+                return size_t{0};
+            }
+            // A wrapped report may reflect real pressure; reclaim before accepting the tracked estimate.
+            LOG_VERBOSE("ignoring impossible free memory report on %s (%.2f MB free / %.2f MB total), using tracked residency",
+                        ggml_backend_name(backend), free_bytes / (1024.0 * 1024.0), total_bytes / (1024.0 * 1024.0));
+            free_bytes                  = total_bytes;
+            result.requires_reclamation = true;
         }
         if (total_bytes > 0) {
             free_bytes = std::min(free_bytes, resident < total_bytes ? total_bytes - resident : 0);
@@ -1803,8 +1810,11 @@ bool ModelManager::fits_compute_backend_capacity(
     const DeviceMemoryRequest& request,
     const std::vector<ggml_tensor*>& required_params) const {
     std::vector<TensorState*> states;
-    return resolve_required_tensor_states(required_params, states, request.compute_backend) &&
-           check_capacity(request, states).fits();
+    if (!resolve_required_tensor_states(required_params, states, request.compute_backend)) {
+        return false;
+    }
+    const auto capacity = check_capacity(request, states);
+    return !capacity.requires_reclamation && capacity.fits();
 }
 
 bool ModelManager::ensure_compute_backend_capacity(
@@ -1826,7 +1836,10 @@ bool ModelManager::ensure_compute_backend_capacity(
         return true;
     }
 
-    auto fits = [&]() { return check_capacity(request, required_states).fits(); };
+    auto fits = [&]() {
+        const auto capacity = check_capacity(request, required_states);
+        return !capacity.requires_reclamation && capacity.fits();
+    };
     if (fits()) {
         return true;
     }
@@ -1866,7 +1879,12 @@ bool ModelManager::ensure_compute_backend_capacity(
             eviction_states.insert(state);
         }
     };
+    size_t last_eviction_count   = 0;
     auto release_eviction_states = [&]() {
+        if (eviction_states.size() == last_eviction_count) {
+            return false;
+        }
+        last_eviction_count = eviction_states.size();
         release_compute_staging_blocks(false, &eviction_states);
         release_params_storage_blocks(false, &eviction_states);
         return fits();
@@ -1947,6 +1965,9 @@ bool ModelManager::ensure_compute_backend_capacity(
     }
     revert_demotions();
 
+    if (check_capacity(request, required_states).fits()) {
+        return true;
+    }
     const auto capacity                = check_capacity(request, required_states, true);
     const std::string available_device = capacity.available_device_bytes == SIZE_MAX
                                              ? "unknown"

src/model_manager.h

@@ -152,6 +152,7 @@ private:
         size_t required_budget_bytes  = 0;
         size_t available_device_bytes = SIZE_MAX;
         size_t available_budget_bytes = SIZE_MAX;
+        bool requires_reclamation     = false;
         bool fits() const {
             return required_device_bytes <= available_device_bytes &&
                    required_budget_bytes <= available_budget_bytes;

src/pipeline/diffusion_engine.cpp

@@ -2735,9 +2735,9 @@ sd::Tensor<float> StableDiffusionGGML::sample(const std::shared_ptr<DiffusionMod
         if (batch_cfg_ok) {
             sd::Tensor<float> batched_out = run_batched_condition(*positive_condition, c_concat_override);
             if (!batched_out.empty() && batched_out.dim() >= 4 && batched_out.shape()[3] == 2) {
-                auto parts    = sd::ops::chunk(batched_out, 2, 3);
-                cond_out      = std::move(parts[0]);
-                uncond_out    = std::move(parts[1]);
+                auto parts = sd::ops::chunk(batched_out, 2, 3);
+                cond_out   = std::move(parts[0]);
+                uncond_out = std::move(parts[1]);
             }
         }