* 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>
276 lines
10 KiB
Python
276 lines
10 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
|
|
|
|
"""Unit tests for the opt-in diffusion speed layer (``diffusion_speed.py``).
|
|
|
|
Hermetic: torch is stubbed via ``sys.modules`` only where a path needs it, so the
|
|
gating logic and the best-effort applier run without a GPU or real diffusers.
|
|
"""
|
|
|
|
from __future__ import annotations
|
|
|
|
import sys
|
|
import types
|
|
|
|
import pytest
|
|
|
|
from core.inference.diffusion_speed import (
|
|
SPEED_DEFAULT,
|
|
SPEED_MAX,
|
|
SPEED_OFF,
|
|
apply_speed_optims,
|
|
compile_eligible,
|
|
normalize_speed_mode,
|
|
resolve_speed_mode,
|
|
restore_backend_flags,
|
|
snapshot_backend_flags,
|
|
)
|
|
|
|
|
|
def _target(
|
|
*,
|
|
device = "cuda",
|
|
dtype = "bfloat16",
|
|
compile_ok = True,
|
|
):
|
|
return types.SimpleNamespace(
|
|
device = device,
|
|
dtype = dtype,
|
|
supports_default_torch_compile = compile_ok,
|
|
)
|
|
|
|
|
|
def _family(*, compile_ok = True):
|
|
return types.SimpleNamespace(supports_torch_compile = compile_ok)
|
|
|
|
|
|
def _stub_torch(monkeypatch):
|
|
torch = types.ModuleType("torch")
|
|
torch.bfloat16 = "bfloat16" # _is_bfloat16 compares by identity then str fallback
|
|
torch.channels_last = "channels_last"
|
|
torch.backends = types.SimpleNamespace(
|
|
cuda = types.SimpleNamespace(matmul = types.SimpleNamespace(allow_tf32 = False)),
|
|
cudnn = types.SimpleNamespace(allow_tf32 = False, benchmark = False),
|
|
)
|
|
monkeypatch.setitem(sys.modules, "torch", torch)
|
|
return torch
|
|
|
|
|
|
# ── normalisation ─────────────────────────────────────────────────────────────
|
|
|
|
|
|
def test_normalize_speed_mode():
|
|
assert normalize_speed_mode(None) == SPEED_OFF
|
|
assert normalize_speed_mode("") == SPEED_OFF
|
|
assert normalize_speed_mode("MAX") == SPEED_MAX
|
|
with pytest.raises(ValueError):
|
|
normalize_speed_mode("ludicrous")
|
|
|
|
|
|
def test_resolve_speed_mode_gguf_auto_default():
|
|
# Unset (None) -> default for GGUF (near-lossless), off for dense.
|
|
assert resolve_speed_mode(None, is_gguf = True) == SPEED_DEFAULT
|
|
assert resolve_speed_mode(None, is_gguf = False) == SPEED_OFF
|
|
# An explicit value is honored verbatim, including an explicit opt-out to off.
|
|
assert resolve_speed_mode("off", is_gguf = True) == SPEED_OFF
|
|
assert resolve_speed_mode("max", is_gguf = True) == SPEED_MAX
|
|
assert resolve_speed_mode("max", is_gguf = False) == SPEED_MAX
|
|
|
|
|
|
# ── compile gating ────────────────────────────────────────────────────────────
|
|
|
|
|
|
def test_compile_eligible_requires_bf16_cuda_friendly(monkeypatch):
|
|
_stub_torch(monkeypatch)
|
|
# The happy path: bf16, CUDA, compile-friendly family.
|
|
assert compile_eligible(_target(), is_gguf = False, family = _family()) is True
|
|
# GGUF is now compile-eligible too (measured ~2.3x, PSNR ~37 dB vs eager).
|
|
assert compile_eligible(_target(), is_gguf = True, family = _family()) is True
|
|
# fp16 (non-bf16) is excluded.
|
|
assert compile_eligible(_target(dtype = "float16"), is_gguf = False, family = _family()) is False
|
|
# A family flagged not compile-friendly is excluded.
|
|
assert compile_eligible(_target(), is_gguf = False, family = _family(compile_ok = False)) is False
|
|
# No compile support (e.g. XPU/MPS) is excluded.
|
|
assert compile_eligible(_target(compile_ok = False), is_gguf = False, family = _family()) is False
|
|
|
|
|
|
# ── backend-flag snapshot / restore (TF32 / cudnn.benchmark leak guard) ────────
|
|
|
|
|
|
def test_snapshot_restore_backend_flags(monkeypatch):
|
|
torch = _stub_torch(monkeypatch)
|
|
snap = snapshot_backend_flags()
|
|
assert snap == {"matmul_tf32": False, "cudnn_tf32": False, "cudnn_benchmark": False}
|
|
# An opt-in max run flips the globals on...
|
|
torch.backends.cuda.matmul.allow_tf32 = True
|
|
torch.backends.cudnn.allow_tf32 = True
|
|
torch.backends.cudnn.benchmark = True
|
|
# ...and restore puts them back, so a later `off` load is bit-identical again.
|
|
restore_backend_flags(snap)
|
|
assert torch.backends.cuda.matmul.allow_tf32 is False
|
|
assert torch.backends.cudnn.allow_tf32 is False
|
|
assert torch.backends.cudnn.benchmark is False
|
|
|
|
|
|
def test_restore_backend_flags_tolerates_none():
|
|
restore_backend_flags(None) # no torch needed, no-op
|
|
|
|
|
|
def test_snapshot_partial_when_some_backends_missing(monkeypatch):
|
|
# A build/platform without cuda.matmul (e.g. CPU/MPS) must still snapshot + restore the
|
|
# flags it does have, rather than skipping the whole snapshot on one missing attribute.
|
|
torch = types.ModuleType("torch")
|
|
torch.backends = types.SimpleNamespace(
|
|
cuda = types.SimpleNamespace(), # no .matmul
|
|
cudnn = types.SimpleNamespace(benchmark = True), # no .allow_tf32
|
|
)
|
|
monkeypatch.setitem(sys.modules, "torch", torch)
|
|
snap = snapshot_backend_flags()
|
|
assert snap == {"cudnn_benchmark": True}
|
|
torch.backends.cudnn.benchmark = False
|
|
restore_backend_flags(snap)
|
|
assert torch.backends.cudnn.benchmark is True
|
|
|
|
|
|
def test_restore_is_independent_per_flag(monkeypatch):
|
|
# A read-only / failing attribute must not abort restoring the remaining flags.
|
|
torch = _stub_torch(monkeypatch)
|
|
|
|
class _NoMatmulSet:
|
|
@property
|
|
def allow_tf32(self):
|
|
return False
|
|
|
|
@allow_tf32.setter
|
|
def allow_tf32(self, value):
|
|
raise RuntimeError("read-only on this build")
|
|
|
|
torch.backends.cuda.matmul = _NoMatmulSet()
|
|
snap = {"matmul_tf32": False, "cudnn_tf32": False, "cudnn_benchmark": False}
|
|
torch.backends.cudnn.benchmark = True
|
|
restore_backend_flags(snap) # matmul setter raises, cudnn still restored
|
|
assert torch.backends.cudnn.benchmark is False
|
|
|
|
|
|
# ── applier ───────────────────────────────────────────────────────────────────
|
|
|
|
|
|
class _Pipe:
|
|
def __init__(
|
|
self,
|
|
*,
|
|
with_compile = False,
|
|
with_fuse = False,
|
|
) -> None:
|
|
self.vae = types.SimpleNamespace(mem_format = None, to = self._vae_to)
|
|
self.transformer = types.SimpleNamespace()
|
|
if with_compile:
|
|
self.transformer.compile_repeated_blocks = self._compile
|
|
if with_fuse:
|
|
self.fuse_qkv_projections = self._fuse
|
|
self.compiled = False
|
|
self.fused = False
|
|
|
|
def _vae_to(self, *, memory_format):
|
|
self.vae.mem_format = memory_format
|
|
|
|
def _compile(self, **kwargs):
|
|
self.compiled = True
|
|
self.compile_kwargs = kwargs
|
|
|
|
def _fuse(self):
|
|
self.fused = True
|
|
|
|
|
|
def test_speed_off_applies_nothing(monkeypatch):
|
|
torch = _stub_torch(monkeypatch)
|
|
pipe = _Pipe(with_compile = True, with_fuse = True)
|
|
applied = apply_speed_optims(
|
|
pipe, _target(), is_gguf = False, family = _family(), speed_mode = SPEED_OFF
|
|
)
|
|
assert applied == {
|
|
"channels_last": False,
|
|
"cudnn_benchmark": False,
|
|
"tf32": False,
|
|
"fused_qkv": False,
|
|
"compiled": False,
|
|
}
|
|
assert pipe.vae.mem_format is None and pipe.compiled is False
|
|
# off must not touch any process-wide flag (bit-identical reference path).
|
|
assert torch.backends.cudnn.benchmark is False
|
|
|
|
|
|
def test_speed_default_channels_last_compile_and_cudnn_benchmark(monkeypatch):
|
|
torch = _stub_torch(monkeypatch)
|
|
pipe = _Pipe(with_compile = True)
|
|
applied = apply_speed_optims(
|
|
pipe, _target(), is_gguf = False, family = _family(), speed_mode = SPEED_DEFAULT
|
|
)
|
|
assert applied["channels_last"] is True and pipe.vae.mem_format == torch.channels_last
|
|
assert applied["compiled"] is True and pipe.compiled is True
|
|
# default compiles with dynamic=True and no autotune mode (fast cold start,
|
|
# resolution-robust, sidesteps the CUDA-graph crash).
|
|
assert pipe.compile_kwargs == {"fullgraph": True, "dynamic": True}
|
|
# default also autotunes the VAE convs but does NOT flip TF32 or fuse QKV.
|
|
assert applied["cudnn_benchmark"] is True and torch.backends.cudnn.benchmark is True
|
|
assert applied["tf32"] is False and applied["fused_qkv"] is False
|
|
|
|
|
|
def test_speed_default_compiles_gguf(monkeypatch):
|
|
_stub_torch(monkeypatch)
|
|
pipe = _Pipe(with_compile = True)
|
|
applied = apply_speed_optims(
|
|
pipe, _target(), is_gguf = True, family = _family(), speed_mode = SPEED_DEFAULT
|
|
)
|
|
assert applied["channels_last"] is True
|
|
# GGUF now compiles (the big near-lossless win).
|
|
assert applied["compiled"] is True and pipe.compiled is True
|
|
|
|
|
|
def test_speed_default_cudnn_benchmark_only_on_cuda(monkeypatch):
|
|
_stub_torch(monkeypatch)
|
|
pipe = _Pipe(with_compile = True)
|
|
applied = apply_speed_optims(
|
|
pipe,
|
|
_target(device = "mps", compile_ok = False),
|
|
is_gguf = True,
|
|
family = _family(),
|
|
speed_mode = SPEED_DEFAULT,
|
|
)
|
|
assert applied["cudnn_benchmark"] is False # not CUDA -> no autotune flip
|
|
|
|
|
|
def test_speed_max_enables_tf32_and_fused_qkv(monkeypatch):
|
|
torch = _stub_torch(monkeypatch)
|
|
pipe = _Pipe(with_compile = True, with_fuse = True)
|
|
applied = apply_speed_optims(
|
|
pipe, _target(), is_gguf = False, family = _family(), speed_mode = SPEED_MAX
|
|
)
|
|
assert applied["tf32"] is True and torch.backends.cuda.matmul.allow_tf32 is True
|
|
assert applied["fused_qkv"] is True and pipe.fused is True
|
|
# max opts into autotuned kernels (static shapes); CUDA-graph modes are avoided.
|
|
assert pipe.compile_kwargs["mode"] == "max-autotune-no-cudagraphs"
|
|
assert pipe.compile_kwargs["dynamic"] is False
|
|
|
|
|
|
def test_speed_max_tf32_only_on_cuda(monkeypatch):
|
|
_stub_torch(monkeypatch)
|
|
pipe = _Pipe()
|
|
applied = apply_speed_optims(
|
|
pipe,
|
|
_target(device = "mps", compile_ok = False),
|
|
is_gguf = True,
|
|
family = _family(),
|
|
speed_mode = SPEED_MAX,
|
|
)
|
|
assert applied["tf32"] is False # not CUDA -> no TF32
|
|
|
|
|
|
def test_apply_tolerates_missing_optims(monkeypatch):
|
|
_stub_torch(monkeypatch)
|
|
# A bare pipe (no vae.to, no compile, no fuse) must not crash.
|
|
bare = types.SimpleNamespace(vae = None, transformer = types.SimpleNamespace())
|
|
applied = apply_speed_optims(
|
|
bare, _target(), is_gguf = False, family = _family(), speed_mode = SPEED_MAX
|
|
)
|
|
assert applied["channels_last"] is False and applied["fused_qkv"] is False
|