unsloth/studio/backend/core/inference/diffusion_speed.py
Daniel Han 6b9b1c72d3
Studio diffusion (Phase 7): accuracy-preserving speed pass (2.2x via GGUF compile) (#6690)
* Studio diffusion: cross-platform device policy, fp16 guard, lock split, validate-before-evict

Phase 1 of porting the richer diffusion stack onto the image-generation backend.

- Add a compartmentalized device/dtype policy module (diffusion_device.py)
  resolving CUDA/ROCm/XPU/MPS/CPU with capability flags. Keeps the NVIDIA
  capability-based bf16 choice; ROCm and XPU are isolated; MPS uses bf16 or
  fp32, never a silent fp16 that renders a black image.
- Add a per-family fp16_incompatible flag (Z-Image) and promote a resolved
  float16 to float32 for those families so they do not produce black images.
- Split the backend locks: a generation holds only _generate_lock, so status,
  unload, and a new load are never blocked by a long denoise. Add per-generation
  cancellation via callback_on_step_end so an eviction or a superseding load
  preempts a running generation; a replacement load waits for it to stop before
  allocating, so two pipelines never sit in VRAM at once.
- Validate a load request before the GPU handoff so an unloadable pick never
  evicts a working chat model, and reject missing local paths up front.
- Add CPU-only tests for the device policy, dtype guard, lock split and
  cancellation, and validate-before-evict, plus a GPU benchmark/regression
  script (scripts/diffusion_bench.py) measuring latency, peak VRAM, and PSNR
  against a saved reference.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 2A): measured-budget memory planner + offload/VAE policy

Add a lean, backend-agnostic memory policy that picks a CPU-offload policy and
VAE tiling/slicing from measured free device memory vs the model's estimated
resident footprint, then applies it to the built pipeline. auto stays resident
when the model fits (byte-identical to the prior resident path), and falls to
whole-module offload when tight; fast/balanced/low_vram are explicit overrides.
Sequential submodule offload is unreliable for GGUF transformers on diffusers
0.38, so it falls back to whole-module offload and status reports the policy
actually engaged.

Verified on Z-Image-Turbo Q4_K_M (B200): auto reproduces the resident image with
no VRAM/latency regression (PSNR inf); balanced/low_vram cut generation peak VRAM
47.9% (15951 -> 8318 MB) with byte-identical output, at the expected latency cost.

73 prior + 35 new CPU tests pass.

* Studio diffusion (Phase 2D): streamed block-level offload + functional VAE tiling

Add a streamed 'group' offload tier (diffusers apply_group_offloading, block_level,
use_stream) that keeps the transformer flowing through the GPU a few blocks at a
time while the text encoder / VAE stay resident, and fix VAE tiling to drive the
VAE submodule (pipelines like Z-Image expose enable_tiling on pipe.vae, not the
pipeline). apply_memory_plan now returns the (policy, tiling) actually engaged so
status never overstates either, and group falls back to whole-module offload when
the transformer can't be streamed.

Measured on Z-Image (B200), all lossless (PSNR inf vs resident): balanced/group
cuts generation peak VRAM 32% (15951 -> 10840 MB) at near-resident speed (2.07 ->
2.99s); low_vram/model cuts it 48% (-> 8318 MB) but is slower (7.99s). Mode names
now match that tradeoff: balanced = stream the transformer, low_vram = offload
every component. auto picks group when the companions fit resident, else model.

112 CPU tests pass.

* Studio diffusion (Phase 5): image quality-vs-quant accuracy harness

Add scripts/diffusion_quality.py, the accuracy analogue of the KLD workflow: hold
prompt + seed fixed, render a grid with a reference quant (default BF16), then render
each candidate quant and measure drift from the reference. Records mean PSNR + SSIM
(pure-numpy, no skimage/scipy) and optional CLIP text-alignment + image-similarity
(transformers, --clip), plus file size, latency, and peak VRAM, then prints a
quality-vs-cost table and recommends the smallest quant within a quality budget.
--selftest validates the metrics on synthetic images with no GPU or model.

Verified on Z-Image (B200): the table degrades monotonically with quant size
(Q8 -> Q4 -> Q2: PSNR 21.7 -> 15.5, SSIM 0.82 -> 0.61), while CLIP-text stays flat
(~0.34) -- quantization erodes fine detail far more than prompt adherence.

* Studio diffusion (Phase 3): opt-in speed layer (channels_last / compile / TF32)

Add a speed_mode knob (off by default, so the render path stays bit-identical):
default applies channels_last VAE + regional torch.compile of the denoiser's
repeated block where eligible; max also enables TF32 matmul and fused QKV. Regional
compile is gated off for the GGUF transformer (dequantises per-op) and for families
flagged not compile-friendly (a new supports_torch_compile flag, False for Z-Image),
so it activates automatically only once a non-GGUF bf16 transformer is loaded. Speed
optims run before placement/offload, per the diffusers composition order. status now
reports speed_mode + the optims actually engaged.

Verified on Z-Image (B200): default -> ['channels_last'], max -> ['channels_last',
'tf32'], compile correctly skipped for GGUF; generation works in every mode.

121 CPU tests pass.

* Studio diffusion (Phase 2B): opt-in fp8 text-encoder layerwise casting

Add a text_encoder_fp8 knob that casts the companion text encoder(s) to fp8 (e4m3)
storage via diffusers apply_layerwise_casting, upcasting per layer to the bf16
compute dtype while normalisations and embeddings stay full precision. Applied
before placement, gated to CUDA + bf16, best-effort (a failure leaves the encoder
dense). status reports which encoders were cast.

Verified on Z-Image (B200, balanced/group mode where the encoder stays resident):
generation peak VRAM dropped 37% (10840 -> 6791 MB, below the lowest-VRAM offload)
at near-resident speed. It is a memory-vs-quality tradeoff, not free -- ~20 dB PSNR
vs the bf16 encoder, a larger shift than one transformer quant step -- so it is off
by default and documented as such, with the Phase 5 harness to size the cost.

127 CPU tests pass.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 2C): NVFP4 text-encoder quant (+ generalise fp8 knob)

Generalise the text-encoder precision knob from a fp8 bool to text_encoder_quant
(fp8 | nvfp4). nvfp4 quantises the companion text encoder to 4-bit via torchao
NVFP4 weight-only (two-level microscaling) on Blackwell's FP4 tensor cores; fp8
stays the broader-hardware path (cc>=8.9). Both are gated, best-effort, and run
before placement; status reports the mode actually engaged. This is the lean
realisation of GGUF-native text-encoder quant: 4-bit on the encoder without the
3045-line port.

Verified on Z-Image (B200, balanced/group where the encoder stays resident), vs the
bf16 encoder: nvfp4 cut generation peak VRAM 48% (10840 -> 5593 MB, the lowest TE
option, below whole-model offload) at near-fp8 quality (16.4 vs 17.1 dB PSNR), and
both quants ran faster than bf16. A memory-vs-quality tradeoff (off by default);
size it per model with the Phase 5 quality harness. diffusion_bench gains
--text-encoder-quant.

129 CPU tests pass.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 4): native stable-diffusion.cpp engine for CPU/Mac

Adds the CPU / Apple-Silicon tier of the two-engine strategy, mirroring the
chat backend's llama.cpp shell-out. Diffusers stays the default on CUDA / ROCm
/ XPU; this covers the hardware diffusers serves poorly, consuming the same
split GGUF assets Studio already curates.

- sd_cpp_args.py: pure sd-cli command builder. Maps the family to its
  text-encoder flag (Z-Image Qwen3 to --llm, Qwen-Image to --qwen2vl, FLUX.1
  CLIP-L + T5), and the diffusers memory policy (none/group/model/sequential)
  to sd.cpp's offload flags (--offload-to-cpu / --clip-on-cpu / --vae-on-cpu /
  --vae-tiling / --diffusion-fa), so one user knob drives both engines.
- sd_cpp_engine.py: SdCppEngine over a located sd-cli. find_sd_cpp_binary()
  with the same precedence as the llama finder (env override, then the Studio
  install root, then in-tree, then PATH), an is_available/version probe, and a
  one-shot subprocess generate that streams progress and returns the PNG.
  runtime_env() prepends the binary's directory to the platform library path
  so a prebuilt's bundled libstable-diffusion.so resolves.
  select_diffusion_engine() is the pure routing decision (GPU backends to
  diffusers, CPU/MPS to native when present).
- install_sd_cpp_prebuilt.py: resolve + download the per-host prebuilt
  (macOS-arm64/Metal, Linux x86_64 CPU, Vulkan/ROCm/Windows variants) into the
  Studio install root. resolve_release_asset() is a pure, unit-tested
  host-to-asset matrix.
- scripts/sd_cpp_smoke.py: end-to-end native generation harness.

Tests (CPU-only, subprocess/filesystem stubbed): 49 new across args, engine,
routing, runtime env, and the installer resolver. Full diffusion suite 166
passing.

Verified on a B200 box: built sd-cli (CUDA) and the prebuilt (CPU) both
generate Z-Image-Turbo Q4_K end to end through SdCppEngine: balanced (group
offload, 5.0s gen), low_vram (full CPU offload + VAE tiling, 13.4s), and the
dynamically-linked CPU prebuilt (50.4s on CPU), all producing coherent images.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 6): img2img / inpaint / edit / LoRA / upscale on the native engine

Builds on Phase 4's native stable-diffusion.cpp engine, extending it from
text-to-image to the wider feature surface, since sd.cpp supports all of these
through the binary already. Pure command-builder additions plus one engine
method, so the txt2img path is unchanged.

- sd_cpp_args.py: SdCppGenParams gains image-conditioning fields. init_img +
  strength make a run img2img, adding mask makes it inpaint, ref_images drives
  FLUX-Kontext / Qwen-Image-Edit style editing (repeated --ref-image), and
  lora_dir + the <lora:name:weight> prompt syntax select LoRAs. New
  SdCppUpscaleParams + build_sd_cpp_upscale_command for the ESRGAN upscale run
  mode (input image + esrgan model, no prompt / text encoders).
- sd_cpp_engine.py: the subprocess runner is factored into a shared _run() so
  generate() (now carrying the conditioning flags) and a new upscale() reuse
  the same streaming / error / output-check path.
- scripts/sd_cpp_smoke.py: --task {txt2img,img2img,upscale} with --init-img /
  --strength / --upscale-model / --upscale-repeats.

Tests: 10 new across the img2img / inpaint / edit / LoRA flag construction, the
upscale builder and its validation, and the engine's img2img + upscale paths.
Full diffusion suite 176 passing.

Verified on a B200 box through SdCppEngine: img2img (Z-Image-Turbo Q4_K, the
init image conditioned at strength 0.6, 4.8s) and ESRGAN upscale
(512x512 -> 2048x2048 via RealESRGAN_x4plus_anime_6B, 2.7s), both producing
coherent images. Video and the diffusers-path feature wiring are deferred.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 7): accuracy-preserving speed pass

Re-review of the diffusion stack (#6675/#6679/#6680) surfaced one real accuracy
bug and a dead-on-arrival speed path; this fixes both and adds the lossless /
near-lossless wins, all measured on a B200.

Correctness:
- TF32 global-state leak (fix). speed_mode=max flipped torch.backends.*.allow_tf32
  process-wide and never restored them, so a later `off` load silently inherited
  TF32 and was no longer bit-identical. Added snapshot_backend_flags /
  restore_backend_flags (TF32 + cudnn.benchmark), captured before the speed layer
  runs and restored on unload. Verified: load max -> unload -> load off is now
  byte-identical (PSNR inf) to a fresh off.
- sd-cli timeout could hang forever. _run() blocked in `for line in stdout` and
  only checked the timeout after EOF, so a child stuck in model load / GPU init
  with no output ignored the timeout. Drained stdout on a reader thread with a
  wall-clock deadline. Added a silent-hang regression test.

Speed (diffusers path), near-lossless, opt-in tiers:
- Regional torch.compile now runs on the GGUF transformer. The is_gguf gate (and
  Z-Image's supports_torch_compile=False) were stale: compile_repeated_blocks
  compiles and runs ~2.2x faster on the GGUF Z-Image transformer on
  torch 2.9.1 / diffusers 0.38 (the per-op dequant stays eager, the rest of the
  block compiles). Measured: off 1.80s -> default 0.82s/gen (+54.7%), PSNR 37.7 dB
  vs eager -- far above the Q4 quant noise floor (~21 dB), so it does not move
  output quality. Gate relaxed; default tier delivers it.
- cudnn.benchmark added to the default tier (autotunes the fixed-shape VAE convs).
- torch.inference_mode() around the pipeline call (lossless, strictly faster than
  the no_grad diffusers uses internally).

Memory path:
- VAE tiling (not bit-identical >1MP) restricted to the model/sequential/CPU tiers;
  the balanced (group) tier keeps exact slicing only, so it is now bit-identical to
  the resident image (verified PSNR inf) and slightly faster.
- Group offload adds non_blocking + record_stream on the CUDA stream path to
  overlap each block's H2D copy with compute (lossless; gated on the installed
  diffusers signature so older versions still work).

Native (sd.cpp) path:
- native_speed_flags: a first-class speed knob (default -> --diffusion-fa, a
  near-lossless CUDA win that was previously only added on offload tiers; max also
  -> --diffusion-conv-direct). conv-direct stays opt-in: measured +45% on CUDA, so
  it is never auto-on. Engine generate() merges it, de-duped against offload flags.

Default profile: a GGUF model with no explicit speed_mode now resolves to the
`default` profile (resolve_speed_mode), since compile's perturbation sits below the
quantisation noise floor and so does not reduce quality versus the dense reference;
out of the box a GGUF Z-Image generation drops from 1.80s to 0.81s. Dense models
stay `off` / bit-identical, and an explicit speed_mode -- including "off" -- is
always honored, so the byte-identical path remains one flag away and is the
regression reference.

Tooling: scripts/compile_probe.py (eager vs compiled GGUF probe), scripts/
perf_verify.py (the B200 verification above), and diffusion_bench.py gains
--speed-mode so the speed tiers are benchmarkable.

Tests: 183 passing (was 166); new coverage for the backend-flag snapshot/restore,
GGUF compile eligibility, the balanced tiling/slicing split, native_speed_flags +
the engine de-dup, and the sd-cli silent-hang timeout.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 7): max tier uses max-autotune-no-cudagraphs + engine/lever benchmarks

The opt-in `max` speed tier now compiles the repeated block with
mode=max-autotune-no-cudagraphs (dynamic=False) instead of the default mode:
Triton autotuning for GEMM/conv-heavier models, gated to the tier where a longer
cold compile is acceptable. CUDA-graph modes (reduce-overhead / max-autotune) are
deliberately avoided -- both crash on the regionally-compiled block (its static
output buffer is overwritten across denoise steps), measured.

Adds two reproducible benchmarks used to validate the optimization research:
- scripts/compare_engines.py: PyTorch (diffusers GGUF) vs native sd.cpp head-to-head.
- scripts/leverage_probe.py: coordinate_descent_tuning + FirstBlockCache probes.

Measured on B200 (Z-Image Q4_K_M, 1024px, 8 steps): default compile 0.80s/gen;
coordinate_descent_tuning 0.79s (within noise, already covered by max-autotune);
FirstBlockCache does not run on Z-Image (diffusers 0.38 block-detection / Dynamo).

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 7): robust backend-flag snapshot/restore and restore on failed speeded load

- snapshot_backend_flags reads each flag defensively (getattr + hasattr), so a build/platform
  missing one (no cuda.matmul on CPU/MPS) still captures the rest instead of skipping the
  whole snapshot. restore_backend_flags restores each flag independently so one failure can't
  leave the others leaked process-wide.
- load_pipeline restores the flags (and clears the GPU cache) when the build fails after
  apply_speed_optims mutated the process-wide flags but before _state captured them for unload
  to restore -- otherwise a failed default/max load left cudnn.benchmark/TF32 on and
  contaminated later off generations.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 4): enforce the sd-cli timeout while reading output

Iterating proc.stdout directly blocks until the stream closes, so a sd-cli that hangs
without producing output (or without closing stdout) would never reach proc.wait and the
wall-clock timeout was silently bypassed. Drain stdout on a daemon thread and wait on the
PROCESS, so the main thread always enforces the timeout and kills a hung process (which
closes the pipe and ends the reader). Add a test that times out even when stdout blocks,
and make the no-binary test hermetic so a host-installed sd-cli can't leak in.

* Studio diffusion (Phase 7) review fixes: offload fallback + bench scripts

- diffusion_memory: when group offload is unavailable and the plan falls back to
  whole-module offload, enable VAE tiling (the group plan left it off, but the fallback
  is the low-VRAM path where the decode spike can OOM). Covers both the group and
  sequential fallback branches.
- perf_verify: include the balanced-vs-off PSNR in the pass/fail condition, so a
  balanced bit-identity regression actually fails the check instead of exiting 0.
- compare_engines: --vae/--llm default to None (were author-absolute /mnt paths), and
  the load-progress poll has a 30 min deadline instead of looping forever on a hang.
- test for the group->model fallback enabling VAE tiling.

* Studio diffusion (Phase 4) review fixes: sd.cpp installer + engine hardening

- install_sd_cpp_prebuilt: download the release archive with urlopen + an explicit
  timeout + copyfileobj (urlretrieve has no timeout and hangs on a stalled socket);
  extract through a per-member containment check (Zip-Slip guard); expanduser the
  --install-dir so a tilde path is not taken literally; and on Windows CUDA also fetch
  the separately-published cudart runtime DLL archive so sd-cli.exe can start.
- sd_cpp_engine: find_sd_cpp_binary honors UNSLOTH_STUDIO_HOME / STUDIO_HOME like the
  installer, so a custom-root install is discovered without UNSLOTH_SD_CPP_PATH; start
  sd-cli with the parent-death child_popen_kwargs so it is not orphaned on a backend
  crash; reap the SIGKILLed child (proc.wait) so a cancel/timeout does not leave a zombie.
- tests: Zip-Slip rejection, normal extraction, studio-home discovery.

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Studio diffusion (Phase 4) review round 2: collect sd-cli batch outputs

Codex review: when batch_count > 1, stable-diffusion.cpp's save_results() writes
the numbered files <stem>_<idx><suffix> (base_0.png, base_1.png, ...) instead of
the literal --output path. SdCppEngine.generate checked only the literal path, so
a batch generation would exit 0 and then raise 'no image' (or return a stale
file). generate now returns the literal path when present and otherwise falls
back to the numbered siblings; single-image behavior is unchanged.

Test: a fake sd-cli that writes img_0.png/img_1.png (not img.png) is collected
without error.

* Studio diffusion (Phase 6) review round 2: img2img source dims + upscale repeats

Codex review on the native engine arg builder:

- build_sd_cpp_command emitted --width/--height unconditionally, so an
  img2img/inpaint/edit run that left dims unset forced a 1024x1024 resize/crop of
  the input. width/height are now Optional (None = unset): an image-conditioned
  run (init_img or ref_images) with unset dims omits the flags so sd.cpp derives
  the size from the input image (set_width_and_height_if_unset); a plain txt2img
  run with unset dims keeps the prior 1024x1024 default; explicit dims are always
  honored. width/height are read only by the builder, so the type change is local.

- build_sd_cpp_upscale_command used a truthiness guard (params.repeats and ...)
  that silently swallowed repeats=0 into sd-cli's default of one pass, turning an
  explicit no-op into a real upscale. It now rejects repeats < 1 with ValueError
  and emits the flag for any explicit value != 1.

Tests: img2img unset dims omit width/height (init_img and ref_images), explicit
dims emitted, txt2img keeps 1024; upscale rejects repeats=0 and omits the flag at
the default. (Two pre-existing binary-discovery tests fail only because a real
sd-cli is installed in this dev environment; unrelated to this change.)

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: oobabooga <112222186+oobabooga@users.noreply.github.com>
2026-07-01 15:31:50 -03:00

304 lines
12 KiB
Python

# SPDX-License-Identifier: AGPL-3.0-only
# Copyright 2026-present the Unsloth AI Inc. team. All rights reserved. See /studio/LICENSE.AGPL-3.0
"""Opt-in speed optimisations for the local diffusion backend.
Off by default, so the default render path stays bit-identical to a plain run (the
property the regression harness checks). When the operator opts in, this applies the
near-lossless speedups in the order the diffusers guides recommend
(channels_last + cudnn.benchmark -> regional compile, with TF32 / fused-QKV under
"max"):
off - nothing (default; bit-identical reference).
default - near-lossless: channels_last VAE memory format + cudnn.benchmark conv
autotune + regional torch.compile of the denoiser's repeated block WHERE
eligible (bf16, CUDA, a compile-friendly family). Compile is the big win
(~2.3x denoise on the GGUF Z-Image transformer, PSNR ~36 dB vs eager,
well above the Q4 quantisation noise floor, so it does not meaningfully
move output quality).
max - default plus near-lossless TF32 matmul and fused QKV projections.
Regional compile used to be gated off for the GGUF transformer, but it compiles and
runs faster on the current diffusers/torch (measured; the GGUF dequant ops stay
eager and the rest of the repeated block compiles), so the GGUF gate is removed; the
per-family ``supports_torch_compile`` flag and the bf16/CUDA checks still apply.
The backend flags this layer flips (TF32, cudnn.benchmark) are PROCESS-WIDE, so
``snapshot_backend_flags`` / ``restore_backend_flags`` let the caller capture the
prior values at load and restore them at unload, keeping a later ``off`` load
bit-identical instead of inheriting a previous ``max`` run's globals. torch is
imported lazily.
"""
from __future__ import annotations
from typing import Any, Optional
SPEED_OFF = "off"
SPEED_DEFAULT = "default"
SPEED_MAX = "max"
SPEED_MODES = (SPEED_OFF, SPEED_DEFAULT, SPEED_MAX)
def snapshot_backend_flags() -> Optional[dict]:
"""Capture the process-wide torch backend flags this layer may mutate, so the
caller can restore them on unload. None if torch is unavailable. Each flag is read
defensively so a build/platform missing one (e.g. no cuda.matmul on CPU/MPS) still
captures the rest -- otherwise a single missing attribute would skip the whole
snapshot and a real mutated flag would leak."""
try:
import torch
except Exception: # noqa: BLE001 — no torch -> nothing to snapshot/restore
return None
state: dict[str, bool] = {}
matmul = getattr(getattr(torch.backends, "cuda", None), "matmul", None)
if matmul is not None and hasattr(matmul, "allow_tf32"):
state["matmul_tf32"] = bool(matmul.allow_tf32)
cudnn = getattr(torch.backends, "cudnn", None)
if cudnn is not None:
if hasattr(cudnn, "allow_tf32"):
state["cudnn_tf32"] = bool(cudnn.allow_tf32)
if hasattr(cudnn, "benchmark"):
state["cudnn_benchmark"] = bool(cudnn.benchmark)
return state
def restore_backend_flags(state: Optional[dict]) -> None:
"""Restore the flags captured by ``snapshot_backend_flags``. No-op on None. Each
flag is restored independently so one failure can't leave the others leaked."""
if not state:
return
try:
import torch
except Exception: # noqa: BLE001 — no torch -> nothing to restore
return
def _set(obj: Any, attr: str, key: str) -> None:
if obj is not None and key in state and hasattr(obj, attr):
try:
setattr(obj, attr, state[key])
except Exception: # noqa: BLE001 — best-effort per-flag restore
pass
_set(
getattr(getattr(torch.backends, "cuda", None), "matmul", None), "allow_tf32", "matmul_tf32"
)
cudnn = getattr(torch.backends, "cudnn", None)
_set(cudnn, "allow_tf32", "cudnn_tf32")
_set(cudnn, "benchmark", "cudnn_benchmark")
def normalize_speed_mode(value: Optional[str]) -> str:
"""Lower/strip a requested speed mode (dashes ok); None / "" -> off."""
if value is None:
return SPEED_OFF
normalized = str(value).strip().lower().replace("-", "_")
if not normalized:
return SPEED_OFF
if normalized not in SPEED_MODES:
raise ValueError(
f"Unsupported diffusion speed_mode '{value}'. Use one of: {', '.join(SPEED_MODES)}."
)
return normalized
def resolve_speed_mode(value: Optional[str], *, is_gguf: bool) -> str:
"""The effective speed mode when the caller leaves it UNSET (``None``).
A GGUF model defaults to ``default``: regional compile is ~2.2x faster and its
numeric perturbation sits well below the quantisation noise floor (measured
PSNR ~37 dB compile-vs-eager versus ~21 dB Q4-vs-bf16), so it does not reduce
output quality relative to the dense reference. A dense (non-GGUF) model stays
``off`` / bit-identical, since there compile would be the only source of drift.
An explicit value -- including ``"off"`` -- is always honored verbatim."""
if value is None:
return SPEED_DEFAULT if is_gguf else SPEED_OFF
return normalize_speed_mode(value)
def compile_eligible(target: Any, *, is_gguf: bool, family: Any) -> bool:
"""Whether the denoiser's repeated block should be regionally compiled.
Only on CUDA (incl. ROCm via supports_default_torch_compile), for a bf16
transformer, on a compile-friendly family. ``is_gguf`` no longer disqualifies:
``compile_repeated_blocks`` runs fine on the GGUF transformer (the per-op
dequant stays eager, the rest of the block compiles) and is ~2.3x faster, so it
is kept only for signature/logging compatibility."""
del is_gguf # GGUF is compile-eligible now; param kept for call-site compat.
if not bool(getattr(target, "supports_default_torch_compile", False)):
return False
if not bool(getattr(family, "supports_torch_compile", True)):
return False
return _is_bfloat16(getattr(target, "dtype", None))
def _is_bfloat16(dtype: Any) -> bool:
try:
import torch
return dtype is torch.bfloat16
except Exception:
return str(dtype).endswith("bfloat16")
def apply_speed_optims(
pipe: Any,
target: Any,
*,
is_gguf: bool,
family: Any,
speed_mode: str = SPEED_OFF,
logger: Any = None,
) -> dict[str, bool]:
"""Apply the opt-in speed optimisations for ``speed_mode`` to a built pipeline,
BEFORE placement / offload. Returns which optimisations actually engaged. Every
step is best-effort: a pipeline that doesn't support one is simply skipped."""
applied = {
"channels_last": False,
"cudnn_benchmark": False,
"tf32": False,
"fused_qkv": False,
"compiled": False,
}
mode = normalize_speed_mode(speed_mode)
# TF32 is the one PROCESS-GLOBAL flag we flip (on max). Restore it whenever this
# load isn't max, so a later default/off diffusion load -- or chat inference in the
# same long-lived process -- doesn't silently inherit a prior max load's TF32 and
# lose the bit-identical default the regression harness checks.
if mode != SPEED_MAX:
_restore_tf32(logger)
if mode == SPEED_OFF:
return applied
# Lossless: a channels-last VAE speeds up its convolutions with no numeric change.
applied["channels_last"] = _vae_channels_last(pipe, logger)
# Near-lossless: let cuDNN autotune the fixed-shape VAE convs (CUDA only). It may
# pick a different conv algorithm, so it is a "default"-tier (not bit-identical) win.
if getattr(target, "device", None) == "cuda":
applied["cudnn_benchmark"] = _enable_cudnn_benchmark(logger)
# Near-lossless and the largest win: regional compile of the repeated denoiser
# block, where eligible (now incl. the GGUF transformer). `max` opts into
# max-autotune (longer compile, autotuned kernels).
if compile_eligible(target, is_gguf = is_gguf, family = family):
applied["compiled"] = _compile_repeated_blocks(pipe, logger, max_autotune = mode == SPEED_MAX)
if mode == SPEED_MAX:
# Near-lossless: TF32 matmul (CUDA only) trades a few mantissa bits for speed.
if getattr(target, "device", None) == "cuda":
applied["tf32"] = _enable_tf32(logger)
applied["fused_qkv"] = _fuse_qkv(pipe, logger)
return applied
def _vae_channels_last(pipe: Any, logger: Any) -> bool:
vae = getattr(pipe, "vae", None)
if vae is None or not hasattr(vae, "to"):
return False
try:
import torch
vae.to(memory_format = torch.channels_last)
return True
except Exception as exc: # noqa: BLE001 — optimisation only
_warn(logger, "channels_last", exc)
return False
def _compile_repeated_blocks(
pipe: Any,
logger: Any,
*,
max_autotune: bool = False,
) -> bool:
transformer = getattr(pipe, "transformer", None)
fn = getattr(transformer, "compile_repeated_blocks", None)
if not callable(fn):
return False
# default: mode="default" + dynamic=True -- fast cold start, robust to resolution
# changes (no recompile). max: mode="max-autotune-no-cudagraphs" + dynamic=False --
# Triton autotuning for a few % more on GEMM/conv-heavy models, at a much longer
# compile and a recompile per new resolution. The CUDA-graph modes (reduce-overhead
# / max-autotune) are deliberately NOT used: they crash on the regionally-compiled
# block because its static output buffer is overwritten across denoise steps.
kwargs: dict[str, Any] = {"fullgraph": True, "dynamic": not max_autotune}
if max_autotune:
kwargs["mode"] = "max-autotune-no-cudagraphs"
try:
fn(**kwargs)
return True
except Exception as exc: # noqa: BLE001 — optimisation only
_warn(logger, "compile_repeated_blocks", exc)
return False
def _enable_cudnn_benchmark(logger: Any) -> bool:
try:
import torch
torch.backends.cudnn.benchmark = True
return True
except Exception as exc: # noqa: BLE001 — optimisation only
_warn(logger, "cudnn_benchmark", exc)
return False
# The TF32 flag values from before the first max load flipped them, so a later
# non-max load / unload can put the process back exactly as it found it (rather than
# forcing a hardcoded default that might clobber another component's choice).
_tf32_prev: Optional[tuple[bool, bool]] = None
def _enable_tf32(logger: Any) -> bool:
global _tf32_prev
try:
import torch
if _tf32_prev is None:
_tf32_prev = (
torch.backends.cuda.matmul.allow_tf32,
torch.backends.cudnn.allow_tf32,
)
torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cudnn.allow_tf32 = True
return True
except Exception as exc: # noqa: BLE001 — optimisation only
_warn(logger, "tf32", exc)
return False
def restore_tf32(logger: Any = None) -> None:
"""Put the process-global TF32 flags back to their pre-max-load values. No-op if
a max load never set them. Called on a non-max load and on unload."""
_restore_tf32(logger)
def _restore_tf32(logger: Any) -> None:
global _tf32_prev
if _tf32_prev is None:
return
try:
import torch
torch.backends.cuda.matmul.allow_tf32 = _tf32_prev[0]
torch.backends.cudnn.allow_tf32 = _tf32_prev[1]
except Exception as exc: # noqa: BLE001 — best-effort restore
_warn(logger, "tf32_restore", exc)
finally:
_tf32_prev = None
def _fuse_qkv(pipe: Any, logger: Any) -> bool:
for owner in (pipe, getattr(pipe, "transformer", None)):
fn = getattr(owner, "fuse_qkv_projections", None)
if callable(fn):
try:
fn()
return True
except Exception as exc: # noqa: BLE001 — optimisation only
_warn(logger, "fuse_qkv_projections", exc)
return False
return False
def _warn(logger: Any, what: str, exc: Exception) -> None:
if logger is not None:
logger.warning("diffusion.speed: %s failed: %s", what, exc)