Skip to content

[MIOpen] Fix MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF build in grouped-conv CK solvers - #7338

Closed
amd-bartgips wants to merge 2 commits into
developfrom
users/bartgips/ai_guard_fix
Closed

amd-bartgips wants to merge 2 commits into
developfrom
users/bartgips/ai_guard_fix

Conversation

@amd-bartgips

Copy link
Copy Markdown
Contributor

Motivation

Building MIOpen with MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF failed to compile with
error: unknown type name 'SolverHeuristicConfig' and
error: use of undeclared identifier 'HeuristicInitState' in six grouped-conv CK solver
files. Both types are declared inside the #if MIOPEN_ENABLE_AI_KERNEL_TUNING guard in
ai_conv_nd_kernel_tuning_utils.hpp (which is itself only #included under that guard),
but the solver files used them unconditionally — the constexpr SolverHeuristicConfig k*SolverConfig declaration in each file's anonymous namespace sat outside the guard, and
five of the six files also used HeuristicInitState outside the guard in HeuristicInit().

This broke any build configuration that disables AI kernel tuning (e.g. the build_bo
configuration used for performance profiling without the fdeep/AI dependency).

Related: PR #6834 fixed an analogous guard issue for MIOPEN_USE_COMPOSABLEKERNEL=OFF in
implicitgemm_ck_util.hpp.

Technical Details

Root cause: In all six solver files, the #include of
ai_conv_nd_kernel_tuning_utils.hpp is wrapped in #if MIOPEN_ENABLE_AI_KERNEL_TUNING,
so SolverHeuristicConfig and HeuristicInitState are not declared when the flag is
OFF. The constexpr SolverHeuristicConfig k*SolverConfig objects were declared in
anonymous namespaces outside that guard, causing the compile failure.

Fix:

  • Move each constexpr SolverHeuristicConfig k*SolverConfig declaration inside a
    #if MIOPEN_ENABLE_AI_KERNEL_TUNING / #endif block (together with its enclosing
    anonymous namespace).
  • For the five files where HeuristicInitState state(...), state.Reset(...), and
    state.SetResult(...) were also outside the guard, move them inside the existing
    #if block alongside RunAIHeuristics().
  • Add an #else branch that calls InitValidKernels(problem) +
    DefaultKernelFromList(ctx) directly for the no-AI case. Both functions are defined
    outside the guard and already implement the correct fallback behavior.
  • The ranked_gemm_3d_grp_{bwd,wrw}* kernel ranking vectors are intentionally left
    outside the guard — they are used by DefaultKernelFromList in the new #else path.
  • conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp had HeuristicInitState already
    correctly inside the guard; only its k3DFwdSolverConfig declaration needed moving.

Files changed:

  • src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
  • src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp
  • src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
  • src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp
  • src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp
  • src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp

Test Plan

  1. Build all six solver files with -DMIOPEN_ENABLE_AI_KERNEL_TUNING=OFF — must produce
    zero errors and zero warnings.
  2. Build all six solver files with -DMIOPEN_ENABLE_AI_KERNEL_TUNING=ON (default) —
    must produce zero errors and zero warnings (no regression).
  3. The #else (no-AI) fallback path calls InitValidKernels + DefaultKernelFromList,
    which is the same code path taken when the CK library fails to load; this path is
    already exercised by existing CI.

Test Results

Configuration Result
MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF (build_bo, gfx942, Release) ✅ All 6 files: 0 errors, 0 warnings
MIOPEN_ENABLE_AI_KERNEL_TUNING=ON (default build, gfx942, Release) ✅ All 6 files: 0 errors, 0 warnings

Submission Checklist

Building with MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF failed with "unknown
type name" errors because SolverHeuristicConfig and HeuristicInitState
are declared inside the #if MIOPEN_ENABLE_AI_KERNEL_TUNING guard in
ai_conv_nd_kernel_tuning_utils.hpp, but five solver files used them
unconditionally.

Fix: move the constexpr SolverHeuristicConfig k*SolverConfig declarations
and the HeuristicInitState/state.Reset/state.SetResult usages inside the
#if MIOPEN_ENABLE_AI_KERNEL_TUNING guard in each file. Add an #else branch
with a bare InitValidKernels + DefaultKernelFromList path for the no-AI
case. The ranked_gemm_3d_grp_{bwd,wrw}* vectors stay outside the guard as
they are consumed by DefaultKernelFromList in the #else path.

Verified: clean build with -DMIOPEN_ENABLE_AI_KERNEL_TUNING=OFF (build_bo),
zero errors and zero warnings.

Co-Authored-By: Claude Sonnet 4 (1M context) <noreply@anthropic.com>
@amd-bartgips amd-bartgips changed the title Fix MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF build in grouped-conv CK solvers [MIOpen] Fix MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF build in grouped-conv CK solvers May 12, 2026
@github-actions

github-actions Bot commented Jun 7, 2026

Copy link
Copy Markdown
Contributor

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

  • Add new commits
  • Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

@github-actions github-actions Bot added the Stale PR has no activity for 25+ days label Jun 7, 2026
@amd-bartgips

Copy link
Copy Markdown
Contributor Author

Already handled in #6800 , closing this one

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

organization: ROCm project: miopen Stale PR has no activity for 25+ days

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant