Adds the opt-in speed path for the GGUF diffusion transformer behind a selectable speed mode (default off, so output is unchanged until a profile is chosen): - diffusion_eager_patches.py: shared eager fast-paths (channels_last, attention/backend selection, fused norms and QKV) installed at load and rolled back on unload or failed load. - diffusion_compile_cache.py / diffusion_gguf_compile.py: a persistent torch.compile cache and the GGUF-transformer compile wiring. - diffusion_arch_patches.py: architecture-specific patches. - diffusion_patch_backend.py: shared install/restore plumbing. - diffusion_speed.py: speed-profile planning. Tests for each module plus the benchmarking and probe scripts used to measure speed, memory, and accuracy of the path.
288 lines
12 KiB
Python
288 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 -> compile, with TF32 / fused-QKV under "max"):
|
|
|
|
off - nothing (default; bit-identical reference).
|
|
eager - everything lossless EXCEPT torch.compile: channels_last VAE +
|
|
cudnn.benchmark + the attention backend + the shared eager monkey-patches
|
|
(fused RMSNorm / AdaLayerNorm + per-arch addcmul fusions, see
|
|
diffusion_eager_patches.py / diffusion_arch_patches.py). The fast first-image
|
|
/ casual-use path -- no compile tax to amortise.
|
|
default - LIGHT compile. For a GGUF model: channels_last + cudnn.benchmark +
|
|
torch.compile of ONLY the dequant op chain
|
|
(``torch.compile(dequantize_gguf_tensor, dynamic=True)``) -- the dequant is
|
|
~70-80% of eager GGUF time, so fusing it gives ~1.24-1.64x for a small
|
|
one-time compile (~7.5-10.4s) and ZERO extra VRAM, resolution-invariant
|
|
(the dequant inputs are fixed-shape weights). For a dense (non-GGUF) model
|
|
there is no dequant, so ``default`` falls back to regional torch.compile of
|
|
the denoiser's repeated block (the only compile lever a dense model has).
|
|
max - the FULL torch.compile: regional max-autotune compile of the denoiser's
|
|
repeated block (which fuses the GGUF dequant AND the matmul/norm/elementwise
|
|
in one graph -- ~3.2x on the GGUF Z-Image transformer, PSNR ~36 dB vs eager,
|
|
well above the Q4 noise floor) plus TF32 matmul and fused QKV projections.
|
|
|
|
Tier rationale: ``default`` is the cheap, always-amortising compile (compile just the
|
|
hot GGUF dequant; the block stays eager) so the first image is fast and VRAM is
|
|
untouched; ``max`` pays the larger regional-compile tax for the bigger warm speedup.
|
|
The compiled dequant is deliberately skipped under ``max`` -- the regional block compile
|
|
subsumes the dequant fusion (a separately-compiled dequant would be traced into that
|
|
graph and break it), so ``max`` runs the stock dequant and lets the block compile it. The
|
|
per-family ``supports_torch_compile`` flag and the bf16/CUDA checks gate regional compile.
|
|
|
|
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
|
|
|
|
from . import diffusion_gguf_compile as gguf_compile
|
|
|
|
SPEED_OFF = "off"
|
|
SPEED_EAGER = "eager"
|
|
SPEED_DEFAULT = "default"
|
|
SPEED_MAX = "max"
|
|
SPEED_MODES = (SPEED_OFF, SPEED_EAGER, 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."""
|
|
try:
|
|
import torch
|
|
return {
|
|
"matmul_tf32": bool(torch.backends.cuda.matmul.allow_tf32),
|
|
"cudnn_tf32": bool(torch.backends.cudnn.allow_tf32),
|
|
"cudnn_benchmark": bool(torch.backends.cudnn.benchmark),
|
|
}
|
|
except Exception: # noqa: BLE001 — best-effort; no snapshot -> no restore
|
|
return None
|
|
|
|
|
|
def restore_backend_flags(state: Optional[dict]) -> None:
|
|
"""Restore the flags captured by ``snapshot_backend_flags``. No-op on None."""
|
|
if not state:
|
|
return
|
|
try:
|
|
import torch
|
|
|
|
torch.backends.cuda.matmul.allow_tf32 = state["matmul_tf32"]
|
|
torch.backends.cudnn.allow_tf32 = state["cudnn_tf32"]
|
|
torch.backends.cudnn.benchmark = state["cudnn_benchmark"]
|
|
except Exception: # noqa: BLE001 — best-effort restore
|
|
return
|
|
|
|
|
|
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``: it compiles only the hot dequant op chain
|
|
(~70-80% of eager GGUF time) for ~1.24-1.64x at a small one-time compile and zero
|
|
extra VRAM -- a cheap, always-amortising win whose numeric perturbation sits well
|
|
below the quantisation noise floor (the dequant graph is unchanged, just
|
|
Inductor-fused). 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,
|
|
cache_active: bool = False,
|
|
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,
|
|
"compiled_dequant": False,
|
|
}
|
|
mode = normalize_speed_mode(speed_mode)
|
|
if mode == SPEED_OFF:
|
|
return applied
|
|
|
|
on_cuda = getattr(target, "device", None) == "cuda"
|
|
family_allows_compile = bool(getattr(family, "supports_torch_compile", True))
|
|
|
|
# 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 on_cuda:
|
|
applied["cudnn_benchmark"] = _enable_cudnn_benchmark(logger)
|
|
|
|
# --- the compile lever, remapped per tier ----------------------------------------
|
|
# default = LIGHT compile: for a GGUF model, compile ONLY the dequant op chain
|
|
# (~70-80% of eager GGUF time) -- cheap, VRAM-free, resolution-invariant; the
|
|
# transformer block stays eager. A dense model has no dequant, so default falls
|
|
# back to the regional block compile (its only compile lever).
|
|
# max = FULL compile: regional max-autotune compile of the repeated denoiser block
|
|
# (fuses dequant + matmul + norm + elementwise in one graph). It subsumes the
|
|
# dequant fusion, so we do NOT also install the standalone compiled dequant here.
|
|
# eager = no compile at all.
|
|
if mode == SPEED_DEFAULT:
|
|
if is_gguf and on_cuda and family_allows_compile:
|
|
applied["compiled_dequant"] = gguf_compile.install_compiled_dequant(logger)
|
|
elif compile_eligible(target, is_gguf = is_gguf, family = family):
|
|
applied["compiled"] = _compile_repeated_blocks(
|
|
pipe, logger, max_autotune = False, cache_active = cache_active
|
|
)
|
|
elif mode == SPEED_MAX and compile_eligible(target, is_gguf = is_gguf, family = family):
|
|
applied["compiled"] = _compile_repeated_blocks(
|
|
pipe, logger, max_autotune = True, cache_active = cache_active
|
|
)
|
|
|
|
if mode == SPEED_MAX:
|
|
# Near-lossless: TF32 matmul (CUDA only) trades a few mantissa bits for speed.
|
|
if on_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,
|
|
cache_active: 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.
|
|
#
|
|
# fullgraph drops to False when a step cache is engaged: FBCache's per-step decision is
|
|
# ``@torch.compiler.disable``d, i.e. a graph break, which fullgraph=True rejects ("Skip
|
|
# inlining torch.compiler.disable()d function"). The break is cheap and the rest of the
|
|
# block still compiles.
|
|
kwargs: dict[str, Any] = {"fullgraph": not cache_active, "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
|
|
|
|
|
|
def _enable_tf32(logger: Any) -> bool:
|
|
try:
|
|
import torch
|
|
|
|
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 _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)
|