[jdcampbe]

The grouped xdlops solvers reference SolverHeuristicConfig and HeuristicInitState unconditionally (including in the non-AI fallback path), but the include of ai_conv_nd_kernel_tuning_utils.hpp was gated behind MIOPEN_ENABLE_AI_KERNEL_TUNING and HeuristicInitState was defined inside the same guard in the header. With the flag OFF, these symbols were undeclared and the build failed.

Move the header include outside the guard in the six affected solver .cpp files and move HeuristicInitState out of the AI-only block in ai_conv_nd_kernel_tuning_utils.hpp. The struct holds only references and trivial accessors with no AI-runtime dependency.

Motivation

Building MIOpen with -DMIOPEN_ENABLE_AI_KERNEL_TUNING=OFF fails to compile. The grouped xdlops
convolution solvers reference AI-tuning helper types (SolverHeuristicConfig, HeuristicInitState) from
code paths that are not guarded by the flag — including the non-AI fallback path that runs regardless of
the flag's value. With the flag OFF, the necessary header is not included and the symbols are
undeclared.

This PR fixes the build so that MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF is a supported configuration again,
without changing any runtime behavior when the flag is ON (the default).

Technical Details

Two distinct gating bugs were introduced by the recent centralized AI heuristics refactor:

Bug 1 — Header include gated, but types used unguarded. In each of the six affected solver .cpp files,
the include of ai_conv_nd_kernel_tuning_utils.hpp was inside #if MIOPEN_ENABLE_AI_KERNEL_TUNING. But
SolverHeuristicConfig (used at file scope to declare each solver's constexpr config) and
HeuristicInitState (used inside HeuristicInit, including the fallback path) are referenced unguarded.
With the flag OFF the header was never parsed and the types were undeclared.

Bug 2 — HeuristicInitState defined inside the AI-only block. In ai_conv_nd_kernel_tuning_utils.hpp,
HeuristicInitState was inside the #if MIOPEN_ENABLE_AI_KERNEL_TUNING guard. The struct is a trivial
state holder (refs to valid_kernels/index/split_k/kernel_id plus Reset() and SetResult()) with no
AI-runtime dependency, and it is used by the non-AI fallback path at the end of HeuristicInit, so it
must compile both ways.

Changes:

• src/include/miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp — moved struct
  HeuristicInitState out of the #if MIOPEN_ENABLE_AI_KERNEL_TUNING block, placing it next to
  SolverHeuristicConfig in the unguarded namespace section. The AI-runtime symbols (RunAIHeuristics,
  RunParameterPredictionModel, etc.) remain inside the guard.
• The following six solver .cpp files — moved #include
  <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp> outside the #if. The other two AI headers
  (ai_heuristics.hpp, ai_candidate_selection.hpp) remain gated since they are only referenced from
  already-guarded code:
  • 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

No functional change when MIOPEN_ENABLE_AI_KERNEL_TUNING=ON (the default):
ai_conv_nd_kernel_tuning_utils.hpp was already being included in that path; this PR just makes the
include unconditional. AI-only code remains gated.

Test Plan

• Configure and build with -DMIOPEN_ENABLE_AI_KERNEL_TUNING=OFF (HIP backend) and confirm the previously
  failing translation units now compile cleanly.
• Configure and build with the default (MIOPEN_ENABLE_AI_KERNEL_TUNING=ON) to confirm no regression to
  the AI-enabled build.
• Run existing AI-related unit tests with the default ON build to confirm AI tuning behavior is
  unchanged: test/gtest/kernel_tuning_net.cpp, test/gtest/conv_ai_candidate_selection_model.cpp,
  test/gtest/conv_ai_nd_kernel_tuning_utils.cpp.
• Smoke-test the affected grouped/3D-grouped xdlops convolution solvers via MIOpenDriver on a supported
  architecture (gfx942/gfx950 or gfx90a) to confirm both the AI path (ON) and the fallback path (OFF)
  produce valid kernels.

Test Result

• Build successful with both MIOPEN_ENABLE_AI_KERNEL_TUNING ON/OFF.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.