feat(miopen): Cross-architecture AI heuristics for Navis (immediate mode)

[jdcampbe]

Motivation

MIOpen's immediate mode (MIOPEN_FIND_MODE=2) selects a convolution solver and
its performance config from hand-written per-arch heuristics or, on a find-db
miss, a generic default. This branch adds a two-layer, data-driven AI heuristic
trained on perf-DB data so immediate-mode selection generalizes across
architectures without per-arch rules:

• Layer 1 (solver selection): a cross-architecture LightGBM ranker picks
  which solver to run.
• Layer 2 (perf-config selection): once a solver is chosen, a per-solver
  LightGBM ranker picks that solver's performance config.

It also fixes two regressions found while validating on RDNA: solver selection
dropping to ConvDirectNaive on out-of-distribution shapes, and a build break
under the stricter develop -Werror policy.

JIRA ID : ALMIOPEN-1195

Technical Details

Layer 1 — cross-arch solver picker. ai::lgbm::PickSolverRanked
(src/conv/heuristics/lgbm_pick.cpp) scores the full solver vocabulary with a
Treelite-compiled LightGBM model and returns solver IDs ranked by predicted
speed; PredictSolver (src/conv/heuristics/ai_heuristics.cpp) short-circuits
the TunaNet path with this list, and GetSolutionsFallback
(src/hip/convolutionocl.cpp) walks it applying IsApplicable lazily. The
always-applicable ConvDirectNaiveConv{Fwd,Bwd,Wrw} fallbacks are in the model
vocab and were being over-ranked on OOD shapes, so the walk committed to naive
over a much faster real kernel. The ranker now demotes naive fallbacks below
every non-naive solver (relative score order preserved), gated by group count
(meta.NaiveGuardMaxGroups()) since naive is genuinely fastest for large group
counts. Naive membership is data-driven via rank.naive_fallback_solvers in the
model metadata, not hardcoded.

Layer 2 — per-solver perf-config picker. A new ai::lgbm::pcfg module
(lgbm_pcfg_pick.cpp, lgbm_pcfg_metadata.cpp) ships one Treelite-compiled
model per tunable solver plus a per-(gfx_id, direction, data_type) bucket
catalog of candidate configs. It hooks the perf-config seam in
FindSolutionImpl (src/include/miopen/find_solution.hpp), firing only for
conv problems when no perf_cfg was passed and the perf-db (+alt) missed and no
search ran — exactly where GetDefaultPerformanceConfig would otherwise be
used, so perf-db records stay authoritative. Rather than committing to the
bucket-argmax (globally good but frequently invalid for a specific OOD shape),
PickConfig returns candidates ranked best→worst and the hook walks them,
taking the first that passes IsValidPerformanceConfig:

for(const auto& desc : ranked)
{
    if(desc.empty()) break;              // model preferred the solver default
    PerformanceConfig cfg{};
    cfg.Deserialize(desc);
    if(s.IsValidPerformanceConfig(context, problem, cfg))
        return s.GetSolution(context, problem, cfg);
}
// exhausted -> falls through to GetDefaultPerformanceConfig (never left invalid)

Each model's generated C is compiled with per-solver symbol renames
(-Dpredict=lgbm_pcfg_<solver>_predict) to avoid ODR collisions; the picker
dispatches by solver name.

Diagnostics. Added MIOPEN_LOG_I2 at every silent decision point in the
AI/lgbm path (metadata-not-ready, gfx_id vocab code, scored/valid/dropped solver
counts, cache-hit short-circuit, applicable-vs-inapplicable counts, and the
function-local-static arch cached-vs-live device) so per-SKU selection
inconsistency (e.g. lgbm engaging on one RDNA SKU but not another) is
diagnosable from MIOPEN_LOG_LEVEL>=6 alone. Logging only; no behavior change.

Build fix. Three throw-only virtual bool stubs in pooling/solvers.hpp
(IsValidValue x2, SetNextValue) never return, which the develop -Werror
policy flags via -Wmissing-noreturn. They are virtual bool in the override
hierarchy so cannot be [[noreturn]]; added an unreachable return false;
after each throw. Pre-existing on develop; surfaces here because the branch is
built for an RDNA target under the stricter flags.

Gating unchanged: the whole feature is behind MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK
and MIOPEN_DEBUG_LGBM_PICK / MIOPEN_DEBUG_LGBM_PCFG; disabling either restores
the prior TunaNet/WTI/default behavior.

Test Plan

• CPU-only gtest fixtures (test/gtest/lgbm_picker.cpp,
  test/gtest/lgbm_pcfg_picker.cpp) replay exported test vectors and assert the
  C++ ranking reproduces the Python model's pick (top-1 and the full ranked head
  for the perf-config walk). No GPU required.
• Per-TU compile verification under the HIPNOGPU Docker image
  (MIOPEN_BACKEND=HIPNOGPU, MIOPEN_USE_COMPOSABLEKERNEL=Off) for the new
  modules, all per-solver generated C, the find_solution hook includers,
  hip/convolutionocl.cpp, and solver.cpp (with -Werror=missing-noreturn
  forced to confirm the pooling fix).
• Runtime validation: build for a target set including an RDNA arch (e.g.
  gfx1100) to exercise the layer-1 path; the OOD hardware perf-eval re-run to
  confirm the naive-selection and invalid-config regressions are resolved.

Test Result

• gtest parity fixtures reproduce the exported picks with zero mismatch (offline,
  CPU).
• Per-TU compile checks pass, including solver.cpp with -Werror=missing-noreturn
  explicitly forced (reproduces and clears the build break).
• Offline gold-val for the naive demotion: naive pick rate 51.6% -> 0.09%,
  coverage 0.51 -> 1.00, geomean on non-naive-best shapes unchanged.
• Full RDNA library build and the OOD hardware perf-eval re-run are pending; will
  update with results.

Risk Assessment

Risk level: Medium

• Impacted components: immediate-mode (FIND_MODE=2) conv solver + perf-config
  selection on the AI-fallback path (ai_heuristics.cpp, lgbm_pick.cpp, new
  ai::lgbm::pcfg module, the FindSolutionImpl perf-config seam,
  hip/convolutionocl.cpp fallback). Non-conv primitives are excluded via
  if constexpr. The pooling build fix touches only three throw-only stubs.
  Normal Find, perf-db hits, and non-AI paths are unaffected.
• Potential side effects:
  • Functional: on the AI path, solver and perf-config choices change relative to
    the prior heuristic (the intended improvement). IsValidPerformanceConfig
    validates every picked config; invalid or exhausted picks fall through to the
    existing default, so behavior degrades to prior rather than to an invalid
    kernel. The naive demotion changes which solver wins on OOD shapes.
  • Performance: per-call cost is a handful of Treelite predicts (sub-ms);
    perf-config candidates are capped (ranked walk depth <=16). Embedded model C
    increases binary size.
  • ABI/API: none (internal heuristics only).
• Mitigation steps:
  • Fully env-gated (MIOPEN_DEBUG_LGBM_PICK=0 / MIOPEN_DEBUG_LGBM_PCFG=0 and
    the MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK compile flag) — instant rollback to
    prior selection with no rebuild required for the env knobs.
  • Diagnostic logging added this branch makes per-SKU/per-arch behavior auditable
    before wider rollout.
  • Draft PR pending the RDNA build + OOD hardware perf-eval before it leaves
    draft; extra perf reviewers recommended for the selection-policy change.

Submission Checklist

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