[JonathanLichtnerAMD]

JIRA ID: ROCM-27359

Motivation

MIOpen's host libMIOpen.so is device-neutral and has no build-time Composable Kernel (CK) dependency: it never links CK device code and instead dlopens per-arch MIOpenCKGroupedConv_<arch>.so plugins at runtime (CkImplLibLoader), degrading gracefully when a plugin is absent.

However, host-side registration of the grouped-conv XDLOPS solvers in mlo_dir_conv.cpp was gated on MIOPEN_USE_COMPOSABLEKERNEL. When a build shard configured the host with MIOPEN_USE_COMPOSABLEKERNEL=OFF (for example a CK-unsupported-only arch), the preprocessor stripped those solvers out of the solver registry entirely. Because the host ships as a single device-neutral _generic artifact, a CK-less libMIOpen.so could become the host for every arch and disable fast convolution everywhere — including gfx942 — producing a roughly 97% performance regression observed in a nightly (ROCM-27359).

This is a backport cherry-pick of upstream commit c534dd1 (PR #8988) onto release/therock-7.14. The functional fix — always registering grouped-conv XDLOPS solvers under MIOPEN_BACKEND_HIP alone — is fully present. See the Mitigation section for the one hunk that was intentionally dropped during backport.

Technical Details

The MIOPEN_USE_COMPOSABLEKERNEL guard is removed from the host solver-registration paths so that the grouped-conv XDLOPS solvers are always compiled (still gated only on MIOPEN_BACKEND_HIP):

• projects/miopen/src/mlo_dir_conv.cpp: GetImplicitGemmSolvers() and GetImplicitGemmWrWSolvers() now register the grouped-conv XDLOPS solvers (ConvHipImplicitGemmGroupFwdXdlops, ConvHipImplicitGemmGroupBwdXdlops, ConvHipImplicitGemm3DGroupFwdXdlops, ConvHipImplicitGemm3DGroupBwdXdlops, ConvHipImplicitGemmGroupWrwXdlops, ConvHipImplicitGemm3DGroupWrwXdlops) under #if MIOPEN_BACKEND_HIP alone, dropping the && MIOPEN_USE_COMPOSABLEKERNEL condition.
• projects/miopen/src/include/miopen/conv/solvers.hpp: removes a now-empty #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL block and its associated comment (the block contained no declarations — it was dead code).
• projects/miopen/CMakeLists.txt: the comment above the MIOPEN_USE_COMPOSABLEKERNEL option is expanded to document this decoupling and explain that the flag now only gates the per-arch CK plugin build (src/ck_impl) and the CK-header-dependent gtest; the host builds cleanly regardless.

This is safe to compile and link with or without a CK package present: the solver classes are declared unconditionally in solvers.hpp, their implementations live in the unconditional MIOpen source list, and they link with no CK headers — they only use the runtime dlopen loader. MIOPEN_USE_COMPOSABLEKERNEL now only gates the per-arch CK plugin build and CK-header-dependent tests.

Backport note: The src/hip/handlehip.cpp hunk from the original upstream PR was intentionally dropped during this cherry-pick. The release/therock-7.14 base branch does not contain the PrefetchComposableKernel() helper or the <miopen/solver/ck_impl_lib_loader.hpp> include that the original commit modified — the CK-prefetch mechanism was never introduced on this branch. The dropped hunk was a startup-latency optimization only and is not part of the ROCM-27359 regression fix. The net diff therefore touches 3 files instead of the original 4.

Risk Assessment

Risk Level: 🟡 Medium

The diff is small (3 files, +14/-11 lines net) but touches the convolution solver registry — a foundational MIOpen path whose misregistration is exactly what caused the ~97% regression in ROCM-27359. The change direction is corrective (restoring solvers that should always be present), which bounds the risk. The dropped handlehip.cpp hunk further reduces the surface relative to the upstream PR.

Impacted Components

• MIOpen convolution solver registration (src/mlo_dir_conv.cpp) — grouped-conv XDLOPS Fwd/Bwd/WrW solvers, 2D and 3D.
• Public/internal solver declarations (src/include/miopen/conv/solvers.hpp) — empty guard block removed.
• Build configuration (CMakeLists.txt) — documentation only; no logic change.
• Indirectly: every arch served by the device-neutral host artifact, particularly gfx942 (the ROCM-27359 repro arch).

Potential Side Effects

• On a build with no CK plugin available at runtime, the grouped-conv XDLOPS solvers are now always registered; they must continue to degrade gracefully (return not-applicable / fall through) via the dlopen loader rather than failing. This is the intended and pre-existing runtime behavior of CkImplLibLoader, but reviewers should confirm no code path assumed the solvers were compiled out on CK-less builds.
• CK-enabled builds (MIOPEN_USE_COMPOSABLEKERNEL=ON) should be completely unaffected, since the guards being removed were already satisfied in that configuration.

Mitigation Steps

• The runtime loader (CkImplLibLoader) already handles absent plugins gracefully, so the always-registered solvers are inert when no plugin is present.
• The handlehip.cpp CK-prefetch hunk was not backported because it does not exist on this branch — this eliminates one change surface and one potential regression vector relative to the upstream PR.
• Revert plan: the change is self-contained to three host source files and can be reverted independently if a regression surfaces.
• Recommended: validate a MIOPEN_USE_COMPOSABLEKERNEL=OFF host build still resolves fast convolution on gfx942 (the ROCM-27359 repro) and confirm MIOPEN_USE_COMPOSABLEKERNEL=ON builds are unchanged.

Test Plan

• Build MIOpen host with MIOPEN_USE_COMPOSABLEKERNEL=OFF and confirm it compiles and links cleanly (the grouped-conv XDLOPS solvers are now unconditionally compiled and must link with no CK headers).
• Cross-artifact repro of ROCM-27359: build the host with MIOPEN_USE_COMPOSABLEKERNEL=OFF, then run a grouped convolution on gfx942 with the per-arch CK plugin supplied via MIOPEN_CK_LIB_PATH; confirm ConvHipImplicitGemmGroupFwdXdlops is registered, the plugin loads, and the fast path is selected (before this fix the solver was absent from the registry).
• Build with MIOPEN_USE_COMPOSABLEKERNEL=ON with a CK package present and verify no behavior change.
• No new automated tests are added; the fix is a build-time registration change validated through existing build-matrix configurations and the runtime cross-artifact repro above.

Test Result

To be confirmed by CI across the CK-supported and CK-unsupported build shards on release/therock-7.14. Expected outcome: the host library builds and links with CK ON or OFF; grouped-conv XDLOPS solvers are always registered so a CK-less device-neutral host no longer disables fast convolution (resolving the ROCM-27359 ~97% regression); CK-enabled builds are unchanged.

• multiarch sweep passes
• CI passes

[marbre]

LGTM. Excluding the original changes to projects/miopen/src/hip/handlehip.cpp make sense. Approving on behalf of dev ops.