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

projects/miopen/CMakeLists.txt

@@ -73,7 +73,7 @@ option(MIOPEN_STANDALONE_BUILD
     ${_MIOPEN_STANDALONE_DEFAULT})
 unset(_MIOPEN_STANDALONE_DEFAULT)
 
-project(MIOpen CXX)
+project(MIOpen C CXX)
 
 set(CMAKE_CXX_STANDARD 20)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)

projects/miopen/src/CMakeLists.txt

@@ -323,6 +323,32 @@ 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)
+    file(GLOB LGBM_RANK_SRCS conv/heuristics/lgbm_models/rank/*.c)
+    list(APPEND MIOpen_Source ${LGBM_RANK_SRCS})
+    # Generated C: silence machine-generated warnings, and rename every
+    # exported symbol with an lgbm_rank_ prefix. The prefixing keeps these
+    # generic Treelite names (predict, quantize, predict_unit0..7, the
+    # metadata accessors, the is_categorical global) out of MIOpen's global
+    # namespace. (v10 is rank-only; the applicability model was dropped in v4.)
+    set(LGBM_EXPORT_SYMS
+        predict postprocess quantize
+        predict_unit0 predict_unit1 predict_unit2 predict_unit3
+        predict_unit4 predict_unit5 predict_unit6 predict_unit7
+        get_num_target get_num_class get_num_feature
+        get_threshold_type get_leaf_output_type
+        is_categorical)
+    set(LGBM_RANK_DEFS "-w")
+    foreach(sym ${LGBM_EXPORT_SYMS})
+        list(APPEND LGBM_RANK_DEFS "-D${sym}=lgbm_rank_${sym}")
+    endforeach()
+    set_source_files_properties(${LGBM_RANK_SRCS} PROPERTIES
+        COMPILE_OPTIONS "${LGBM_RANK_DEFS}")
+endif()
+
 list(APPEND MIOpen_Source tmp_dir.cpp binary_cache.cpp md5.cpp)
 if(MIOPEN_ENABLE_SQLITE)
     list(APPEND MIOpen_Source sqlite_db.cpp)
@@ -951,7 +977,9 @@ 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)
     if(NOT ENABLE_ASAN_PACKAGING )
         install(FILES ${MODEL_FILES} DESTINATION ${DATABASE_INSTALL_DIR})
     endif()

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>
@@ -40,6 +44,7 @@
 #include <any>
 
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_AI_FDEEP_USE_SINGLE_THREAD_PREDICT)
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_LGBM_PICK)
 
 // 3D AI heuristics - now declared properly in header
 // No need for local forward declarations since we include the header
@@ -881,6 +886,26 @@ std::vector<uint64_t> PredictSolver(const conv::ProblemDescription& problem,
             return cached_result;
     }
 
+#if MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK
+    // LGBM-based short-circuit: cross-arch dispatcher trained on perf-DB data.
+    // Preempts the ND TunaNet model. 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.
+    if(!env::disabled(MIOPEN_DEBUG_LGBM_PICK))
+    {
+        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, falling through to TunaNet");
+    }
+#endif
+
     // 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)

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

@@ -0,0 +1,95 @@
+#include <miopen/config.h>
+#if MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK
+
+#include <miopen/conv/heuristics/lgbm_metadata.hpp>
+#include <miopen/conv/heuristics/ai_heuristics.hpp>
+#include <miopen/db_path.hpp>
+#include <miopen/errors.hpp>
+#include <miopen/logger.hpp>
+
+#include <nlohmann/json.hpp>
+
+namespace miopen {
+namespace ai {
+namespace lgbm {
+
+const LgbmMetadata& LgbmMetadata::Get()
+{
+    static const LgbmMetadata instance;
+    return instance;
+}
+
+LgbmMetadata::LgbmMetadata()
+{
+    const auto meta_path = GetSystemDbPath() / "lgbm_model_meta.json";
+
+    try
+    {
+        auto meta = ai::common::LoadJSON(meta_path);
+
+        // The rank block holds the v10 feature schema: categorical_vocab +
+        // solvers. (v10 has no triple_vocab masking, so it is not loaded.)
+        const auto& rank = meta.at("rank");
+
+        for(auto it = rank.at("categorical_vocab").begin();
+            it != rank.at("categorical_vocab").end();
+            ++it)
+        {
+            categorical_vocab[it.key()] = it.value().get<std::vector<std::string>>();
+        }
+        solvers = rank.at("solvers").get<std::vector<std::string>>();
+        for(int i = 0; i < static_cast<int>(solvers.size()); ++i)
+            solver_index[solvers[i]] = i;
+
+        // Always-applicable naive fallbacks; demoted to the tail of the ranked
+        // pick list (for low-group convs only). Optional key (older bundles
+        // omit it).
+        if(rank.contains("naive_fallback_solvers"))
+        {
+            for(const auto& name : rank.at("naive_fallback_solvers"))
+                naive_fallback.insert(name.get<std::string>());
+        }
+        // Group-count threshold for the naive demotion; defaults to 64.
+        if(rank.contains("naive_guard_max_groups"))
+            naive_guard_max_groups = rank.at("naive_guard_max_groups").get<int>();
+
+        ready = true;
+        MIOPEN_LOG_I2("LGBM metadata loaded: " << solvers.size() << " solvers");
+    }
+    catch(const std::exception& e)
+    {
+        MIOPEN_LOG_W("LGBM metadata load failed (" << e.what() << "); picker will abstain");
+        ready = false;
+    }
+}
+
+int LgbmMetadata::CategoricalCode(const std::string& column, const std::string& value) const
+{
+    const auto col_it = categorical_vocab.find(column);
+    if(col_it == categorical_vocab.end())
+        return -1;
+    const auto& vocab = col_it->second;
+    for(int i = 0; i < static_cast<int>(vocab.size()); ++i)
+    {
+        if(vocab[i] == value)
+            return i;
+    }
+    return -1;
+}
+
+int LgbmMetadata::SolverCode(const std::string& solver_name) const
+{
+    const auto it = solver_index.find(solver_name);
+    return it != solver_index.end() ? it->second : -1;
+}
+
+bool LgbmMetadata::IsNaiveFallback(const std::string& solver_name) const
+{
+    return naive_fallback.count(solver_name) != 0;
+}
+
+} // namespace lgbm
+} // namespace ai
+} // namespace miopen
+
+#endif // MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK

projects/miopen/src/conv/heuristics/lgbm_models/rank/Makefile

@@ -0,0 +1,22 @@
+lgbm_rank.so: main.o quantize.o tu0.o tu1.o tu2.o tu3.o tu4.o tu5.o tu6.o tu7.o
+	gcc -shared -O3 -o lgbm_rank.so main.o quantize.o tu0.o tu1.o tu2.o tu3.o tu4.o tu5.o tu6.o tu7.o -std=c99 
+main.o: main.c
+	gcc -c -O3 -o main.o main.c -fPIC -std=c99 
+quantize.o: quantize.c
+	gcc -c -O3 -o quantize.o quantize.c -fPIC -std=c99 
+tu0.o: tu0.c
+	gcc -c -O3 -o tu0.o tu0.c -fPIC -std=c99 
+tu1.o: tu1.c
+	gcc -c -O3 -o tu1.o tu1.c -fPIC -std=c99 
+tu2.o: tu2.c
+	gcc -c -O3 -o tu2.o tu2.c -fPIC -std=c99 
+tu3.o: tu3.c
+	gcc -c -O3 -o tu3.o tu3.c -fPIC -std=c99 
+tu4.o: tu4.c
+	gcc -c -O3 -o tu4.o tu4.c -fPIC -std=c99 
+tu5.o: tu5.c
+	gcc -c -O3 -o tu5.o tu5.c -fPIC -std=c99 
+tu6.o: tu6.c
+	gcc -c -O3 -o tu6.o tu6.c -fPIC -std=c99 
+tu7.o: tu7.c
+	gcc -c -O3 -o tu7.o tu7.c -fPIC -std=c99 

projects/miopen/src/conv/heuristics/lgbm_models/rank/header.h

@@ -0,0 +1,42 @@
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <float.h>
+#include <math.h>
+#include <stdint.h>
+
+#if defined(__clang__) || defined(__GNUC__)
+#define LIKELY(x)   __builtin_expect(!!(x), 1)
+#define UNLIKELY(x) __builtin_expect(!!(x), 0)
+#else
+#define LIKELY(x)   (x)
+#define UNLIKELY(x) (x)
+#endif
+
+#define N_TARGET 1
+#define MAX_N_CLASS 1
+
+union Entry {
+  int missing;
+  double fvalue;
+  int qvalue;
+};
+
+int32_t get_num_target(void);
+void get_num_class(int32_t* out);
+int32_t get_num_feature(void);
+const char* get_threshold_type(void);
+const char* get_leaf_output_type(void);
+void predict(union Entry* data, int pred_margin, double* result);
+void postprocess(double* result);
+int quantize(double val, unsigned fid);
+void predict_unit0(union Entry* data, double* result);
+void predict_unit1(union Entry* data, double* result);
+void predict_unit2(union Entry* data, double* result);
+void predict_unit3(union Entry* data, double* result);
+void predict_unit4(union Entry* data, double* result);
+void predict_unit5(union Entry* data, double* result);
+void predict_unit6(union Entry* data, double* result);
+void predict_unit7(union Entry* data, double* result);
+

projects/miopen/src/conv/heuristics/lgbm_models/rank/main.c

@@ -0,0 +1,56 @@
+
+#include "header.h"
+
+const unsigned char is_categorical[] = {  0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 
+  0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 
+  1, 0, 0, 0, 0, 0, 0, 1, 1, };
+static const int32_t num_class[] = {  1, };
+
+int32_t get_num_target(void) {
+  return N_TARGET;
+}
+void get_num_class(int32_t* out) {
+  for (int i = 0; i < N_TARGET; ++i) {
+    out[i] = num_class[i];
+  }
+}
+int32_t get_num_feature(void) {
+  return 61;
+}
+const char* get_threshold_type(void) {
+  return "float64";
+}
+const char* get_leaf_output_type(void) {
+  return "float64";
+}
+
+void predict(union Entry* data, int pred_margin, double* result) {
+
+  // Quantize data
+  for (int i = 0; i < 61; ++i) {
+    if (data[i].missing != -1 && !is_categorical[i]) {
+      data[i].qvalue = quantize(data[i].fvalue, i);
+    }
+  }
+
+  unsigned int tmp;
+  predict_unit0(data, result);
+  predict_unit1(data, result);
+  predict_unit2(data, result);
+  predict_unit3(data, result);
+  predict_unit4(data, result);
+  predict_unit5(data, result);
+  predict_unit6(data, result);
+  predict_unit7(data, result);
+
+  // Apply base_scores
+  result[0] += 0;
+
+  // Apply postprocessor
+  if (!pred_margin) { postprocess(result); }
+}
+
+void postprocess(double* result) {
+  // Do nothing
+}
+