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

projects/miopen/CHANGELOG.md

@@ -10,6 +10,7 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 ### Resolved Issues
 * [RNN] Fix RNN workspace tensor descriptor int overflow
 * [Conv] Enabled grouped Composable Kernel (CK) xdlops fwd, bwd, and wrw convolution (2D and 3D) for tensors whose strides exceed the int32 range.
+* [Conv] Fixed `miopenStatusInternalError` thrown by Find on depthwise NHWC grouped convolutions under `MIOPEN_FIND_MODE=NORMAL`.
 
 ## MIOpen 3.5.1 for ROCm 7.13.0
 ### Added

projects/miopen/src/include/miopen/mlo_internal.hpp

@@ -247,6 +247,13 @@ FindAllFFTSolutions(const miopen::ExecutionContext& ctx,
                     const miopen::conv::ProblemDescription& problem,
                     const miopen::AnyInvokeParams& invoke_ctx);
 
+// Returns the SolverDbId of every convolution solver listed in the Find solver
+// containers defined in mlo_dir_conv.cpp. Used by a unit test to assert that
+// every solver reachable by Find is also present in the solver Id registry,
+// guarding against find-list/registry drift.
+MIOPEN_INTERNALS_EXPORT
+std::vector<std::string> GetAllFindSolverDbIds();
+
 struct mlo_construct_base
 {
     mlo_construct_base(miopen::conv::Direction dir) : _problem(dir)

projects/miopen/src/mlo_dir_conv.cpp

@@ -121,7 +121,6 @@ static auto GetWindogradSolvers()
         miopen::solver::conv::TransposedConvBinWinogradRxS,
         miopen::solver::conv::TransposedConvBinWinogradRxSf2x3g1,
         miopen::solver::conv::TransposedConvWinoFuryRxS<2, 3>,
-        miopen::solver::conv::TransposedConvWinoRageRxS<2, 3>,
         miopen::solver::conv::TransposedConvMPBidirectWinograd<2, 3>,
         miopen::solver::conv::TransposedConvMPBidirectWinograd<3, 3>,
         miopen::solver::conv::TransposedConvMPBidirectWinograd<4, 3>,
@@ -181,7 +180,6 @@ static auto GetWindogradWrWSolvers()
         miopen::solver::conv::TransposedConvBinWinogradRxS,
         miopen::solver::conv::TransposedConvBinWinogradRxSf2x3g1,
         miopen::solver::conv::TransposedConvWinoFuryRxS<2, 3>,
-        miopen::solver::conv::TransposedConvWinoRageRxS<2, 3>,
         miopen::solver::conv::TransposedConvWinograd3x3MultipassWrW<3, 2>,
         miopen::solver::conv::TransposedConvWinograd3x3MultipassWrW<3, 3>,
         miopen::solver::conv::TransposedConvWinograd3x3MultipassWrW<3, 4>,
@@ -354,6 +352,23 @@ AllFFTForwardBackwardDataWorkspaceSize(const miopen::ExecutionContext& ctx,
     return GetFFTSolvers().GetWorkspaceSizes(ctx, problem);
 }
 
+std::vector<std::string> GetAllFindSolverDbIds()
+{
+    std::vector<std::string> ids;
+    const auto collect = [&](auto container) {
+        container.Foreach([&](auto solver) { ids.push_back(solver.SolverDbId()); });
+    };
+    collect(GetGemmSolvers());
+    collect(GetDirectSolvers());
+    collect(GetImplicitGemmSolvers());
+    collect(GetWindogradSolvers());
+    collect(GetImplicitGemmWrWSolvers());
+    collect(GetWindogradWrWSolvers());
+    collect(GetBwdWrW2DSolvers());
+    collect(GetFFTSolvers());
+    return ids;
+}
+
 void mlo_construct_activ_lrn_pooling_common::setupFloats()
 {
     if(_problem.GetInDataType() == miopenFloat && _problem.GetOutDataType() == miopenFloat)

projects/miopen/test/gtest/solver_find_list_registry.cpp

@@ -0,0 +1,40 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier:  MIT
+
+#include <gtest/gtest.h>
+
+#include <miopen/mlo_internal.hpp>
+#include <miopen/solver_id.hpp>
+
+// Regression guard against Find-list/registry drift.
+//
+// Every convolution solver reachable through the Find solver lists in
+// mlo_dir_conv.cpp must also be registered in the solver Id registry
+// (solver.cpp). If a solver is dropped from the registry (e.g. by a revert) but
+// left behind in a Find list, Find can still select it and MIOpen then throws
+// miopenStatusInternalError from solver::Id::GetAlgo() because the id is
+// unknown. This test fails fast on that drift instead of letting it ship and
+// surface as a runtime crash in a workload.
+TEST(CPU_SolverFindListRegistry_NONE, AllFindSolversAreRegistered)
+{
+    const auto db_ids = GetAllFindSolverDbIds();
+    ASSERT_FALSE(db_ids.empty()) << "Find solver lists returned no solvers; test wiring is broken.";
+
+    for(const auto& db_id : db_ids)
+    {
+        const miopen::solver::Id id{db_id};
+        EXPECT_TRUE(id.IsValid())
+            << "Solver '" << db_id
+            << "' is present in a Find solver list (mlo_dir_conv.cpp) but is not registered in the "
+               "solver Id registry (solver.cpp). This drift makes Find throw "
+               "miopenStatusInternalError.";
+
+        // A registered id must also resolve to an algorithm; GetAlgo() is the
+        // exact call that throws in the field when the id is missing.
+        if(id.IsValid())
+        {
+            EXPECT_NO_THROW((void)id.GetAlgo())
+                << "Solver '" << db_id << "': solver::Id::GetAlgo() threw.";
+        }
+    }
+}