[ialex2004]

Proposed changes

Kokoro-82M TTS inference on gfx1201 (Navi48) is dominated by the bidirectional nn.LSTM in the decoder. At batch=1 the per-timestep recurrence issues roughly 45k tiny serial GEMM + elementwise dispatches, leaving the GPU idle ~94% of the LSTM and making the workload launch-latency-bound rather than compute-bound.

This PR adds an experimental cooperative-grid fused single-layer LSTM forward-inference kernel (one launch per direction) that walks the whole sequence on-device with a grid-wide barrier per timestep. It reuses MIOpen's existing batched input-projection workspace and adds only the recurrent W_hh * h_{t-1} term and the LSTM cell update, collapsing the per-timestep host dispatch loop in RNNForwardInferencePacked.

Safety / scope

• Strictly opt-in behind MIOPEN_DEBUG_RNN_FUSED_INFERENCE. Default behavior is unchanged — the flag defaults off because enabling such fusion unconditionally could change results for other models.
• Gated to the validated case only: single layer, LSTM, miopenRNNdefault, fp32, no dropout, packed, linear input, uniform batch. Every other configuration falls back to the existing path.
• The kernel honors the caller's initial hidden/cell state (hx/cx) and writes the final hy/cy, matching the stock path. Uniform batch > 1 is supported.

Testing

test/gtest/lstm_fused_inference_test.cpp drives miopenRNNForwardInference twice in-process — flag off (stock reference) and flag on (fused) — and requires the GPU output plus hy/cy to agree to 1e-5. Because both paths share the input-projection GEMM and differ only in the recurrence, the tolerance is tight and any logic bug (gate order, cell update, direction offset, batch indexing, initial-state handling) trips the assertion. Cases cover unidirectional + bidirectional, N=1 and uniform batch > 1, fp32, including the single-timestep edge case.

The debug flag is registered in the MIOPEN_BUILD_TESTING env registry (env_debug.cpp) so ScopedEnvironment can toggle it from the test. Real Kokoro LSTM outputs match the stock path to ~1e-6 end-to-end.

All 12 gtest cases pass on gfx1201; they fail prior to the kernel fixes, confirming the test has teeth.

Related

JIRA ID : AIRADSW-142

• Companion PR (hipBLASLt gfx1201 Tensile tunings for the same Kokoro shapes): feat(hipblaslt): gfx1201 Tensile tunings for Kokoro-82M GEMM shapes #8916

Checklist

• I have added automated tests relevant to the introduced functionality
• I have sufficient test coverage for the changes, and code coverage hasn't decreased as a result of my PR
• I have ran the tests, and they are all passing locally
• I have added relevant documentation for the changes
• I have removed the stale documentation which is no longer relevant after this pull request
• I have ran make format & make check_format to ensure the changes have been formatted

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[BrianHarrisonAMD]

The changes are properly behind a flag, and ill approve to allow this to move forward.
I think the comments should be fixed before merging, or as a separate follow-up if this needs to land right away.

[ialex2004]

Addressed all seven review threads in f199760ce49.

Validation on Radeon RX 7900 XTX / ROCm 7.13:

• MIOpen core + miopen_gtest built successfully (CK disabled only because the container's installed CK headers are incompatible with current develop; the touched RNN path is independent).
• *LSTM_fused_infer*: 12/12 passed. The strengthened test clears/checks the exact fused-kernel cache key, proving the HIPRTC cooperative kernel actually dispatched.
• Serial paired Kokoro fp32 evaluation (5 OFF/ON pairs): OFF median 143.47 samples/min, ON median 143.22 samples/min. Performance is neutral within run variance; flag-off behavior remains unchanged.

[ialex2004]

Additional combined evaluation after deploying the freshly rebuilt current #8222 hipBLASLt (.so + nested gfx1100 .dat.zlib) with the reviewed #8915 MIOpen build:

• feat(miopen): fused single-layer LSTM forward-inference kernel (flagged) #8915 flag OFF: 144.312, 150.873, 145.350, 143.546, 145.139 → median 145.139 samples/min
• feat(miopen): fused single-layer LSTM forward-inference kernel (flagged) #8915 flag ON: 146.375, 146.262, 144.694, 144.797, 146.677 → median 146.262 samples/min

The +0.77% median difference is within observed run variance and paired deltas are mixed, so I consider the feature performance-neutral rather than claiming an uplift. There is no reliable regression, and the correctness/dispatch tests remain 12/12 passing.

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

[assistant-librarian]

TheRock Submodule Bump Activity

Newest first

• 2026-09-01 — Included in TheRock via ROCm/TheRock#7766.