Qwen-Image-2.1 on gfx1151: cold vs warm, MIOpen kernel-db finding

README.md

@@ -0,0 +1,140 @@
+---
+license: apache-2.0
+base_model: Qwen/Qwen-Image-2.1
+pipeline_tag: text-to-image
+library_name: diffusers
+tags:
+  - qwen-image
+  - rocm
+  - amd
+  - gfx1151
+  - strix-halo
+  - benchmark
+  - documentation
+---
+
+# Qwen-Image-2.1 on AMD Strix Halo (gfx1151, ROCm 7.13)
+
+**This repository contains no weights.** It records what it takes to run
+[`Qwen/Qwen-Image-2.1`](https://huggingface.co/Qwen/Qwen-Image-2.1) on an AMD Ryzen AI Max+
+395 (Radeon 8060S, gfx1151) under ROCm 7.13, and the measured cost of doing so.
+
+It runs. The first image takes **19.8 minutes** and the second takes **6.8 minutes**, and the
+difference is not the model.
+
+## Measured
+
+Radeon 8060S / gfx1151, ROCm 7.13, torch 2.10.0 (hip 7.13.99004), diffusers 0.41.0.dev0,
+bf16, 1024x1024, 24 steps, box otherwise idle. Two runs, same resolution and step count,
+different prompt and seed.
+
+| phase | run 1 (cold) | run 2 (warm) |
+| --- | ---: | ---: |
+| pipeline load to GPU | 7.3 s | — |
+| sampler, 24 steps | 365.0 s (15.22 s/it) | 390.4 s (16.28 s/it) |
+| **VAE decode** | **825.6 s** | **20.5 s** |
+| total (`GEN_TIME_S`) | **1190.6 s** | **410.9 s** |
+| peak GPU allocated | 40.55 GB | 36.89 GB |
+| peak GPU reserved | 42.82 GB | 38.60 GB |
+
+**The VAE decode got 40x faster on the second run.** Nothing about the model or the code
+changed between them.
+
+## ⛔ Why: ROCm ships no MIOpen kernel database for gfx1151
+
+```
+MIOpen(HIP): Warning [ParseAndLoadDb] File is unreadable:
+"/opt/rocm/share/miopen/db/gfx1151_20.HIP.fdb.txt"
+```
+
+MIOpen ships pretuned kernel databases per GPU architecture. ROCm 7.13 has none for gfx1151,
+so every convolution shape is **auto-tuned at runtime the first time it is seen**. The
+transformer is unaffected (it is GEMM work), which is why the sampler rate is stable across
+both runs. The VAE is convolutional, so it pays the entire tuning bill on run 1.
+
+The tuned results persist in `~/.cache/miopen`. Evidence that this is the mechanism: the
+cache was **221,184 bytes after run 1 and 221,184 bytes after run 2** — byte-identical, zero
+new tuning, and the decode collapsed from 825.6 s to 20.5 s.
+
+### It looks exactly like a hang
+
+For ~13 minutes there is no log output, one host thread sits at 100%, and the process appears
+stuck. It is not. Distinguish them by:
+
+| signal | tuning | genuinely hung |
+| --- | --- | --- |
+| `/sys/class/drm/card0/device/gpu_busy_percent` | **94-100%** | ~0% |
+| `~/.cache/miopen` size over a 20 s window | **growing** | static |
+| process state | `Rl` | `D` / `S` |
+
+⭐ **Do not quote a first-run timing as this model's speed on this hardware.** Warm the cache,
+then measure. If a first run must be quick, `MIOPEN_FIND_MODE=FAST` shortens the search at
+some cost in kernel quality — unset it for the run you actually report.
+
+## The trap that silently costs you the GPU
+
+This machine's **system** `python3` already had a working ROCm PyTorch
+(`torch 2.10.0`, `torch.version.hip 7.13.99004`, `torch.cuda.is_available() True`), while
+other virtualenvs on the box carried **CPU-only** torch. Letting pip resolve `torch` freshly,
+or reusing the wrong venv, runs the entire 33 GB pipeline on CPU without ever erroring.
+
+Build the venv so it inherits the working install, and check afterwards:
+
+```bash
+python3 -m venv --system-site-packages ~/build/qwen-image/venv
+~/build/qwen-image/venv/bin/python -c \
+  "import torch; assert torch.cuda.is_available(); print(torch.__version__, torch.version.hip)"
+```
+
+⭐ Make the generator **refuse to run on CPU** rather than fall back — that turns a silent
+30x slowdown into an immediate, obvious failure:
+
+```python
+if not torch.cuda.is_available():
+    print("FATAL: HIP not available - refusing to run on CPU"); sys.exit(2)
+```
+
+## Dependencies
+
+`QwenImage21Pipeline` is newer than any diffusers release on PyPI — it needs
+**diffusers >= 0.37.0.dev0**. Installed here from GitHub main via the archive zip (the build
+box has no `git`):
+
+```bash
+pip install "diffusers @ https://github.com/huggingface/diffusers/archive/refs/heads/main.zip"
+pip install "transformers>=5.17" accelerate safetensors
+```
+
+The text encoder is `Qwen3VLForConditionalGeneration`, which is why transformers 5.17+ is
+required. Component sizes: text_encoder 17.53 GB, transformer 14.23 GB
+(`QwenImage21Transformer2DModel`, 32 layers, 32 heads x 128), vae 1.35 GB — **33.1 GB** total
+in bf16, all of which is resident on the GPU during generation.
+
+## Files
+
+| file | what it is |
+| --- | --- |
+| `README.md` | this document |
+| `gen.py` | the generation script used for both runs (GPU-or-die, prints receipts) |
+
+## Reproduction
+
+```bash
+python3 -m venv --system-site-packages ~/build/qwen-image/venv
+. ~/build/qwen-image/venv/bin/activate
+pip install "diffusers @ https://github.com/huggingface/diffusers/archive/refs/heads/main.zip" \
+            "transformers>=5.17" accelerate safetensors
+
+HF_HUB_DISABLE_XET=1 hf download Qwen/Qwen-Image-2.1 --local-dir ~/models/qwen-image-2.1
+
+python gen.py --model ~/models/qwen-image-2.1 --out out/a.png \
+  --prompt "A photorealistic red-tailed hawk perched on a saguaro cactus at golden hour" \
+  --steps 24 --width 1024 --height 1024 --seed 42
+```
+
+Run it twice. The second run is the honest number.
+
+## Licence
+
+Apache-2.0. Weights belong to [Qwen](https://huggingface.co/Qwen/Qwen-Image-2.1) under their
+own licence; this repository contains only measurements and a script.