[MIOpen] Add cross-architecture LGBM-based immediate-mode solver picker

[jdcampbe]

Motivation

MIOpen's immediate-mode solver selection uses per-architecture TunaNet models, which need a separately trained model per arch and cover only a subset of the GPU lineup. This PR adds a single cross-architecture solver picker driven by a gradient-boosted (LightGBM) rank model that selects a convolution solver from one model across gfx90a, gfx942, gfx950, gfx1151, gfx1201, and the gfx110x family. The picker reads everything it needs about the GPU directly from the device at runtime, so it requires no curated per-SKU data and stays correct on the actual silicon.

Technical Details

The picker is a self-contained inference path with no runtime dependency on libLightGBM or Treelite. The rank model is compiled to plain C via Treelite --source-only and checked in under projects/miopen/src/conv/heuristics/lgbm_models/rank/; each generated translation unit's exported symbols (predict, quantize, predict_unit0..7, the metadata accessors, the is_categorical global) are renamed with an lgbm_rank_ prefix via per-file -D defines in projects/miopen/src/CMakeLists.txt to keep them out of MIOpen's global namespace. The project enables the C language (project(MIOpen C CXX)) to compile these.

Inference (conv/heuristics/lgbm_pick.cpp):

• Reads gfx_id from Handle::GetDeviceName(); abstains if the arch is not in the model vocab / arch-constants table.
• Builds a 51-feature row, scores every solver in the vocabulary, and returns the argmax. There is no candidate masking, margin gate, or applicability second model.
• Maps the chosen name to a solver::Id and abstains if this MIOpen build doesn't know it.

GPU features are sourced two ways:

• Live device values via the Handle: cu_count, wave_size, l2_cache_total_kb, boost_clock_mhz, lds_size_per_workgroup_kb, vram_bytes. This adds two small Handle getters — GetL2CacheSize() and GetClockRateKhz() — implemented for HIP (hipDeviceGetAttribute), OpenCL (clGetDeviceInfo), and stubbed for NOGPU, mirroring the existing attribute-query getters in src/hip/handlehip.cpp.
• Arch-invariant values via a compiled-in gfx_id → constants table (conv/heuristics/lgbm_models/gpu_constants.h).

Model metadata (lgbm_model_meta.json) installs to DATABASE_INSTALL_DIR and loads at runtime via GetSystemDbPath(). The hook is an early short-circuit in ai_heuristics.cpp::PredictSolver, gated by MIOPEN_DEBUG_LGBM_PICK (on by default) under MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK, mirroring the existing TunaNet dispatch.

The 6 derived workload features (flop_cnt/bytes_*/gflops/bandwidth_gbps) are fed as NaN: the values these were trained on are direction-aware in a way a runtime textbook estimate cannot reproduce, and validation showed NaN reproduces more reference picks than an estimate.

LightGBM treats NaN as a real branch direction.

Not changed: the TunaNet path and WTI fallback. When the picker abstains (unknown arch, or model recommends an unavailable solver) it falls through to the existing selection path unchanged.

Test Plan

• CPU-only gtests in projects/miopen/test/gtest/lgbm_picker.cpp (target test_lgbm_picker):
  • CPU_LgbmMetadata — metadata loads, categorical encoding resolves, solver code matches vocab index.
  • CPU_LgbmPickerFixture.ReproducesReferenceArgmax — feeds each encoded feature row from lgbm_test_vectors.json through the scoring path and asserts the argmax solver matches the reference exactly.
• Runtime check on a real GPU: build MIOpen and run a conv with MIOPEN_LOG_LEVEL=6, and confirm the lgbm picked …
  log line.

Test Result

• Reference-fixture replay reproduces the expected argmax solver on 50/50 vectors.
• Per-TU compile is clean for the picker, metadata loader, regenerated rank C, and the NOGPU handle. The HIP and OpenCL handle TUs are not built in the NOGPU backend, but the new getters mirror the existing attribute-query getters and the required hipDeviceAttribute* / CL_DEVICE_* enums were confirmed present.
• A real-GPU smoke run are pending hardware; will update with results.

Risk Assessment

Risk level: Low

• Impacted components: new files under conv/heuristics/lgbm_* and conv/heuristics/lgbm_models/; the PredictSolver hook in ai_heuristics.cpp; two new Handle getters across the HIP/OpenCL/NOGPU backends; src/CMakeLists.txt (sources, install, C-language enable); a new gtest. The TunaNet and WTI paths are untouched, and behavior is identical when MIOPEN_DEBUG_LGBM_PICK is disabled.
• Potential side effects:
  • Functional: when enabled, immediate-mode picks on supported archs come from this model instead of TunaNet; unknown archs abstain and fall through.
  • Performance: one cached model load plus one predict per solver candidate; sub-millisecond per call, negligible versus tuning/search.
  • ABI/API: none (the new Handle getters are additive internal methods).
• Mitigation steps:
  • Gated behind MIOPEN_DEBUG_LGBM_PICK (runtime toggle, no rebuild) and compiled only under MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK.
  • The shipped lgbm_test_vectors.json fixture makes picks reproducible and checkable pre-merge.
  • Self-contained and easily revertible (new files plus one guarded hook and additive getters).

Submission Checklist

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

[astrelsky]

Motivation

MIOpen's immediate-mode solver selection uses per-architecture TunaNet models, which need a separately trained model per arch and cover only a subset of the GPU lineup. This PR adds a single cross-architecture solver picker driven by a gradient-boosted (LightGBM) rank model that selects a convolution solver from one model across gfx90a, gfx942, gfx950, gfx1151, gfx1201, and the gfx110x family.

This is still only a subset of the supported GPU lineup though, unless I misunderstand something.

[jdcampbe]

This is still only a subset of the supported GPU lineup though, unless I misunderstand something.

You're correct. In theory it should work across all architectures, but there is an artificial gate to include architectures that we have some data for. The artificial gates will be removed as we test on other architectures.

[yassinsolim]

I went through the latest CI run (v16, commit 47cf542). Good news is the failures are all mechanical or packaging related, nothing in the model logic itself. Here is the breakdown.

1. miopen_gtest fails to link (breaks HIP Package, gfx94X, gfx950, and gfx1151)

All four of those red checks die at the same place, linking bin/miopen_gtest:

ld.lld: error: undefined symbol: miopen::ai::lgbm::LgbmMetadata::Get()
ld.lld: error: undefined symbol: miopen::ai::lgbm::LgbmMetadata::CategoricalCode(...)
ld.lld: error: undefined symbol: miopen::ai::lgbm::LgbmMetadata::SolverCode(...)
ld.lld: error: undefined symbol: miopen::ai::lgbm::ScoreRowArgmaxForTest(...)

libMIOpen compiles fine, so this is a symbol visibility problem. The library is built with CXX_VISIBILITY_PRESET hidden and VISIBILITY_INLINES_HIDDEN (src/CMakeLists.txt), so anything not explicitly exported stays hidden, and lgbm_picker.cpp cannot resolve those symbols at link time. The declarations in lgbm_metadata.hpp and lgbm_pick.hpp are missing MIOPEN_INTERNALS_EXPORT. ai_candidate_selection.hpp is a good reference since it annotates its test-facing methods for exactly this reason.

Marking LgbmMetadata::Get, CategoricalCode, SolverCode, and ScoreRowArgmaxForTest with MIOPEN_INTERNALS_EXPORT should clear the package build and all three arch builds in one shot.

2. clang-tidy

Two separate things:

• ai_heuristics.cpp:749 trips readability-redundant-preprocessor (nested redundant #if). The new #if MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK block is already inside a block guarded by the same macro.
• The generated Treelite C files (lgbm_models/rank/tu0.c through tu7.c, quantize.c) fail under -warnings-as-errors='*' with readability-function-size (76k statements vs the 800 threshold) and bugprone-assignment-in-if-condition from the LIKELY/UNLIKELY macros. The -w flag silences the compiler but not clang-tidy, so the generated TUs need to be excluded from the tidy target, for example a .clang-tidy with Checks: '-*' dropped into the rank/ directory.

3. pre-commit

Still failing on trailing whitespace, end-of-file-fixer, black, and clang-format (mostly the generated model files plus gen_gpu_constants_header.py), and the GTest naming hook is rejecting the CPU_LgbmMetadata and CPU_LgbmPickerFixture fixtures. Running pre-commit run --all-files locally should clean up most of it.

A few design questions, separate from CI

Since this changes the default immediate-mode selection path on supported archs, I would like to understand a few things before it lands:

• The 6 derived workload features are fed as NaN. Can you share the validation behind NaN reproducing more reference picks than a textbook estimate?
• The picker takes a pure argmax over the full solver vocab and only checks IsApplicable afterward in the hook. How often does the top pick get rejected and fall through to TunaNet, and is that cost acceptable on the hot path?
• The real GPU smoke validation is still marked pending. I would want to see that result before approving.

On the arch coverage point above, thanks for clarifying that the gating is intentional and temporary. It might be worth a short note in the code or PR description saying the gfx_id allowlist will expand as more architectures get validated, just so it does not read as a hard limitation later.

[jdcampbe]

This PR will be merged in with #9680