Skip to content

Studio: avoid slow cold MIOpen searches on gfx1151 - #12359

Merged
danielhanchen merged 3 commits into
unslothai:mainfrom
oobabooga:fix/amd-miopen-search-cutoff
Oct 1, 2026
Merged

danielhanchen merged 3 commits into
unslothai:mainfrom
oobabooga:fix/amd-miopen-search-cutoff

Conversation

@oobabooga

Copy link
Copy Markdown
Member

Qwen-Image-2.1's first VAE decode can spend over 13 minutes benchmarking MIOpen convolution solvers on Radeon 8060S, while the next decode takes about two seconds. Cache isolation traced the delay to solver search: retaining the solver database removes it, but retaining only compiled kernels does not.

Fix

Set MIOPEN_SEARCH_CUTOFF=1 during Studio's existing lazy ROCm hardware detection, before model convolutions. MIOpen can skip excessively slow search candidates while retaining its naive fallback solver.

Enable the default only when every visible Torch GPU reports gfx1151. Preserve explicit environment values, including 0; mixed, unknown and other architectures keep their existing defaults. Users on other ROCm GPUs can opt in with the same environment variable. The setting is process-wide and inherited by workers.

Tests

Live A/B against main at 4e8503e437 on Radeon 8060S / gfx1151, Windows and Linux, PyTorch 2.11 / ROCm 7.13. Each OS runs stock and the actual patch on the same runner with separate empty MIOpen caches. Qwen-Image-2.1 Q4_K_M, Auto memory, 1024x1024, 40 steps, two prompts, seed 42, guidance 1.0 and a 19 GiB Torch allocation cap.

First generation Main This PR
Linux VAE decode 800.30 s 19.94 s
Windows VAE decode 939.37 s 23.86 s
Linux full generation 978.47 s 193.40 s
Windows full generation 1249.08 s 239.14 s

All four paired images and decoder input tensors are bit-identical. Cold peak Torch GPU allocation stays 14.92 GiB; warm decode stays around 1.5-2 seconds. The probes confirm that startup sets the flag automatically, with no external override in the patched cases. Load/download time is excluded. Denoiser timing also varied, so the VAE phase is the direct measure of this fix.

Live actual-patch A/B

Separate solver validation passed 468 CPU-reference forward/gradient comparisons across FP32, FP16, BF16, grouped/depthwise/pointwise/3D convolutions and multiple layouts. Solver validation

Rebased onto main at b1c522a4ea; 308 related hardware and startup tests passed, 36 skipped. The live A/B above was run on the stated earlier base and was not repeated after this rebase. Coverage includes explicit overrides, mixed and unknown architectures, probe failures, ROCm SDK version detection, CUDA exclusion and lazy startup. Ruff, formatting and diff checks passed.

Scope

Test hardware is Radeon 8060S / gfx1151 with shared memory. The allocation cap does not emulate a discrete GPU's transfer costs. RX 7900 XT remains untested. This changes cold solver-search behavior, not Auto placement, model-weight memory use or vision-attention correctness.

@oobabooga

Copy link
Copy Markdown
Member Author

Reached Codex review convergence at oobabooga#321.

@danielhanchen

Copy link
Copy Markdown
Member

Review record for #12359, carried over from the mirror it was reviewed on: danielhanchen/unsloth-staging-review#269.

Findings are quoted as posted and attributed to the account that posted them. The reactions shown are the triage recorded on the mirror; none were re-applied here. Commits are named as the mirror's, with this PR's equivalent where one was found.

Round 1: reviewed mirror bbfd543d65 (this PR's fbe234442)

chatgpt-codex-connector[bot] (on the mirror)

Codex Review: Didn't find any major issues. What shall we delve into next?

Reviewed commit: bbfd543d65

No triage reaction was recorded on the mirror

Verdict

The round converged at mirror bbfd543d65 (this PR's fbe234442).

This is the review as it stood at that commit. Anything pushed to this PR afterwards was not part of it.

@danielhanchen
danielhanchen merged commit 9aae51e into unslothai:main Oct 1, 2026
35 of 41 checks passed
@LeoBorcherding

LeoBorcherding commented Oct 1, 2026 •

Copy link
Copy Markdown
Collaborator

Tested on a discrete RDNA4 card with the gate forced on and off. The cold-decode problem isn't limited to gfx1151, so the cutoff should apply more widely.

First 1024² generation, empty MIOpen cache MIOPEN_SEARCH_CUTOFF unset =1
VAE decode 400.0 s 33.5 s
Full render 411.0 s 45.9 s

The images are near-identical (max pixel difference 2/255, mean 0.026), not bit-identical as on gfx1151. The stock numbers are from a rerun once the model was cached; an earlier stock run, confounded by the model download and a different memory plan, showed a similar 412 s decode. One clean run per side; the gap is far larger than run-to-run noise.

Setup: Radeon AI PRO R9700 (gfx1201), Windows 11, PyTorch 2.11 / ROCm 7.13, main at 364a42f0, Qwen-Image-2.1 Q4_K_M, Auto memory, 4 steps, seed 42, using full_probe.py from the amd-qwen21-memory-probe workflow with no allocator cap.

Recommend dropping the all-gfx1151 check in _configure_rocm_miopen (studio/backend/utils/hardware/hardware.py:1416 on main) so the cutoff defaults on for every ROCm GPU, still respecting an explicit value. I can open the follow-up.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants