Skip to content

ROCm performance: disable MIOpen, which compiles a kernel per shape. - #518

Merged
remsky merged 2 commits into
remsky:masterfrom
s-kerdel:fix/rocm-miopen-per-shape-compilation
Sep 5, 2026
Merged

ROCm performance: disable MIOpen, which compiles a kernel per shape.#518
remsky merged 2 commits into
remsky:masterfrom
s-kerdel:fix/rocm-miopen-per-shape-compilation

Conversation

@s-kerdel

@s-kerdel s-kerdel commented Aug 31, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Problem

MIOpen compiles a kernel per unseen tensor shape. Kokoro's decoder length is derived from predicted durations, so it changes with the input text and nearly every request compiles.

gfx1100 (RX 7900 XTX), ROCm 7.2, torch 2.9.1, 6 novel texts per run:

per generation repeat of same text
MIOpen on 2433 ms 92 ms
MIOpen off 268 ms 218 ms

That is 9x faster on new text, which is the normal case.
MIOpen only wins when the exact same text repeats and its shape is already compiled.

Profile of one 2.26s generation: torch.instance_norm is 1.57s over 70 calls, which dispatches to MIOpen batch-norm on ROCm.

Fix

torch.backends.cudnn.enabled = False on ROCm, guarded by torch.version.hip so CUDA, CPU and MPS return early. KOKORO_ENABLE_MIOPEN=1 restores the old behaviour.

Why warmup_miopen.py does not cover this

It sweeps phoneme counts 1-340, but shape is not a function of phoneme count:
20 random texts at a fixed word count gave 20 distinct output lengths.

I also got no cross-process cache reuse. After clearing both MIOpen dirs and warming, a new process ran the same texts cold (687 ms vs 690 ms). A full FIND_MODE=3 plus FIND_ENFORCE=3 pass then a restart made no difference either.
May differ on gfx1151, where the #454 recipe came from.

Output unchanged

Audio output is not changed.
Sample count and RMS match, and the on/off correlation (0.9941) is the same as between two MIOpen runs (0.9954).

Changes

  • kokoro_v1.py: disable MIOpen on ROCm, with KOKORO_ENABLE_MIOPEN override
  • test_kokoro_v1.py: 3 tests, patching torch.version.hip so no GPU needed
  • docs/troubleshooting.md: warmup section scoped behind KOKORO_ENABLE_MIOPEN=1
  • docs/configuration.md: new variable in the env table

pytest api/tests: 315 passed. ruff format and ruff check clean.

Caveat

Tested on gfx1100 only, and on a self-built ROCm 7.2 image rather than the repo's docker/rocm (ROCm 6.4.4).
The change is only in kokoro_v1.py behind the HIP check, so it does not depend on that image.
This flips a default for all AMD users based on one card, so numbers from other hardware are welcome.
Happy to make it opt-in instead.

Note

Investigation and drafting were AI-assisted, per AGENTS.md.
The measurements are from my own hardware and I reviewed, tested and understand my change.

@s-kerdel
ROCm performance: disable MIOpen, which compiles a kernel per shape.
9f90472 
Kokoro's decoder length varies with the input text, so MIOpen recompiles on
nearly every request. Measured on gfx1100 / ROCm 7.2: 2433ms per generation
with MIOpen, 268ms without. Gated on torch.version.hip; KOKORO_ENABLE_MIOPEN=1
restores the old behaviour.

Updates the ROCm troubleshooting section, whose warmup procedure now only
applies when MIOpen is explicitly re-enabled.
@remsky

remsky commented Sep 1, 2026 •
edited
Loading

Copy link
Copy Markdown
Owner

Default off makes sense to me as the more common case. Thanks for the PR.

  1. Can we make it a settings field instead, enable_miopen: bool = False in config.py, then settings.enable_miopen in kokoro_v1.py; then the var can become ENABLE_MIOPEN like the other settings

  2. Did the MIOpen-on baseline have MIOPEN_FIND_MODE=2 set? Mostly just curious

@s-kerdel

s-kerdel commented Sep 1, 2026

Copy link
Copy Markdown
Contributor Author

Happy to contribute!

Good suggestion, I've implemented the settings field: enable_miopen: bool = False in config.py, read as settings.enable_miopen, documented as ENABLE_MIOPEN. Both paths verified on gfx1100 / ROCm 7.2.

Regarding FIND_MODE: yes, MIOPEN_FIND_MODE=2 was set on both sides of the benchmark. It comes from the ENV in the ROCm Dockerfile, so my image inherits it.

It does not help, for two reasons:

  • FIND_MODE only picks which algorithm to use. It does not skip compiling the kernel, and the compile is the slow part.
  • The prebaked kernels in the image cover CDNA and gfx1030. There is nothing for gfx1100.

I also measured the reuse question directly, same test script for every row:

Case MIOpen ON MIOpen OFF
Novel sentence 2859 ms 257 ms
Same sentence again, same process 91 ms 219 ms
Novel sentence, new process, warm on-disk cache 2400 ms -

So the cache does work, it just needs the identical shape. MIOpen is genuinely faster once warm, 91ms against 219ms. But staying warm means repeating the exact input text, and a second process only got 2859ms down to 2400ms from the cache on disk, so ordinary traffic almost never lands on a warm shape.

@remsky

remsky commented Sep 2, 2026

Copy link
Copy Markdown
Owner

Looks good to me and the logic checks out, I'm just going to courtesy ping some of the previous contributors/AMD users who've submitted PR's or issues, as I don't have the hardware myself.

Not sure if they're still actively using, but if no objections pop up will merge in on the next release. Thanks again.

@realugbun @kprinssu @Cirvaazny @projects-land

@kprinssu

kprinssu commented Sep 2, 2026

Copy link
Copy Markdown
Contributor

Looks good to me, I have since sold my MI50s and have switched over to Nvidia.

@remsky
remsky merged commit 5fc3ba3 into remsky:master Sep 5, 2026
2 checks passed
remsky added a commit that referenced this pull request Sep 5, 2026
@remsky
docs(changelog): backfill 0.8.2 entries for #518, #520, #522, #523
a7c7d61
github-actions Bot pushed a commit to autoeditor-video-or-audio/kokoro-tts-server that referenced this pull request Oct 1, 2026
@semantic-release-bot
chore(release): 0.6.0 [skip ci]
09889b4 
# [0.6.0](v0.5.0...v0.6.0) (2026-10-01)

### Bug Fixes

* **api:** turn off CORS credentials for starlette 1.x ([abcb9cf](abcb9cf))
* **audio:** drop WAV close()-time junk trailer that caused chunk-end click ([remsky#467](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/467)) ([7649d86](7649d86)), closes [remsky#463](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/463)
* **audio:** implement audio file source swapping for stream playback and seeking ([remsky#475](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/475)) ([2153529](2153529))
* **audio:** improve error handling for aborted audio streams ([106cbd4](106cbd4))
* **audio:** read finalize buffer before close except for ogg, fixes flac tail loss ([remsky#497](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/497)) ([ef13810](ef13810))
* **changelog:** update model validation details for consistency and error handling ([3eb9606](3eb9606))
* **chunking:** replace regex sentence splitter with UAX [remsky#29](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/29) segmentation ([remsky#308](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/308), adapts [remsky#415](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/415)) ([remsky#520](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/520)) ([d2aae10](d2aae10))
* **ci:** run the test job with the cpu extra ([c6398b9](c6398b9))
* **debug-endpoints:** gate /debug/* introspection routes behind ENABLE_DEBUG_ENDPOINTS, documentation ([316aa67](316aa67))
* **dependencies:** update fastapi to 0.116.1 and pin starlette dependency to avoid CVE ([c2687ec](c2687ec))
* **dev-unload:** Prevent POST /dev/unload behind an opt-in setting ([remsky#483](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/483)) ([24ac58e](24ac58e)), closes [remsky#478](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/478)
* **docker:** aligns cu128 base image with wheel (12.9.1 → 12.8.1), minimum toolkit with Blackwell (sm_120) support (CUDA 12.8) ([0d6b4d8](0d6b4d8)), closes [remsky#306](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/306) [remsky#320](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/320) [remsky#389](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/389) [remsky#443](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/443) [remsky#365](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/365)
* **docker:** compose defaults to idempotent model download; harden download_model scripts, static release artifact link ([0cf239e](0cf239e))
* **docker:** honour HOST/PORT in entrypoint for IPv6 binds ([remsky#408](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/408)), download_model defaults; drop api dev mount cruft ([remsky#522](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/522)) ([21a78e0](21a78e0))
* **download_model:** checksum verification for downloaded model files, avoid clobbering pre-existing model files ([aab59af](aab59af))
* honour default_voice_code in _split_multi_voice (issue [remsky#514](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/514)) ([remsky#515](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/515)) ([5ba470b](5ba470b))
* int16 overflow in silence-trim, invalid-escape in URL regex ([64cb068](64cb068))
* **mappings:** /v1/audio/voices returns id/name objects ([remsky#462](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/462)) to align with OpenWebUI ([7bbd3a8](7bbd3a8))
* **mappings:** update voice mappings for consistency, add missing audio file for 'bf_isabella' ([remsky#480](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/480)) ([c5ccfa1](c5ccfa1))
* **normalizer:** 'F' is normalised to "farad" instead of "fahrenheit" ([remsky#526](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/526)) ([f901ed8](f901ed8))
* **normalizer:** anchor decimal regex to prevent backtracking on digit floods ([af6e444](af6e444))
* **normalizer:** expand apostrophe-less "'re" contractions ([7746fbe](7746fbe)), closes [remsky#488](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/488) [remsky#488](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/488)
* **normalizer:** expand wh-word 're contractions espeak voices as -ray ([remsky#488](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/488)) ([f3a88ab](f3a88ab))
* **normalizer:** handle double dashes in text normalization and improve token processing ([4f826d6](4f826d6))
* **normalizer:** handle version-like numbers before NUMBER_PATTERN splits ([a736569](a736569))
* **normalizer:** move range sub before NUMBER_PATTERN to fix minus handling ([6c5773d](6c5773d))
* **normalizer:** treat underscore as a word boundary in caps detection ([123580c](123580c))
* **paths:** allow operator symlinks in search roots, keep traversal containment ([9dac665](9dac665))
* **paths:** harden _find_file to stay contained to its search roots ([0cb3a0d](0cb3a0d))
* phonemes default voice code ([remsky#516](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/516)) ([9138d89](9138d89))
* preserve paragraph breaks as sentence boundaries in normalizer ([remsky#525](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/525)) ([fbbaa75](fbbaa75)), closes [remsky#519](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/519)
* prevent ReDoS in text normalizer (URL/email/acronym patterns) ([remsky#489](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/489)) ([5fb4753](5fb4753))
* **rocm:** persist MIOpen cache and add FIND_MODE=2 default +warmup script ([remsky#469](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/469)) ([8ce0821](8ce0821))
* **schema, pause overflow:** caps total pause duration, input length, volume_multiplier, validates lang_code ([remsky#511](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/511)) ([316409d](316409d))
* **server_addons:** run KOKORO_SYNC_ON_START sync inside the upstream lifespan ([39f1f6d](39f1f6d)), closes [#1](#1)
* **ssml, web:**  SSML root validation, rate validation in VoiceSelector ([5d54647](5d54647))
* **ssml:** keep block boundaries, drop <desc>, add provider corpus, translate documents with an xml prolog, refuse DTDs, 400 on deep nesting,  cap nesting depth ([ffb4040](ffb4040))
* **streaming:** close audio writer and generator on client disconnect ([remsky#288](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/288), [remsky#337](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/337)) ([remsky#533](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/533)) ([cf8d71e](cf8d71e))
* **tests:** lazy import integration reqs, exclude on CI ([94f729f](94f729f))
* **tuning:** decode reference clips in a child process, serialized tune requests ([e498565](e498565))
* **uv direct run deps:** pyopenjtalk-plus on native Windows, avoids MSVC ([remsky#508](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/508)); pending swapover after multi-platform image testing ([8b0a51c](8b0a51c))
* **voices:** combine endpoint accepts weighted syntax, matching speech endpoints ([remsky#285](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/285)) ([remsky#512](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/512)) ([fdbcf96](fdbcf96))
* **web download:**  voice-timestamp as download names, associated url handling ([5ba5190](5ba5190))
* **web:** add Firefox-safe audio playback fallback when MSE MP3 is unsupported ([5e55b5c](5e55b5c))
* **web:** anchor the cast menu to its Options button, float it in the top layer ([86662e3](86662e3))
* **web:** assorted UI tweaks, responsiveness, animation ([91cd44d](91cd44d))
* **web:** drop auto-insert voice tag, enter key to save name, update related tests ([1b66dde](1b66dde))
* **web:** filtering, edit mode opens name as well ([db46865](db46865))
* **web:** firefox/supported playback handling and msging, bmac link, waveform bugfix ([a10fc75](a10fc75))
* **web:** pin playbar to bottom on slim view, responsive fixes, redesign tweaks ([328b816](328b816))
* **web:** playback slider step size; smoothness ([89850a9](89850a9))
* **webui:** bounded buffer to resolve 10 minute crashout, player state short ciruit on no-up updates to preserve responsiveness on long sessions. ([remsky#471](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/471)) ([3f64dc2](3f64dc2))
* **web:** voice search dropdown keys, clearable alias name, focus restore guard ([4aa5f1a](4aa5f1a))
* word timestamps for non-English (espeak) voices in captioned_speech ([remsky#484](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/484)) ([c0acb91](c0acb91)), closes [remsky#331](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/331)

### Features

* **api, web ui:** include voice grades on the voice listing endpoint ([remsky#517](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/517)) ([5fb71ea](5fb71ea))
* **api:** add POST /dev/unload to release model from GPU VRAM ([remsky#474](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/474)) ([ff6efaf](ff6efaf)), closes [remsky#473](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/473)
* **api:** optional model auto-unload after idle timeout ([remsky#506](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/506)) ([8307b0e](8307b0e))
* **benchmarks:** add model unload benchmarking script and update README with VRAM reclaim details ([6b58b53](6b58b53))
* **gpu:** adds clear cu128 opt-in via compose, setup for Blackwell ([f08279f](f08279f))
* **gpu:** adds clear cu128 opt-in via compose, setup for Blackwell / RTX 50-series release image ([65ef3b8](65ef3b8))
* improve voice selector and tag handling ([e65f1ba](e65f1ba))
* **multi-speaker:** add inline multi-speaker support, voice alias mappings, read-along scrobbing in web ui ([remsky#502](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/502)) ([5775958](5775958)), closes [remsky#272](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/272)
* **normalizer:** caps_normalization option, spoken phone numbers, hertz plurals, drop the 'S pass, emoji tag flags and joiners, web toggles ([remsky#353](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/353)) ([794a27b](794a27b))
* **normalizer:** opt-in remove_emoji, 400 before streaming on unspeakable input ([remsky#353](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/353)) ([c3cb020](c3cb020))
* **player:**  duration handling and update UI controls for audio playback ([a559254](a559254))
* **ssml, cleanup:** ENABLE_SSML kill switch, drop legacy settings, documentation of config variables) ([60d6862](60d6862))
* **ssml, speed/rate tag:** initial crack at integrating speed param per voice, ssml compliance ([8e4810a](8e4810a))
* **ssml:** adjust rate tag semantics to scale speaking pace and introduce baserate tag ([536bc76](536bc76))
* **ssml:** dedicated ssml router, shared rate clamp contract, alias rate visible + synced in web editor ([e282b40](e282b40))
* **ssml:** ssml flag on the speech endpoints, one call to translate and synthesize ([719f47a](719f47a))
* **time-to-first-audio:** lazy phonemization in get_sentence_info ([remsky#513](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/513)) ([61be98d](61be98d))
* **voice-clone-tune:** integrate Inno tuner, improve web ui ([e3049b2](e3049b2))
* **web,accessibility:** enhance keyboard navigation and ARIA attributes for improved accessibility ([8322a9b](8322a9b))
* **web:** add cast pins, resize insert button and ordering ([ffa52e1](ffa52e1))
* **web:** add duplicate to options, simplify suggested cast name rate to 3 digits ([fe9c7ae](fe9c7ae))
* **web:** add text normalization toggle to web UI settings ([remsky#523](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/523)) ([e57216c](e57216c))
* **web:** improve styling, responsiveness ([c53efd5](c53efd5))
* **web:** normalization options menu in player ([de18072](de18072))
* **web:** seamless auto swap-to-file ([c82946a](c82946a))
* **web:** show server version badge in the UI footer ([8e1bafe](8e1bafe))
* **web:** swap to the finished file as soon as it exists, allow minor skip on swap/ ([8cb69ac](8cb69ac))
* **web:** voice list sorted by grade then name, DEFAULT_VOICE honoured ([5fec3ad](5fec3ad))

### Performance Improvements

* **rocm:** disable MIOpen by default, add ENABLE_MIOPEN opt-in ([remsky#518](https://github.com/autoeditor-video-or-audio/kokoro-tts-server/issues/518)) ([5fc3ba3](5fc3ba3))
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@s-kerdel @remsky @kprinssu