From 51a98c9e39bc8a226987a10c54eaaba36bfa8e8e Mon Sep 17 00:00:00 2001 From: Vladimir Mandic Date: Wed, 18 Feb 2026 09:14:02 +0100 Subject: [PATCH] update Signed-off-by: Vladimir Mandic --- AMD-MIOpen.md | 264 ++++++++++++++++++++++++++++++++++++++++++++++++++ AMD-ROCm.md | 11 ++- 2 files changed, 271 insertions(+), 4 deletions(-) create mode 100644 AMD-MIOpen.md diff --git a/AMD-MIOpen.md b/AMD-MIOpen.md new file mode 100644 index 0000000..ed152c8 --- /dev/null +++ b/AMD-MIOpen.md @@ -0,0 +1,264 @@ +# Making Your SD.Next Faster with MIOpen Tuning + +If you're running SD.Next on an AMD GPU, this guide will help you unlock speed improvements that make your image generation noticeably faster. It also explains what MIOpen is and how it works under the hood, so the tuning process makes sense rather than feeling like magic. + +--- + +## What is MIOpen? + +MIOpen is AMD's open-source library for deep learning primitives — AMD's equivalent of NVIDIA's cuDNN. It provides optimized implementations of the operations that neural networks rely on: convolutions, pooling, normalization, activations, and more. + +When SD.Next generates an image, it doesn't talk to your GPU directly. The request travels through several layers: + +``` +┌─────────────────────────────────┐ +│ SD.Next │ ← User Interface +└────────────┬────────────────────┘ + │ +┌────────────▼────────────────────┐ +│ PyTorch (with ROCm backend) │ ← Deep Learning Framework +└────────────┬────────────────────┘ + │ +┌────────────▼────────────────────┐ +│ MIOpen Library │ ← Optimized GPU Kernels +└────────────┬────────────────────┘ + │ +┌────────────▼────────────────────┐ +│ AMD GPU (7900 XTX, 9070 XT…) │ ← Hardware +└─────────────────────────────────┘ +``` + +MIOpen sits between PyTorch and your GPU. Every time a convolution runs in the U-Net or VAE, MIOpen is the one deciding *how* to execute it. + +--- + +## How MIOpen Selects an Algorithm + +For any given operation, MIOpen has multiple **solvers** (algorithm implementations) that can produce the same result, but at very different speeds. For example, a single 3×3 convolution might be handled by: + +- **Winograd** — fast for small kernels, low memory usage +- **Implicit GEMM** — very fast on modern AMD GPUs using matrix cores +- **Direct GEMM** — general purpose, moderate speed +- **Naive Direct** — always works, extremely slow (100× slower in some cases) + +The "right" choice depends on your specific GPU, the tensor dimensions, and the data type. MIOpen's job is to pick the fastest one. + +### The Solver Selection Flow + +``` + You generate an image + │ + ▼ + ┌────────────────────────────────┐ + │ MIOpen receives GPU operation │ + │ (e.g. U-Net convolution) │ + └────────────────┬───────────────┘ + │ + ▼ + ┌────────────────────────────────┐ + │ Check local tuning cache │ + │ (~/.miopen/ or AppData) │ + └──────────┬─────────────────────┘ + │ + ┌──────────┴──────────┐ + │ │ + FOUND ✓ NOT FOUND ✗ + (warm run) (cold run) + │ │ + ▼ ▼ +┌─────────────────┐ ┌────────────────────────────┐ +│ Load optimal │ │ Benchmark all solvers: │ +│ algorithm from │ │ │ +│ cache instantly│ │ Winograd: 0.12 ms ✓ │ +└────────┬────────┘ │ Implicit GEMM: 0.15 ms │ + │ │ Direct GEMM: 2.15 ms │ + │ │ Naive Direct: 156.30 ms │ + │ │ │ + │ │ → Save winner to cache │ + │ └────────────────┬───────────┘ + │ │ + └─────────────┬───────────────┘ + │ + ▼ + ┌─────────────────────────────┐ + │ Execute optimal GPU kernel │ + └─────────────────────────────┘ +``` + +**This is why the first run is slow.** MIOpen has to benchmark every solver for every unique convolution it encounters — and a single image generation touches hundreds of different convolution configurations across the text encoder, U-Net, and VAE. Each benchmark takes a fraction of a second, but they add up. + +Once the cache is built, every subsequent run skips benchmarking entirely and jumps straight to the winner. + +--- + +## What to Expect: Before and After + +### First Run (Cold Start — Slow) + +- Your first generation will be **much slower than usual** — potentially several minutes +- The GPU will be active but no image progress is visible yet +- SD.Next may appear "stuck" +- Console messages like `Find Start`, `Evaluating Solver`, or `compiling kernels` are normal + +**This only happens once per unique operation shape.** A "shape" is defined by the combination of: +- The model (SD 1.5, SDXL, Flux, etc.) +- The image resolution (512×512 and 768×768 each need their own tuning) +- The batch size + +### After First Run (Warm Cache — Fast) + +- Image generation is noticeably faster +- Performance is consistent +- The cache is reused automatically on every subsequent launch — you don't need to redo it + +--- + +## How to Enable Tuning + +Set these two environment variables before launching SD.Next: + +### MIOPEN_FIND_MODE + +Controls *how* MIOpen searches for the best solver: + +| Value | Mode | Behavior | +|-------|------|----------| +| `2` | Fast | Uses heuristics — guesses a good solver without testing. Fast startup, suboptimal performance. | +| `3` | Hybrid | Checks the cache first; if not found, benchmarks all solvers and saves the winner. **Recommended.** | + +### MIOPEN_FIND_ENFORCE + +Controls *when* MIOpen is allowed to update the cache: + +| Value | Behavior | +|-------|----------| +| `1` | Never update the cache (read-only) | +| `3` | Update the cache when new results are found. **Recommended for tuning.** | + +### Commands + +**Linux — temporary (current session only):** +```bash +export MIOPEN_FIND_MODE=3 +export MIOPEN_FIND_ENFORCE=3 +``` + +**Linux — permanent (add to `~/.bashrc` or `~/.profile`):** +```bash +echo 'export MIOPEN_FIND_MODE=3' >> ~/.bashrc +echo 'export MIOPEN_FIND_ENFORCE=3' >> ~/.bashrc +source ~/.bashrc +``` + +**Windows — Command Prompt (temporary):** +```cmd +set MIOPEN_FIND_MODE=3 +set MIOPEN_FIND_ENFORCE=3 +``` + +**Windows — PowerShell (temporary):** +```powershell +$env:MIOPEN_FIND_MODE = "3" +$env:MIOPEN_FIND_ENFORCE = "3" +``` + +**Windows — Permanent (System Environment Variables):** + +1. Open **Start** → search **"Edit the system environment variables"** +2. Click **Environment Variables** +3. Under **User variables**, click **New** +4. Add `MIOPEN_FIND_MODE` = `3`, then repeat for `MIOPEN_FIND_ENFORCE` = `3` + +**Windows — Alternatively, edit `webui.bat`** and add the `set` lines before the launch command: +```bat +set MIOPEN_FIND_MODE=3 +set MIOPEN_FIND_ENFORCE=3 +.\python_embeded\python.exe -m streamlit run ... +``` + +Then launch SD.Next as normal. On Windows with ROCm: +```powershell +.\webui.bat --use-rocm +``` + +--- + +## The Tuning Cache + +MIOpen stores its results in a local database so tuning survives across restarts. + +**Default cache location:** +- **Linux:** `~/.config/miopen/` +- **Windows:** `%USERPROFILE%\.config\miopen\` + +The cache contains two file types: +- **`.udb`** (User Database) — stores the winning solver per operation. Used for fast lookup on every run. +- **`.ufdb`** (User Find Database) — stores full benchmark results for all tested solvers. Used for analysis. + +**You can override the cache location** with: +```bash +export MIOPEN_USER_DB_PATH=/path/to/your/cache +``` + +This is useful if you want to share a pre-built cache across multiple installs, or keep it on a fast drive. + +**To reset the cache** (forces re-tuning from scratch), delete the `.udb` and `.ufdb` files in the cache directory. + +--- + +## When Does Re-Tuning Happen? + +MIOpen caches results per unique operation shape. You'll see a slow first run again when: + +| Trigger | Why | +|---------|-----| +| New model (SD 1.5 → SDXL) | Different layer sizes = new convolution shapes | +| Different resolution | Tensor dimensions change | +| Different batch size | Affects solver selection | +| Cache was deleted | Fresh start | +| Model updated/modified | Some shapes may differ | + +Switching between two resolutions you've already tuned (e.g. 512 and 1024) will not re-tune — results for both are already cached. + +--- + +## Mode Comparison + +| | Find Mode 2 (Fast) | Find Mode 3 + Enforce 3 | +|---|---|---| +| **First run speed** | Normal | Slow (benchmarking) | +| **Subsequent speed** | Suboptimal | Optimal | +| **Cache written** | No | Yes | +| **Best for** | Debugging, stability issues | Production use | + +**Use Mode 2 if:** you're seeing crashes, driver timeouts, or unusually long hangs during the first run. Mode 2 never benchmarks, so it can work around instability. + +**Use Mode 3 if:** you want the best possible performance and can tolerate a slow first run per model/resolution. + +--- + +## Tips + +- **Be patient on first run.** With larger models (SDXL, Flux), the cold start can take 5–20 minutes. As long as GPU activity is visible in Task Manager or `rocm-smi`, it's working. +- **Tune at the resolution you actually use.** If you mostly generate at 1024×1024, run your first image at that size. A cache built at 512×512 won't cover the larger shapes. +- **Back up your cache.** Once built, the cache is reusable across SD.Next updates. Copy the `miopen/` directory somewhere safe. +- **Enable logging to watch progress.** Add `MIOPEN_LOG_LEVEL=4` temporarily to see solver benchmarking in the console. Remove it afterward to reduce noise. + +--- + +## Troubleshooting + +**"It's been 10+ minutes and nothing's happening!"** +This is normal for the first run with large models. Check that your GPU is active (`rocm-smi` on Linux, Task Manager on Windows). As long as GPU load or memory usage is nonzero, MIOpen is working. + +**"Second run is still slow"** +The cache may not have been written. Verify `MIOPEN_FIND_ENFORCE=3` is set (not just `MIOPEN_FIND_MODE`). Check that the cache directory exists and is writable. + +**"Performance got worse after a driver update"** +Driver updates can invalidate cached kernels. Delete the `.udb`/`.ufdb` files to force a fresh tuning session. + +**"It crashes during the first run"** +Switch to `MIOPEN_FIND_MODE=2` to skip benchmarking entirely. This avoids the long solver search that can trigger driver timeouts on some systems. + +**"I switched resolutions and it's slow again"** +Expected — each resolution has its own set of convolution shapes that need tuning. Run one generation at the new resolution to build the cache for it. diff --git a/AMD-ROCm.md b/AMD-ROCm.md index cfbef44..ade2a79 100644 --- a/AMD-ROCm.md +++ b/AMD-ROCm.md @@ -8,11 +8,12 @@ To use **AMD ROCm** with **SD.Next** you need to > AMD ROCm is officially supported for specific AMD GPUs. > [!IMPORTANT] -> Currently, PyTorch support on Windows is not officially maintained by PyTorch team. See [AMD's announcement](https://www.amd.com/en/resources/support-articles/release-notes/RN-AMDGPU-WINDOWS-PYTORCH-PREVIEW.html) for more information. +> Currently, PyTorch support on Windows is not officially maintained by PyTorch team. +> See [AMD's announcement](https://www.amd.com/en/resources/support-articles/release-notes/RN-AMDGPU-WINDOWS-PYTORCH-PREVIEW.html) for more information. > [!WARNING] > Unofficial support for other platforms is provided by the community and **SD.Next** does not guarantee it will work. -> Use of any third-party libraries is at your own risk. +> Use of any third-party libraries is at your own risk. - For preview support on Windows platform, see [ROCm on Windows](#rocm-on-windows) section. - For unofficial support for Windows platform, see [ZLUDA](https://github.com/vladmandic/sdnext/wiki/ZLUDA) page. @@ -99,7 +100,7 @@ Install ROCm SDK: sudo pacman -S rocm-hip-sdk libxml2-legacy gcc14 gcc14-libs ``` -### Running SD.Next with ROCm locally +### Running SD.Next with ROCm locally Open the terminal in a folder you want to install SD.Next and install SD.Next from Github with this command: ```shell @@ -149,7 +150,9 @@ sudo docker run -it \ ## ROCm performance tuning on Linux -### MIOpen database tuning +For more details, see [AMD-MIOpen Guide](AMD-MIOpen.md) page. + +### MIOpen database tuning On first use, when using a resolution for the first time, or when upgrading Pytorch versions, ROCm runs a series of benchmarks to select the most efficient approach. This can lead to slow (up to 5-8 minutes) startup times, in particular if you use a refine pass at high resolution, but it only happens once per resolution (subsequent runs are much faster). If for any case this is undesirable, you can set the environment variable `MIOPEN_FIND_MODE` to `FAST`. This will reduce a lot the startup time on first use at the price of worse performance when generating. On the other hand, for best performance during generation (but slower startup on first use), you can set the variable `MIOPEN_FIND_ENFORCE` to `SEARCH`.