Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels… by JonathanLichtnerAMD · Pull Request #6196 · ROCm/rocm-libraries

JonathanLichtnerAMD · 2026-04-07T01:07:41Z

This reverts commit 7f4ad49.

Motivation

The Winograd kernels that were integrated caused several pytorch models to fail with memory access faults and so this change is being reverted.

Test Plan

One of the failing models was run with the revert, and it was verified that that model no longer hits the memory access faults.

Test Result

Manual testing on a failing model passed

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

ROCM-20758

…#2429)" This reverts commit 7f4ad49.

SreecharanGundaboluAMD

LGTM.

AnaghaRaoAMD · 2026-04-07T16:10:13Z

Would it be easier to return not applicable for all platforms, with bug ticket mentioned in comment?
(and update/disable tests, may also need some heuristics disabled)
Would be a smaller change and this would avoid reverting/removing files from 2 commits

JonathanLichtnerAMD · 2026-04-07T18:27:24Z

Note that this PR should also address ROCM-20758

JonathanLichtnerAMD · 2026-04-07T19:03:30Z

Would it be easier to return not applicable for all platforms, with bug ticket mentioned in comment? (and update/disable tests, may also need some heuristics disabled) Would be a smaller change and this would avoid reverting/removing files from 2 commits

Possibly, but that leaves temporarily dead code in the code base, and now if I make a change that touches the dead code, how do I test that? We can do special reverts, but that shouldn't be a common practice.

+ @cderb and @jdcampbe for the heuristics concern, and I'll also mention this to the Silo AI guys since this might affect the 2D heuristics that they're doing

jdcampbe · 2026-04-07T20:04:30Z

Would it be easier to return not applicable for all platforms, with bug ticket mentioned in comment? (and update/disable tests, may also need some heuristics disabled) Would be a smaller change and this would avoid reverting/removing files from 2 commits

Possibly, but that leaves temporarily dead code in the code base, and now if I make a change that touches the dead code, how do I test that? We can do special reverts, but that shouldn't be a common practice.

@cderb and @jdcampbe for the heuristics concern, and I'll also mention this to the Silo AI guys since this might affect the 2D heuristics that they're doing

Thanks for the heads up. I believe the data used for the AI models was collected before the original PR was merged so it should be alright (but we can check with Silo). This would affect some of the data we gathered. I'll see if it is necessary to ignore any of the data or update it.

AnaghaRaoAMD

we are going ahead with disabling transposed solver, even though it is still valid for fp16 case on gfx942 only. Once the bug is resolved we need to ensure transposed solver is added back.

This reverts commit ddb31d4.

ROCm#6196) … (ROCm#2429)" This reverts commit 7f4ad49. ## Motivation The Winograd kernels that were integrated caused several pytorch models to fail with memory access faults and so this change is being reverted. ## Test Plan One of the failing models was run with the revert, and it was verified that that model no longer hits the memory access faults. ## Test Result Manual testing on a failing model passed ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…v (ROCM-26918) (#8722) ## Motivation PyTorch CvT / FAMBench training on gfx942 (MI325) threw `RuntimeError: miopenStatusInternalError` on a depthwise NHWC fp16 3x3 grouped convolution. The crash traced to `solver::Id::GetAlgo()` in `src/solver.cpp` doing a registry lookup for a solver that no longer existed in the registry. Root cause: PR #6196 ("Revert Integrate several Wino Rage kernels") removed `TransposedConvWinoRageRxS<2,3>` from the solver Id registry (replacing it with a `++id;` placeholder to keep downstream solver Ids stable) but left that solver in both Winograd Find solver lists in `src/mlo_dir_conv.cpp` (`GetWindogradSolvers` and `GetWindogradWrWSolvers`). Under NORMAL find mode, Find selects the now-unregistered solver; the subsequent `GetAlgo()` call finds no registry entry and throws `miopenStatusInternalError`. Immediate and AI-heuristic find modes avoid the code path, which is why the bug only reproduced via PyTorch's DYNAMIC_HYBRID path (which falls back to NORMAL for uncached convolutions). ## Technical Details **Fix (`src/mlo_dir_conv.cpp`):** Remove `TransposedConvWinoRageRxS<2,3>` from `GetWindogradSolvers()` and `GetWindogradWrWSolvers()`, completing the revert started in PR #6196. The Rage kernels were originally reverted due to memory-access faults, so re-registration is not appropriate here. **Guard (`src/include/miopen/mlo_internal.hpp` + `src/mlo_dir_conv.cpp`):** Add `GetAllFindSolverDbIds()` — a new exported helper that enumerates every solver across all eight Find containers (`GetGemmSolvers`, `GetDirectSolvers`, `GetImplicitGemmSolvers`, `GetWindogradSolvers`, `GetImplicitGemmWrWSolvers`, `GetWindogradWrWSolvers`, `GetBwdWrW2DSolvers`, `GetFFTSolvers`) and returns their `SolverDbId` strings. This surfaces the complete Find-reachable solver set to tests without exposing internal container types. **Regression test (`test/gtest/solver_find_list_registry.cpp`):** `CPU_SolverFindListRegistry_NONE.AllFindSolversAreRegistered` iterates every id returned by `GetAllFindSolverDbIds()`, asserts `solver::Id::IsValid()` is true for each, and asserts `GetAlgo()` does not throw — the exact call that crashed in production. The test is CPU-only (no GPU required) and runs in NONE find mode, so it catches find-list/registry drift at build time rather than in a field workload. No other behavior is changed. The solver Id numbering in the registry is not disturbed (the `++id;` placeholder introduced by PR #6196 is left in place). ## Risk Assessment **Risk Level: 🟢 Low** ### Impacted Components - `projects/miopen/src/mlo_dir_conv.cpp` — two one-line removals from the Winograd Find solver lists - `projects/miopen/src/include/miopen/mlo_internal.hpp` — new exported declaration for `GetAllFindSolverDbIds()` - `projects/miopen/test/gtest/solver_find_list_registry.cpp` — new CPU-only gtest; no production code path ### Potential Side Effects Removing `TransposedConvWinoRageRxS<2,3>` from the Find lists means Find will never select it. This is intentional and was the desired outcome of PR #6196: the Rage kernels were reverted because they cause memory-access faults. Any workload that was silently falling back to another solver (due to applicability checks rejecting the Rage solver) will continue to do so unchanged. No performance regression is expected because the Rage solver was not usable in practice. ### Mitigation Steps - The fix is a strict subset of the already-merged PR #6196 revert; it does not introduce new code paths. - The new regression test will fail on any future find-list/registry drift, preventing this class of bug from re-occurring silently. - The fix was verified end-to-end on a gfx942 host before submission (see Test Result). ## Test Plan 1. **New regression gtest** (`CPU_SolverFindListRegistry_NONE.AllFindSolversAreRegistered`): builds without GPU, confirms every solver reachable through Find is registered and that `GetAlgo()` does not throw. 2. **Manual repro** on gfx942 with `MIOPEN_FIND_MODE=NORMAL`: ``` MIOPEN_FIND_MODE=NORMAL ./bin/MIOpenDriver convfp16 \ -n 256 -c 64 -H 56 -W 56 -k 64 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 \ --in_layout NHWC --fil_layout NHWC --out_layout NHWC -g 64 -F 1 ``` 3. **No Winograd-disable workaround applied** — the test was run with the full solver set available, exercising the exact code path that crashed in production. ## Test Result - The new gtest **FAILS on the unfixed tree**, naming `TransposedConvWinoRageRxS<2-3>` as the unregistered solver, and **PASSES after the fix**. - The `MIOpenDriver` repro command goes from exiting with code 5 (crash at `solver.cpp:198`) to `"Forward Convolution Verifies OK on GPU reference"` (exit 0) after the fix. - No known limitations or follow-up work. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> Co-authored-by: Matt Williams <matt.williams@amd.com>

## Motivation Previous Rage update #2429 was reverted by #6196 due to memory faults found on several models. Support for gfx12 was lost in Rage solver leading to a performance drop on gfx12. ## Technical Details The reverted PR was split into dedicated steps. This step brings gfx12 support back. ## Test Plan - `madengine` Tests - Unit Tests - Speedup Check for StableDiffusion - CI ## Test Result - `madengine` Tests | Index | Model | Topology | Launcher | Deployment | GPU Arch | Performance | Metric | Status | Duration | Data Name | Data Provider | |--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:| | 1 | pyt_detectron2_retinanet_singlenode | 1N×8G | docker | local | gfx942 | 2.03 | samples_per_second | ✅ Success | 19.4m | coco2017 | minio | | 2 | pyt_detectron2_retinanet_singlenode | 1N×8G | docker | local | gfx942 | 2.42 | samples_per_second | ✅ Success | 16.9m | coco2017 | minio | - unittest `miopen_gtest --gtest_filter=*UnitTestConvSolverWinoRage2x3*` ``` [==========] Running 17 tests from 4 test suites. [----------] Global test environment set-up. [----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16 [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0 warning: xnack 'Off' was requested for a processor that does not support it! [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0 (2483 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2 warning: xnack 'Off' was requested for a processor that does not support it! [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2 (37 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3 (1 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms) [----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16 (2526 ms total) [----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16 [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0 (10 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2 (18 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3 (2 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms) [----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16 (34 ms total) [----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16 [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0 (0 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1 (0 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2 (0 ms) [ RUN ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3 [ OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3 (1 ms) [----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16 (2 ms total) [----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE [ RUN ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0 [ OK ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0 (67 ms) [----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE (67 ms total) [----------] Global test environment tear-down [==========] 17 tests from 4 test suites ran. (2632 ms total) [ PASSED ] 17 tests. ``` - Speedup check for StableDiffusion The average gain on certain shapes is around 44% | N | C | H | W | K | R | S | Time (ms) develop→PR | GFLOPs develop→PR | Speedup | |--:|---:|--:|--:|----:|--:|--:|---------------------:|-------------------:|--------:| | 2 | 1152 | 32 | 32 | 768 | 1 | 1 | 0.206 → **0.201** | 18K → **18K** | **1.02x** | | 2 | 1152 | 32 | 32 | 768 | 3 | 3 | 0.317 → **0.205** | 103K → **159K** | **1.54x** | | 2 | 1536 | 32 | 32 | 1536 | 3 | 3 | 1.221 → **0.485** | 71K → **179K** | **2.52x** | | 2 | 1536 | 32 | 32 | 768 | 1 | 1 | 0.281 → **0.275** | 17K → **18K** | **1.02x** | | 2 | 1536 | 32 | 32 | 768 | 3 | 3 | 0.405 → **0.264** | 107K → **165K** | **1.53x** | | 2 | 2304 | 32 | 32 | 768 | 1 | 1 | 0.402 → **0.362** | 18K → **20K** | **1.11x** | | 2 | 2304 | 32 | 32 | 768 | 3 | 3 | 0.592 → **0.396** | 110K → **165K** | **1.50x** | | 2 | 384 | 32 | 32 | 768 | 1 | 1 | 0.080 → **0.077** | 15K → **16K** | **1.04x** | | 2 | 384 | 32 | 32 | 768 | 3 | 3 | 0.125 → **0.089** | 87K → **123K** | **1.41x** | | 2 | 768 | 32 | 32 | 768 | 3 | 3 | 0.219 → **0.142** | 99K → **153K** | **1.54x** | | 2 | 768 | 64 | 64 | 768 | 3 | 3 | 0.777 → **0.497** | 112K → **175K** | **1.56x** | | | | | | | | | | **Average** | **1.44x** | ## Submission Checklist - [x] `madengine` Tests - [x] Unit Tests - [x] Speedup Check for StableDiffusion - [x] CI - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Artem Kuzmitckii <artem.kuzmitckii@amd.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

JonathanLichtnerAMD self-assigned this Apr 7, 2026

JonathanLichtnerAMD requested review from a team as code owners April 7, 2026 01:07

github-actions Bot added documentation project: miopen labels Apr 7, 2026

assistant-librarian Bot added the organization: ROCm label Apr 7, 2026

JonathanLichtnerAMD requested review from AnaghaRaoAMD, SreecharanGundaboluAMD, randyspauldingamd and scerzh April 7, 2026 02:02

JonathanLichtnerAMD force-pushed the users/jlichtne/ROCM-20997-revert-winograd-pr2429 branch from 8e39d9b to 2203cb8 Compare April 7, 2026 02:07

Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels (…

7b805c6

…#2429)" This reverts commit 7f4ad49.

JonathanLichtnerAMD force-pushed the users/jlichtne/ROCM-20997-revert-winograd-pr2429 branch from 2203cb8 to 7b805c6 Compare April 7, 2026 02:36

SreecharanGundaboluAMD approved these changes Apr 7, 2026

View reviewed changes

Reformat

ad4a0cd

AnaghaRaoAMD reviewed Apr 7, 2026

View reviewed changes

Comment thread projects/miopen/test/gtest/unit_conv_solver_ConvWinoRageRxS.cpp Resolved

iq136boy approved these changes Apr 7, 2026

View reviewed changes

AnaghaRaoAMD requested changes Apr 7, 2026

View reviewed changes

Also remove the transpose based rage solver from the registry

b06192e

AnaghaRaoAMD reviewed Apr 7, 2026

View reviewed changes

Comment thread projects/miopen/src/solver.cpp Resolved

AnaghaRaoAMD self-requested a review April 7, 2026 19:18

AnaghaRaoAMD approved these changes Apr 7, 2026

View reviewed changes

JonathanLichtnerAMD added 2 commits April 7, 2026 20:34

TMP: limit threads for building docker

ddb31d4

Revert "TMP: limit threads for building docker"

a506e86

This reverts commit ddb31d4.

mattwill-amd approved these changes Apr 8, 2026

View reviewed changes

JonathanLichtnerAMD enabled auto-merge (squash) April 8, 2026 04:57

JonathanLichtnerAMD merged commit a691b15 into develop Apr 8, 2026
38 checks passed

JonathanLichtnerAMD deleted the users/jlichtne/ROCM-20997-revert-winograd-pr2429 branch April 8, 2026 07:18

scerzh mentioned this pull request Apr 21, 2026

feat(miopen): Add gfx120x support to Winograd Rage #6439

Merged

5 tasks done

JonathanLichtnerAMD mentioned this pull request Jun 23, 2026

[MIOpen] Fix miopenStatusInternalError on depthwise NHWC Winograd conv (ROCM-26918) #8722

Merged

Navigation Menu

Repository navigation

Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels… - #6196

Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels…#6196
JonathanLichtnerAMD merged 5 commits into
developfrom
users/jlichtne/ROCM-20997-revert-winograd-pr2429

JonathanLichtnerAMD commented Apr 7, 2026 •

edited by tony-davis

Loading

Uh oh!

SreecharanGundaboluAMD left a comment

Uh oh!

Uh oh!

AnaghaRaoAMD commented Apr 7, 2026

Uh oh!

JonathanLichtnerAMD commented Apr 7, 2026

Uh oh!

Uh oh!

JonathanLichtnerAMD commented Apr 7, 2026 •

edited

Loading

Uh oh!

jdcampbe commented Apr 7, 2026

Uh oh!

AnaghaRaoAMD left a comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Conversation

JonathanLichtnerAMD commented Apr 7, 2026 • edited by tony-davis Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Test Plan

Test Result

Submission Checklist

Uh oh!

SreecharanGundaboluAMD left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

AnaghaRaoAMD commented Apr 7, 2026

Uh oh!

JonathanLichtnerAMD commented Apr 7, 2026

Uh oh!

Uh oh!

JonathanLichtnerAMD commented Apr 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jdcampbe commented Apr 7, 2026

Uh oh!

AnaghaRaoAMD left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

JonathanLichtnerAMD commented Apr 7, 2026 •

edited by tony-davis

Loading

JonathanLichtnerAMD commented Apr 7, 2026 •

edited

Loading