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

.pre-commit-config.yaml

@@ -55,6 +55,16 @@ exclude: |
     projects/miopen/.*\.s$|
     projects/miopen/.*\.inc$|
     projects/miopen/.*\.model$|
+    # Machine-generated LGBM assets: native LightGBM text models + baked
+    # catalog/metadata + golden parity fixtures (large, not human-authored;
+    # exceed check-added-large-files and are not lintable/end-of-file-fixable).
+    projects/miopen/src/conv/heuristics/lgbm_models/.*|
+    projects/miopen/src/conv/heuristics/lgbm_pcfg_models/.*|
+    projects/miopen/src/kernels/lgbm_pcfg_catalog\.json|
+    projects/miopen/src/kernels/lgbm_rank_model\.txt|
+    projects/miopen/src/kernels/lgbm_pcfg_.*_model\.txt|
+    projects/miopen/src/kernels/lgbm_rank_golden\.json|
+    projects/miopen/src/kernels/lgbm_pcfg_golden\.json|
     # RPP exclusion list
     projects/rpp/utilities/test_suite/TEST_IMAGES/three_images_yuv/.*|
     projects/rpp/utilities/test_suite/REFERENCE_OUTPUT/yuv_to_rgb/.*|

projects/miopen/CMakeLists.txt

@@ -793,6 +793,12 @@ enable_cppcheck(
     SUPPRESS
         ConfigurationNotChecked
         constStatement
+        # cppcheck cannot resolve a transitive header for the LGBM picker gtests
+        # in project mode, so it aborts parsing the whole TU with a spurious
+        # syntaxError at the first TEST_F. The files compile fine with the real
+        # toolchain; suppress the parse-abort for just these two.
+        syntaxError:*test/gtest/lgbm_picker.cpp
+        syntaxError:*test/gtest/lgbm_pcfg_picker.cpp
         # There is no ODR violation because of using separate executables,
         # but cppcheck doesn't understand that as it assumes everything
         # will be compiled together in one binary.

projects/miopen/src/CMakeLists.txt

@@ -324,6 +324,25 @@ if(MIOPEN_ENABLE_AI_KERNEL_TUNING OR MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK)
     list(APPEND MIOpen_Source anyramdb.cpp)
 endif()
 
+if(MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK)
+    list(APPEND MIOpen_Source
+        conv/heuristics/lgbm_pick.cpp
+        conv/heuristics/lgbm_metadata.cpp
+        conv/heuristics/lgbm_forest.cpp)
+    # The rank model ships as its native LightGBM text dump
+    # (kernels/lgbm_rank_model.txt, installed with the other model assets
+    # below) and is walked at runtime by lgbm_forest.cpp -- no generated C.
+
+    # Perf-config picker: one LightGBM lambdarank model PER tuneable solver.
+    # Each model ships as its native text dump (kernels/lgbm_pcfg_<solver>_
+    # model.txt, installed with the other model assets below) and is loaded +
+    # walked at runtime by lgbm_forest.cpp via the pcfg metadata loader. No
+    # generated C, so no per-solver symbol-rename / ODR machinery is needed.
+    list(APPEND MIOpen_Source
+        conv/heuristics/lgbm_pcfg_pick.cpp
+        conv/heuristics/lgbm_pcfg_metadata.cpp)
+endif()
+
 list(APPEND MIOpen_Source tmp_dir.cpp binary_cache.cpp md5.cpp)
 if(MIOPEN_ENABLE_SQLITE)
     list(APPEND MIOpen_Source sqlite_db.cpp)
@@ -948,10 +967,26 @@ if(MIOPEN_ENABLE_AI_KERNEL_TUNING OR MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK)
         add_library(nlohmann_json INTERFACE IMPORTED GLOBAL)
         target_link_libraries(nlohmann_json INTERFACE nlohmann_json::nlohmann_json)
     endif()
-    file(GLOB MODEL_FILES CONFIGURE_DEPENDS kernels/*.model)
+    file(GLOB MODEL_FILES CONFIGURE_DEPENDS
+        kernels/*.model
+        kernels/lgbm_model_meta.json
+        kernels/lgbm_pcfg_model_meta.json
+        kernels/lgbm_pcfg_catalog.json
+        kernels/lgbm_rank_model.txt
+        kernels/lgbm_pcfg_*_model.txt)
     if(NOT ENABLE_ASAN_PACKAGING )
         install(FILES ${MODEL_FILES} DESTINATION ${DATABASE_INSTALL_DIR})
     endif()
+    # Test-only golden-vector fixtures for the LGBM forest walker parity gates
+    # (test/gtest/lgbm_picker.cpp, lgbm_pcfg_picker.cpp). Copied next to the
+    # model assets so the tests find them via GetSystemDbPath(), but NOT
+    # installed into the package.
+    if(BUILD_TESTING)
+        configure_file("kernels/lgbm_rank_golden.json"
+            "${KERNELS_BINARY_DIR}/lgbm_rank_golden.json" COPYONLY)
+        configure_file("kernels/lgbm_pcfg_golden.json"
+            "${KERNELS_BINARY_DIR}/lgbm_pcfg_golden.json" COPYONLY)
+    endif()
     foreach(MODEL_FILE ${MODEL_FILES})
         get_filename_component(MODEL_FILE_FILENAME "${MODEL_FILE}" NAME)
         configure_file("${MODEL_FILE}" "${KERNELS_BINARY_DIR}/${MODEL_FILE_FILENAME}" COPYONLY)

projects/miopen/src/conv/heuristics/ai_heuristics.cpp

@@ -32,6 +32,10 @@
 //   4. Kernel tuning AI models (sequential prediction of kernel parameters)
 
 #include <miopen/conv/heuristics/ai_heuristics.hpp>
+#if MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK
+#include <miopen/conv/heuristics/lgbm_pick.hpp>
+#include <miopen/any_solver.hpp>
+#endif
 #if MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK || MIOPEN_ENABLE_AI_KERNEL_TUNING
 #include <fdeep/fdeep.hpp>
 #include <miopen/filesystem.hpp>
@@ -41,6 +45,7 @@
 #include <mutex>
 
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_AI_FDEEP_USE_SINGLE_THREAD_PREDICT)
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_ENABLE_LGBM_SELECTOR)
 
 // 3D AI heuristics - now declared properly in header
 // No need for local forward declarations since we include the header
@@ -879,22 +884,32 @@ std::vector<uint64_t> PredictSolver(const conv::ProblemDescription& problem,
     {
         auto cached_result = GetCachedPrediction(problem, device, is3d);
         if(!cached_result.empty())
+        {
+            // A cache hit short-circuits BEFORE the TunaNet/lgbm blocks, so this
+            // run shows no model log lines even though either could have produced
+            // the cached entry on an earlier call for this problem+device.
+            MIOPEN_LOG_I2("AI predict: cache hit for " << device << " (" << cached_result.size()
+                                                       << " solvers); TunaNet/lgbm not re-run");
             return cached_result;
+        }
     }
 
-    // Strategy:
-    // 1. Try ND model first (for gfx942/gfx950, supports both 2D and 3D)
-    // 2. Fall back to legacy model (for gfx908/gfx90a, 2D only)
-    // 3. Return empty vector to trigger WTI fallback for unsupported architectures
-    //
+    // Selection order: TunaNet -> LGBM -> WTI.
+    // TunaNet (the mature per-arch fdeep models) is preferred where it has a model
+    // that supports the problem. The cross-arch LGBM selector is the fallback: it
+    // covers architectures/problems TunaNet cannot serve (no per-arch model, or the
+    // problem is outside the model's supported set). Only if neither produces a
+    // prediction do we fall back to the non-AI WTI heuristic.
+
+    // 1. TunaNet.
     // Any failure inside this block (including a model that fdeep cannot load -- e.g. a
     // model exported in an incompatible format, which throws a non-miopen std::exception)
-    // is swallowed and turned into an empty result, so the caller degrades to the non-AI
-    // heuristic instead of failing the convolution. This is the predictor's contract: it
-    // returns a prediction or nothing, but never throws.
+    // is swallowed so we fall through to the LGBM selector instead of failing the
+    // convolution. This is the predictor's contract: it returns a prediction or nothing,
+    // but never throws.
     try
     {
-        // Try ND model first (preferred for gfx942/gfx950)
+        // ND model (for gfx942/gfx950, supports both 2D and 3D).
         if((is2d || is3d) && HasNDTunaNetSupport(device))
         {
             int dim                        = is3d ? 3 : 2;
@@ -907,13 +922,11 @@ std::vector<uint64_t> PredictSolver(const conv::ProblemDescription& problem,
                 return ProcessAndCachePredictions(
                     problem, device, true, predictions, model->GetSolverMap());
             }
-            // If ND model failed for this architecture, don't try legacy - go to WTI
-            MIOPEN_LOG_I2("ND TunaNet not applicable for this problem on " << device);
-            return {};
+            MIOPEN_LOG_I2("ND TunaNet not applicable for this problem on " << device
+                                                                           << "; trying LGBM");
         }
-
-        // Fall back to legacy 2D model (for gfx908/gfx90a only)
-        if(is2d && HasLegacyTunaNetSupport(device))
+        // Legacy 2D model (for gfx908/gfx90a only).
+        else if(is2d && HasLegacyTunaNetSupport(device))
         {
             std::unique_ptr<Model> model = GetModel(device);
 
@@ -924,20 +937,41 @@ std::vector<uint64_t> PredictSolver(const conv::ProblemDescription& problem,
                 return ProcessAndCachePredictions(
                     problem, device, false, predictions, model->metadata.solver_map);
             }
-            MIOPEN_LOG_I2("Legacy TunaNet not applicable for this problem on " << device);
-            return {};
+            MIOPEN_LOG_I2("Legacy TunaNet not applicable for this problem on " << device
+                                                                               << "; trying LGBM");
         }
     }
     catch(const std::exception& e)
     {
-        MIOPEN_LOG_W("TunaNet prediction failed (" << e.what()
-                                                   << "); falling back to non-AI heuristic");
-        return {};
+        MIOPEN_LOG_W("TunaNet prediction failed (" << e.what() << "); trying LGBM selector");
+    }
+
+    // 2. LGBM selector fallback: cross-arch dispatcher trained on perf-DB data.
+    // Returns the full solver vocabulary ranked by predicted speed; the caller
+    // (GetSolutionsFallback) walks this list and applies IsApplicable lazily, exactly
+    // like the TunaNet path -- so no applicability check is done here. Enabled by
+    // default (active for all architectures the model covers); set
+    // MIOPEN_ENABLE_LGBM_SELECTOR=0 to force selection straight to WTI.
+    if(!env::disabled(MIOPEN_ENABLE_LGBM_SELECTOR))
+    {
+        const auto ranked = ai::lgbm::PickSolverRanked(problem, ctx.GetStream());
+        if(!ranked.empty())
+        {
+            MIOPEN_LOG_I2("lgbm: returning " << ranked.size() << " ranked solvers");
+            std::vector<std::any> any_sol(ranked.begin(), ranked.end());
+            StorePredictionCache(problem, device, any_sol);
+            return ranked;
+        }
+        MIOPEN_LOG_I2("lgbm: abstained for " << device << ", falling back to WTI");
+    }
+    else
+    {
+        MIOPEN_LOG_I2("lgbm: disabled via MIOPEN_ENABLE_LGBM_SELECTOR=0 for " << device
+                                                                              << ", using WTI");
     }
 
-    // No TunaNet model available for this device/problem combination
-    // Return empty vector to trigger WTI fallback
-    MIOPEN_LOG_I2("No TunaNet model available for " << device << ", falling back to WTI");
+    // 3. WTI last: no AI prediction available, trigger the non-AI heuristic fallback.
+    MIOPEN_LOG_I2("No AI prediction for " << device << ", falling back to WTI");
     return {};
 }