Repository navigation
[MIOpen] Register TransposedConvWinoRageRxS in the Find solver lists (ALMIOPEN-2606) - #11949
Merged
DarylHawkinsAMD merged 7 commits intoSep 14, 2026
Merged
Conversation
TransposedConvWinoRageRxS<2,3> exists, is registered in the solver registry, is explicitly instantiated and has unit tests, but was absent from both SolverContainer lists in mlo_dir_conv.cpp. Find enumerates exactly those two containers, so the solver was never enumerated, never got a workspace size, and could never be selected. Its sibling TransposedConvWinoFuryRxS<2,3> is present in both lists.
The shipped guard checks Find-list -> registry. The defect fixed in this branch was the opposite direction: a solver registered in solver.cpp but absent from every Find solver list in mlo_dir_conv.cpp, which Find can never enumerate or select. AllRegisteredConvSolversAreFindReachable asserts that direction. It fails on the unpatched tree, reporting id 191 TransposedConvWinoRageRxS<2-3> as the only unreachable convolution solver, and passes with the fix.
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
DarylHawkinsAMD
deleted the
users/dahawkin/register-transposed-wino-rage
branch
September 14, 2026 16:44
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task
JonathanLichtnerAMD
pushed a commit
that referenced
this pull request
Oct 15, 2026
) ## Motivation The code behind these changelog entries merged after the ROCm 10.1 snap, so they did not ship in 10.1. They belong in the next release section, not in the released section. This moves them and changes no text. ## Technical Details Moves 2 entries to `MIOpen 3.6.2 for ROCm 10.2.0`, each under its matching category heading: - [#11949](#11949): [Conv] Fixed `TransposedConvWinoRageRxS<2,3>` being unreachable: it was registered in the solver ... - Merge commit `4084dbd194` is not in the snap (GitHub compare status `diverged`). - 1 line(s) the PR added to `projects/miopen/src/mlo_dir_conv.cpp` are missing from it at the pin, e.g. `miopen::solver::conv::TransposedConvWinoRageRxS<2, 3>,`. - [#11809](#11809): [Conv] Relaxed the ASM-GTC NHWC large-tensor guard for forward and backward-data. Those kernels ... - Merge commit `5a28fa27a6` is not in the snap (GitHub compare status `diverged`). - 3 line(s) the PR added to `projects/miopen/src/include/miopen/conv/asm_implicit_gemm.hpp` are missing from it at the pin, e.g. `const auto group_count = static_cast<std::size_t>(problem.GetGroupCount());`. If an entry should stay where it is, say so in a comment and it will be dropped from this PR. ## Test Plan Check each merge commit against the ROCm 10.1 snap (`release/therock-10.1`, `rocm-libraries` pin `39d8d60365d2272a5366f76ccb3bfba24b219eaf`): ``` git fetch origin 39d8d60 git merge-base --is-ancestor 4084dbd 39d8d60; echo $? # 1 means not in the snap git merge-base --is-ancestor 5a28fa2 39d8d60; echo $? # 1 means not in the snap ``` ## Test Result - Each merge commit is not an ancestor of the pin (compare status `diverged` or `ahead`). - The code each PR added is absent from the files at the pin (see Technical Details). - The patch applies cleanly to develop, and the non-blank lines of the changelog are identical before and after, so this is a pure move. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-libraries/blob/develop/CONTRIBUTING.md.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
TransposedConvWinoRageRxS<2,3>was registered in the MIOpen solver registry (id 191) andexplicitly instantiated, but was missing from both
SolverContainerlists inmlo_dir_conv.cpp. Find enumerates exactly those two containers, so the solver was neverenumerated, never received a workspace size, and could never be selected — and because
TransposingSolverreportsIsDynamic() == false, immediate mode could not reach it either.This PR adds it to both lists beside its sibling
TransposedConvWinoFuryRxS<2,3>, making theNHWC transposed Winograd Rage path reachable for forward, backward-data and backward-weights.
JIRA ID : ALMIOPEN-2606
Risk Assessment
Risk level 3. The diff is two lines of solver-list wiring, one regression test and a CHANGELOG entry — no solver
logic, applicability, kernel or perf-db change. But it widens the Find candidate set on every
NHWC convolution on
gfx942/gfx950/gfx120X(the solver's arch gate), which is a dispatch-surface change rather than pure plumbing. Measured effect on gfx950: cold Find latency +16–25 %
on affected NHWC problems, zero steady-state execution change, zero workspace change, and the
selected winner is unchanged in 20/20 paired trials on every case measured.
Related
ASIC Coverage
Support-surface-expanding: a solver that could previously never be selected can now be
enumerated and selected. The solver's arch gate is a prefix test on
gfx942,gfx950andgfx120(src/solver/conv/conv_wino_rage_RxS.cpp:317-322), so the affected families aregfx94X, gfx950 and gfx120X.
gfx1250is not covered — it does not match thegfx120prefix, and the Rage suites skip there rather than fail (verified by spoofing
MIOPEN_DEVICE_ARCH=gfx1250).Testing Summary
Baseline (
origin/developsnapshot) vs. this branch, two fresh identically-configured Ninjabuilds on MI355X (gfx950) / ROCm 10.1.0. Restricted to MIOpen source the two revisions differ
only by this fix, so upstream drift cannot be mistaken for it. All nine blocking exit criteria
passed.
Forcing the solver by name fails cleanly on the baseline with No suitable algorithm was
found — no HIP error, no crash — and verifies numerically here in all 4 cases.
AllRegisteredConvSolversAreFindReachablewalks registry →Find-list, the direction the shipped guard did not cover. On the baseline it reports
191 TransposedConvWinoRageRxS<2-3>as the only unreachable convolution solver (122registered vs. 121 reachable); here the two describe the same 122.
benchmarked, inserted into FindDb and numerically verified in all three directions.
MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3=0andMIOPEN_DEBUG_CONV_WINOGRAD=0each suppress the solver with a verifying fallback. NCHW is untouched — Not applicable, with
identical winners on both builds. Reported workspace always covers the selected solution.
skipped, 25 disabled, 0 failed), differenced by individual test identity: zero differences.
(+16.1 % to +24.9 %, disjoint IQRs) — the expected cost of benchmarking one more candidate on
an empty cache. Execution median within +0.5 %, workspace delta 0 B. Winner unchanged in 20/20
trials: Rage is benchmarked but never wins on gfx950, so Find correctly declines it.
The gfx950 baseline-vs-feature comparison above was run on the two-line fix alone, before the
regression test was added; the test is CPU-only and changes no product code.
Testing Checklist
test_solver_find_list_registry- Devices: gfx950 (MI355X) - Status: Passedtest_solver_find_list_registry --gtest_filter=*AllRegisteredConvSolversAreFindReachable*- Status: Pending (no MIOpen build on the authoring host; verified as a standalone probe on gfx950, where it fails on the baseline and passes with the fix)Flags / Guardrails
No new flag. The solver stays behind the two existing kill switches,
MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3andMIOPEN_DEBUG_CONV_WINOGRAD, both verified tosuppress it.
Technical Changes
projects/miopen/src/mlo_dir_conv.cpp: addTransposedConvWinoRageRxS<2, 3>toGetWindogradSolvers()(forward, backward-data) andGetWindogradWrWSolvers()(backward-weights), beside the existing
TransposedConvWinoFuryRxS<2, 3>.projects/miopen/test/gtest/solver_find_list_registry.cpp: addAllRegisteredConvSolversAreFindReachable, asserting every registered convolution solverappears in a Find list.
projects/miopen/CHANGELOG.md: record the fix under Resolved Issues.