Skip to content

fix(miopen): always register CK grouped-conv solvers in host (ROCM-27359) #9105

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
marbre merged 1 commit into from
Jul 6, 2026
Merged

fix(miopen): always register CK grouped-conv solvers in host (ROCM-27359) #9105

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 10 additions & 1 deletion projects/miopen/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -285,7 +285,16 @@ if( DEFINED MIOPEN_OVERRIDE_HIP_VERSION_PATCH )
message(STATUS "MIOPEN_hip_VERSION_PATCH overriden with ${MIOPEN_OVERRIDE_HIP_VERSION_PATCH}")
endif()

# Depend on Composable Kernels
# Composable Kernel (CK) is decoupled from the MIOpen *host* library at build time: the host never
# links CK device code -- it dlopens per-arch MIOpenCKGroupedConv_<arch>.so plugins at runtime
# (CkImplLibLoader) and degrades gracefully when a plugin is absent. Host-side CK solver
# registration and the dlopen prefetch are compiled UNCONDITIONALLY (no MIOPEN_USE_COMPOSABLEKERNEL
# guards remain in mlo_dir_conv.cpp / handlehip.cpp), so the grouped-conv XDLOPS solvers are always
# present in the device-neutral libMIOpen.so regardless of this flag. That is what prevents a
# CK-less host from silently disabling fast convolution on every arch (ROCM-27359) -- it is enforced
# in code, not by forcing this flag on. This flag now only gates the per-arch CK plugin build
# (src/ck_impl, itself arch-gated) and the CK-header-dependent gtest; the host configures and builds
# cleanly whether or not a CK package is found, so no per-arch build gating of this flag is needed.
option(MIOPEN_USE_COMPOSABLEKERNEL "Enable MIOpen to use composable kernels for various operations" On)
if(MIOPEN_BACKEND_OPENCL)
set(MIOPEN_USE_COMPOSABLEKERNEL OFF)
Expand Down
6 changes: 0 additions & 6 deletions projects/miopen/src/include/miopen/conv/solvers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4807,12 +4807,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvDepthwiseFwd3D final : ConvSolver
const miopen::conv::ProblemDescription&) const override;
};

// Test helper functions for metadata validation
// These functions return all CK kernel TypeStrings without problem-based filtering
// Declared here but implemented in the respective solver .cpp files
#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
#endif

} // namespace conv
} // namespace solver
} // namespace miopen
8 changes: 4 additions & 4 deletions projects/miopen/src/mlo_dir_conv.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,12 +87,12 @@ static auto GetImplicitGemmSolvers()
miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC,
miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC,
miopen::solver::conv::ConvCkIgemmFwdV6r1DlopsNchw,
#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
#if MIOPEN_BACKEND_HIP
miopen::solver::conv::ConvHipImplicitGemmGroupFwdXdlops,
miopen::solver::conv::ConvHipImplicitGemmGroupBwdXdlops,
miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops,
miopen::solver::conv::ConvHipImplicitGemm3DGroupBwdXdlops,
#endif // MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
#endif // MIOPEN_BACKEND_HIP
miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC>{};
}

Expand Down Expand Up @@ -146,10 +146,10 @@ static auto GetImplicitGemmWrWSolvers()
miopen::solver::conv::ConvMlirIgemmWrWXdlops,
miopen::solver::conv::ConvMlirIgemmWrW,
miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicWrwXdlops,
#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
#if MIOPEN_BACKEND_HIP
miopen::solver::conv::ConvHipImplicitGemmGroupWrwXdlops,
miopen::solver::conv::ConvHipImplicitGemm3DGroupWrwXdlops,
#endif // MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
#endif // MIOPEN_BACKEND_HIP
miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC>{};
}

Expand Down
Loading