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

README.md

@@ -20,7 +20,8 @@ tags:
 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.
+difference is not the model — it is MIOpen tuning convolution kernels once, on a ROCm build
+that has no pretuned database for this GPU.
 
 ## Measured
 
@@ -40,17 +41,33 @@ different prompt and seed.
 **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
+## ⛔ Why: no MIOpen *tuning database* for gfx1151 on this ROCm build
 
 ```
 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.
+MIOpen ships pretuned performance databases per GPU architecture. On **this install
+(ROCm 7.13)** there is none for gfx1151: `/opt/rocm/share/miopen/db/` holds 81 files covering
+gfx908, gfx90a and gfx942, and **zero** matching `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.
+
+⭐ **Two separate things, and only one is missing here.** The CK grouped-convolution *kernel
+library* `libMIOpenCKGroupedConv_gfx1151.so` **is present** on this system. Only the *tuning
+database* is absent. That distinction matters: the related upstream report
+[ROCm/TheRock#5105](https://github.com/ROCm/TheRock/issues/5105) describes a worse case where
+**both** were missing and convolutions fell back to the `GemmFwdRest` solver, making an
+inference workload 3-5x slower than CPU. If you see
+`CK grouped conv library not found for device gfx1151` in addition to the fdb warning, you
+have that bug, not this one.
+
+⚠️ **Version scope.** These numbers are ROCm 7.13. AMD has since moved gfx1151 packaging
+forward (ROCm 10.0 publishes dedicated `device-gfx1151` payloads). **We have not tested
+ROCm 10 on this hardware**, so do not read this page as a claim about current ROCm — it is a
+measurement of one released version, and the *method* below is what generalises.
 
 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
@@ -71,6 +88,25 @@ stuck. It is not. Distinguish them by:
 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.
 
+### You can build the tuning database yourself
+
+A missing *system* performance database is not a dead end — MIOpen consults a user PerfDb
+that overrides it, and AMD documents generating one
+([tuning performance databases](https://rocm.docs.amd.com/projects/MIOpen/en/develop/conceptual/tuningdb.html)).
+Exercise your real shapes once with search enabled:
+
+```bash
+export MIOPEN_USER_DB_PATH="$HOME/.config/miopen"
+export MIOPEN_FIND_MODE=NORMAL
+export MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE
+python gen.py ...            # run the resolutions you actually use
+unset MIOPEN_FIND_MODE MIOPEN_FIND_ENFORCE
+```
+
+Subsequent runs read the tuned entries. This is what the warm run above is doing implicitly —
+explicit tuning just lets you front-load it deliberately, per shape, instead of paying it
+inside a user-facing request.
+
 ## The trap that silently costs you the GPU
 
 This machine's **system** `python3` already had a working ROCm PyTorch