Skip to content

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

Merged
adickin-amd merged 5 commits into
developfrom
users/jlichtne/ROCM-26918-fix-incomplete-winograd-rage-revert
Jun 24, 2026
Merged

adickin-amd merged 5 commits into
developfrom
users/jlichtne/ROCM-26918-fix-incomplete-winograd-rage-revert

Conversation

@JonathanLichtnerAMD

Copy link
Copy Markdown
Contributor

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 Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels… #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

…v (ROCM-26918)

Revert #6196 commented out the TransposedConvWinoRageRxS<2,3> registration in
solver.cpp (keeping a ++id placeholder so later solver ids stay stable) but left
that solver in the two Winograd Find solver lists in mlo_dir_conv.cpp. Under
NORMAL find, Find selects the unregistered solver and then solver::Id::GetAlgo()
(solver.cpp:198) throws miopenStatusInternalError, which PyTorch surfaces as a
RuntimeError on CvT's depthwise NHWC fp16 3x3 group convolution.

Remove TransposedConvWinoRageRxS<2,3> from both Find lists (GetWindogradSolvers
and GetWindogradWrWSolvers) to complete the revert and make the Find lists
consistent with the Id registry. No solver.cpp change is needed.

Also add a regression-guard gtest (CPU_SolverFindListRegistry_NONE) that asserts
every solver in the mlo_dir_conv.cpp Find containers is present in the Id
registry, so this class of find-list/registry drift cannot silently ship again.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@JonathanLichtnerAMD
JonathanLichtnerAMD requested a review from a team as a code owner June 23, 2026 16:33
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Keep the fix self-documenting without the internal ticket id; the defect is
referenced in the PR instead.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Comment thread projects/miopen/CHANGELOG.md Resolved

@BrianHarrisonAMD BrianHarrisonAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM.

The production Find paths that triggered the original failure
(ShrinkToFind10Results, EvaluateConvSolutions, FillFindReturnParameters in
convolutionocl.cpp) call solver::Id::GetAlgo() on found solver ids without an
IsValid() guard. Drop the IsValid() guard around the GetAlgo() call in the test
so it exercises that exact unguarded path instead of skipping it for the very
ids the test exists to catch.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@adickin-amd

Copy link
Copy Markdown
Contributor

issuing override merge. Confirmed with developer the change is low risk

@adickin-amd
adickin-amd merged commit a0300dc into develop Jun 24, 2026
68 of 72 checks passed
@adickin-amd
adickin-amd deleted the users/jlichtne/ROCM-26918-fix-incomplete-winograd-rage-revert branch June 24, 2026 14:26
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants