[Find-2.0] Serialize kernel perf configs (#1876)

ROCm · Dec 6, 2022 · d437a79 · d437a79
1 parent 2d87f83
commit d437a79
Show file tree

Hide file tree

Showing 7 changed files with 80 additions and 35 deletions.
diff --git a/fin/CMakeLists.txt b/fin/CMakeLists.txt
@@ -63,8 +63,8 @@ if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
 endif()
 
 ############################################################
-# require C++14
-add_compile_options(-std=c++14)
+# require C++17
+add_compile_options(-std=c++17)
 
 ############################################################
 # OPTION - MIOpen Backend

diff --git a/fin/src/include/conv_fin.hpp b/fin/src/include/conv_fin.hpp
@@ -691,6 +691,7 @@ int ConvFin<Tgpu, Tref>::MIOpenPerfEval()
                     std::cerr << solver_name << " Finished Search WRW" << std::endl;
                     kern_objs = BuildJsonKernelList(h, solution.construction_params);
                     SolutionHasProgram(h, solution);
+                    // todo: proper handling of empty values
                     params = s.GetPerfCfgParams(ctx, db);
 
                     const auto invoker =

diff --git a/src/include/miopen/any_solver.hpp b/src/include/miopen/any_solver.hpp
@@ -42,8 +42,6 @@ namespace solver {
 
 struct AnySolver
 {
-    using Db = decltype(std::declval<mlo_construct_base>().GetDb());
-
     AnySolver() : ptr_value(nullptr){};
     template <class U>
     AnySolver(U src) : ptr_value(new AnySolver_tmpl<U>(std::forward<U>(src))){};
@@ -85,13 +83,14 @@ struct AnySolver
     };
     bool IsEmpty() const { return ptr_value == nullptr; };
     ConvSolution FindSolution(const ConvolutionContext& ctx,
-                              Db& db,
-                              const miopen::AnyInvokeParams& invoke_ctx) const
+                              PerformanceDb& db,
+                              const miopen::AnyInvokeParams& invoke_ctx,
+                              const std::string& perf_cfg = "") const
     {
         assert(ptr_value != nullptr);
-        return ptr_value->FindSolution(ctx, db, invoke_ctx);
+        return ptr_value->FindSolution(ctx, db, invoke_ctx, perf_cfg);
     };
-    std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const
+    std::string GetPerfCfgParams(const ConvolutionContext& ctx, PerformanceDb& db) const
     {
         assert(ptr_value != nullptr);
         return ptr_value->GetPerfCfgParams(ctx, db);
@@ -130,9 +129,11 @@ struct AnySolver
         virtual const std::type_info& Type() const                                             = 0;
         virtual std::string GetSolverDbId() const                                              = 0;
         virtual ConvSolution FindSolution(const ConvolutionContext& ctx,
-                                          Db& db,
-                                          const miopen::AnyInvokeParams& invoke_ctx) const     = 0;
-        virtual std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const      = 0;
+                                          PerformanceDb& db,
+                                          const miopen::AnyInvokeParams& invoke_ctx,
+                                          const std::string& perf_cfg) const                   = 0;
+        virtual std::string GetPerfCfgParams(const ConvolutionContext& ctx,
+                                             PerformanceDb& db) const                          = 0;
         virtual size_t GetWorkspaceSize(const ConvolutionContext& ctx) const                   = 0;
         virtual bool MayNeedWorkspace() const                                                  = 0;
     };
@@ -244,51 +245,56 @@ struct AnySolver
         float GetWti(const ConvolutionContext& ctx) const override { return value.GetWti(ctx); }
 
         ConvSolution FindSolution(const ConvolutionContext& ctx,
-                                  Db& db,
-                                  const miopen::AnyInvokeParams& invoke_ctx) const override
+                                  PerformanceDb& db,
+                                  const miopen::AnyInvokeParams& invoke_ctx,
+                                  const std::string& perf_cfg) const override
         {
-            return miopen::solver::FindSolution(value, ctx, db, invoke_ctx);
+            return miopen::solver::FindSolution(value, ctx, db, invoke_ctx, perf_cfg);
         };
 
-        std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db, std::true_type) const
+        std::string
+        GetPerfCfgParams(const ConvolutionContext& ctx, PerformanceDb& db, std::true_type) const
         {
             using PerformanceConfig = decltype(value.GetDefaultPerformanceConfig(ctx));
             PerformanceConfig config{};
             if(db.Load(ctx.problem, value.SolverDbId(), config))
             {
-                MIOPEN_LOG_I2("Perf Db: Record Loaded: " << value.SolverDbId());
+                MIOPEN_LOG_I2("PerformanceDb: Record Loaded: " << value.SolverDbId());
                 if(value.IsValidPerformanceConfig(ctx, config))
                 {
                     return config.ToString();
                 }
-                MIOPEN_LOG_I2("Perf Db: Invalid Config: " << value.SolverDbId());
+                MIOPEN_LOG_I2("PerformanceDb: Invalid Config: " << value.SolverDbId());
             }
             else if(!value.AltSolverDbId().empty() &&
                     db.Load(ctx.problem, value.AltSolverDbId(), config))
             {
-                MIOPEN_LOG_I("Perf Db: alternate record loaded: " << value.AltSolverDbId());
+                MIOPEN_LOG_I("PerformanceDb: alternate record loaded: " << value.AltSolverDbId());
                 if(value.IsValidPerformanceConfig(ctx, config))
                 {
                     return config.ToString();
                 }
-                MIOPEN_LOG_I2("Perf Db: Invalid alternate record from Perf Db: "
-                              << value.AltSolverDbId() << ": " << config);
+                MIOPEN_LOG_I2("PerformanceDb: Invalid alternate record: " << value.AltSolverDbId()
+                                                                          << ": " << config);
             }
 
-            MIOPEN_LOG_I2("Perf Db: Failed Loading, Using Default: " << value.SolverDbId());
+            MIOPEN_LOG_I2("PerformanceDb: Failed Loading, Using Default: " << value.SolverDbId());
             config = value.GetDefaultPerformanceConfig(ctx);
             return config.ToString();
         }
-        std::string
-        GetPerfCfgParams(const ConvolutionContext& ctx, const Db& db, std::false_type) const
+
+        std::string GetPerfCfgParams(const ConvolutionContext& ctx,
+                                     const PerformanceDb& db,
+                                     std::false_type) const
         {
-            MIOPEN_LOG_I2("Perf Db: No Config: " << value.SolverDbId());
+            MIOPEN_LOG_I2("PerformanceDb: No Config: " << value.SolverDbId());
             std::ignore = ctx;
             std::ignore = db;
             return "";
         }
 
-        std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const override
+        std::string GetPerfCfgParams(const ConvolutionContext& ctx,
+                                     PerformanceDb& db) const override
         {
             return GetPerfCfgParams(ctx, db, std::integral_constant<bool, TunableSolver::Is>());
         }

diff --git a/src/include/miopen/find_solution.hpp b/src/include/miopen/find_solution.hpp
@@ -45,8 +45,12 @@ struct AnyInvokeParams;
 namespace solver {
 
 template <class Solver, class Context, class Db>
-auto FindSolutionImpl(
-    rank<1>, Solver s, const Context& context, Db& db, const AnyInvokeParams& invoke_ctx)
+auto FindSolutionImpl(rank<1>,
+                      Solver s,
+                      const Context& context,
+                      Db& db,
+                      const AnyInvokeParams& invoke_ctx,
+                      const std::string& perf_cfg)
     -> decltype(s.GetSolution(context, s.Search(context, invoke_ctx)))
 {
     const FindEnforce enforce;
@@ -94,6 +98,16 @@ auto FindSolutionImpl(
                               << s.AltSolverDbId() << ": " << config
                               << ". Performance may degrade.");
             }
+            else if(!perf_cfg.empty())
+            {
+                config.Deserialize(perf_cfg);
+                if(s.IsValidPerformanceConfig(context, config))
+                {
+                    return s.GetSolution(context, config);
+                }
+                MIOPEN_LOG_WE("Invalid config loaded from Perf Db: "
+                              << s.SolverDbId() << ": " << config << ". Performance may degrade.");
+            }
             else
             {
                 MIOPEN_LOG_I("Perf Db: record not found for: " << s.SolverDbId());
@@ -120,7 +134,8 @@ auto FindSolutionImpl(
 }
 
 template <class Solver, class Context, class Db>
-auto FindSolutionImpl(rank<0>, Solver s, const Context& context, Db&, const AnyInvokeParams&)
+auto FindSolutionImpl(
+    rank<0>, Solver s, const Context& context, Db&, const AnyInvokeParams&, const std::string&)
     -> decltype(s.GetSolution(context))
 {
     MIOPEN_LOG_I(s.SolverDbId() << " (not searchable)");
@@ -134,13 +149,16 @@ auto FindSolutionImpl(rank<0>, Solver s, const Context& context, Db&, const AnyI
 /// Could take long if an exhaustive search is requested/performed.
 /// May read/write perfDb.
 template <class Solver, class Context, class Db>
-ConvSolution
-FindSolution(Solver s, const Context& context, Db& db, const AnyInvokeParams& invoke_ctx)
+ConvSolution FindSolution(Solver s,
+                          const Context& context,
+                          Db& db,
+                          const AnyInvokeParams& invoke_ctx,
+                          const std::string& perf_cfg = "")
 {
     static_assert(sizeof(Solver) == sizeof(SolverBase), "Solver must be stateless");
     static_assert(std::is_base_of<SolverBase, Solver>{}, "Not derived class of SolverBase");
     // TODO: This assumes all solutions are ConvSolution
-    auto solution      = FindSolutionImpl(rank<1>{}, s, context, db, invoke_ctx);
+    auto solution      = FindSolutionImpl(rank<1>{}, s, context, db, invoke_ctx, perf_cfg);
     solution.solver_id = s.SolverDbId();
     return solution;
 }

diff --git a/src/include/miopen/solution.hpp b/src/include/miopen/solution.hpp
@@ -38,6 +38,7 @@
 
 #include <boost/optional.hpp>
 
+#include <optional>
 #include <unordered_map>
 
 namespace miopen {
@@ -81,6 +82,7 @@ struct Solution : miopenSolution
     void SetWorkspaceSize(std::size_t value) { workspace_required = value; }
     const solver::Id& GetSolver() const { return solver; }
     void SetSolver(solver::Id value) { solver = value; }
+    void SetPerfConfig(const std::optional<std::string>& cfg) { perf_cfg = cfg; }
     const Problem& GetProblem() const { return problem; }
     void SetProblem(Problem value) { problem = std::move(value); }
 
@@ -97,6 +99,7 @@ struct Solution : miopenSolution
     std::size_t workspace_required = 0;
     solver::Id solver;
     Problem problem;
+    std::optional<std::string> perf_cfg = std::nullopt;
 
     void RunImpl(Handle& handle,
                  const std::unordered_map<miopenTensorArgumentId_t, RunInput>& inputs,

diff --git a/src/problem.cpp b/src/problem.cpp
@@ -32,6 +32,8 @@
 #include <miopen/datatype.hpp>
 #include <miopen/execution_context.hpp>
 #include <miopen/handle.hpp>
+#include <miopen/any_solver.hpp>
+#include <miopen/mlo_internal.hpp>
 #include <miopen/solution.hpp>
 #include <miopen/search_options.hpp>
 #include <miopen/tensor_ops.hpp>
@@ -340,7 +342,13 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
                                                : conv::Direction::Forward;
     })();
 
-    const auto netcfg = actual.AsConvolution().BuildConfKey();
+    const auto conv_problem = actual.AsConvolution();
+    const auto netcfg       = conv_problem.BuildConfKey();
+    auto conv_ctx           = ConvolutionContext{conv_problem, {&handle}};
+    conv_ctx.DetectRocm();
+    conv_ctx.SetupFloats();
+
+    decltype(auto) db = GetDb(conv_ctx);
 
     for(auto i = 0; i < found; ++i)
     {
@@ -351,8 +359,8 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
         solution.SetTime(find1_solutions[i].time);
         solution.SetWorkspaceSize(find1_solutions[i].memory);
         solution.SetSolver(handle.GetFound1_0SolverId(netcfg, AlgorithmName{algo}).value());
+        solution.SetPerfConfig(solution.GetSolver().GetSolver().GetPerfCfgParams(conv_ctx, db));
         solution.SetProblem(*this);
-
         MIOPEN_LOG_I("Found solvution: " << solution.GetSolver().ToString() << " , "
                                          << solution.GetWorkspaceSize() << ", "
                                          << solution.GetTime());

diff --git a/src/solution.cpp b/src/solution.cpp
@@ -165,8 +165,9 @@ void Solution::RunImpl(Handle& handle,
     conv_ctx.DetectRocm();
     conv_ctx.SetupFloats();
 
-    decltype(auto) db        = GetDb(conv_ctx);
-    const auto conv_solution = GetSolver().GetSolver().FindSolution(conv_ctx, db, invoke_ctx);
+    decltype(auto) db = GetDb(conv_ctx);
+    const auto conv_solution =
+        GetSolver().GetSolver().FindSolution(conv_ctx, db, invoke_ctx, perf_cfg.value_or(""));
     decltype(auto) invoker =
         handle.PrepareInvoker(*conv_solution.invoker_factory, conv_solution.construction_params);
     handle.RegisterInvoker(invoker, net_cfg, GetSolver().ToString());
@@ -212,6 +213,9 @@ void to_json(nlohmann::json& json, const Solution& solution)
         {"solver", solution.solver.ToString()},
         {"problem", solution.problem},
     };
+
+    if(solution.perf_cfg.has_value())
+        json["perf_cfg"] = *solution.perf_cfg;
 }
 
 void from_json(const nlohmann::json& json, Solution& solution)
@@ -233,5 +237,10 @@ void from_json(const nlohmann::json& json, Solution& solution)
     json.at("workspace").get_to(solution.workspace_required);
     solution.solver = json.at("solver").get<std::string>();
     json.at("problem").get_to(solution.problem);
+
+    const auto perf_cfg_json = json.find("perf_cfg");
+    solution.perf_cfg        = perf_cfg_json != json.end()
+                                   ? std::optional{perf_cfg_json->get<std::string>()}
+                                   : std::nullopt;
 }
 } // namespace miopen