Fin Perf Eval handles (reboot)

src/binary_cache.cpp

@@ -169,7 +169,7 @@ std::string LoadBinary(const TargetProperties& target,
     auto record = db.FindRecord(cfg);
     if(record)
     {
-        MIOPEN_LOG_I2("Sucessfully loaded binary for: " << verbose_name << "; args: " << args);
+        MIOPEN_LOG_I2("Successfully loaded binary for: " << verbose_name << "; args: " << args);
         return record.get();
     }
     else

src/include/miopen/any_solver.hpp

@@ -31,6 +31,8 @@
 #include <miopen/find_solution.hpp>
 #include <miopen/mlo_internal.hpp>
 
+#include <miopen/generic_search.hpp>
+
 #include <cassert>
 #include <memory>
 #include <typeinfo>
@@ -61,6 +63,11 @@ struct AnySolver
         assert(ptr_value != nullptr);
         return ptr_value->TestSysDbRecord(record);
     };
+    std::vector<ConvSolution> GetAllSolutions(const ConvolutionContext& ctx) const
+    {
+        assert(ptr_value != nullptr);
+        return ptr_value->GetAllSolutions(ctx);
+    };
     bool IsDynamic() const
     {
         assert(ptr_value != nullptr);
@@ -84,6 +91,11 @@ struct AnySolver
         assert(ptr_value != nullptr);
         return ptr_value->FindSolution(ctx, db, invoke_ctx);
     };
+    std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const
+    {
+        assert(ptr_value != nullptr);
+        return ptr_value->GetPerfCfgParams(ctx, db);
+    };
     std::string GetSolverDbId() const
     {
         assert(ptr_value != nullptr);
@@ -108,18 +120,20 @@ struct AnySolver
         using ptr = std::shared_ptr<const AnySolver_base>;
 
         virtual ~AnySolver_base(){};
-        virtual bool IsApplicable(const ConvolutionContext& ctx) const                     = 0;
-        virtual bool IsTunable() const                                                     = 0;
-        virtual bool TestSysDbRecord(const DbRecord& record) const                         = 0;
-        virtual bool IsDynamic() const                                                     = 0;
-        virtual float GetWti(const ConvolutionContext& ctx) const                          = 0;
-        virtual const std::type_info& Type() const                                         = 0;
-        virtual std::string GetSolverDbId() const                                          = 0;
+        virtual bool IsApplicable(const ConvolutionContext& ctx) const                         = 0;
+        virtual bool IsTunable() const                                                         = 0;
+        virtual bool TestSysDbRecord(const DbRecord& record) const                             = 0;
+        virtual std::vector<ConvSolution> GetAllSolutions(const ConvolutionContext& ctx) const = 0;
+        virtual bool IsDynamic() const                                                         = 0;
+        virtual float GetWti(const ConvolutionContext& ctx) const                              = 0;
+        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 size_t GetWorkspaceSize(const ConvolutionContext& ctx) const               = 0;
-        virtual bool MayNeedWorkspace() const                                              = 0;
+                                          const miopen::AnyInvokeParams& invoke_ctx) const     = 0;
+        virtual std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const      = 0;
+        virtual size_t GetWorkspaceSize(const ConvolutionContext& ctx) const                   = 0;
+        virtual bool MayNeedWorkspace() const                                                  = 0;
     };
 
     // templated derived class
@@ -140,6 +154,22 @@ struct AnySolver
             static constexpr bool Is = type::value;
         };
 
+        struct LegacySolver
+        {
+            template <typename U>
+            static constexpr auto Test(U*) ->
+                typename std::is_same<ConvSolution,
+                                      decltype(std::declval<U>().GetSolution(
+                                          std::declval<const ConvolutionContext&>(),
+                                          std::declval<const LegacyPerformanceConfig&>()))>::type;
+
+            template <typename U>
+            static constexpr std::false_type Test(...);
+
+            using type               = decltype(Test<T>(nullptr));
+            static constexpr bool Is = type::value;
+        };
+
         bool TestSysDbRecord(const DbRecord& record, std::true_type) const
         {
             using PerformanceConfig = decltype(
@@ -149,7 +179,7 @@ struct AnySolver
         }
         bool TestSysDbRecord(const DbRecord& record, std::false_type) const
         {
-            (void)(record);
+            std::ignore = record;
             return false;
         }
 
@@ -158,6 +188,43 @@ struct AnySolver
             return TestSysDbRecord(record, std::integral_constant<bool, TunableSolver::Is>());
         }
 
+        // tunable legacy solver
+        std::vector<ConvSolution>
+        GetAllSolutions(const ConvolutionContext&, std::true_type, std::true_type) const
+        {
+            MIOPEN_THROW("No solutions returned for Legacy Solvers.");
+        }
+
+        // tunable solver, not legacy
+        std::vector<ConvSolution>
+        GetAllSolutions(const ConvolutionContext& ctx, std::true_type, std::false_type) const
+        {
+            return miopen::solver::GetAllSolutions(value, ctx);
+        }
+
+        // non tunable solver
+        std::vector<ConvSolution>
+        GetAllSolutions(const ConvolutionContext& ctx, std::false_type, std::true_type) const
+        {
+            std::vector<ConvSolution> solutions;
+            solutions.push_back(value.GetSolution(ctx));
+            return solutions;
+        }
+        std::vector<ConvSolution>
+        GetAllSolutions(const ConvolutionContext& ctx, std::false_type, std::false_type) const
+        {
+            std::vector<ConvSolution> solutions;
+            solutions.push_back(value.GetSolution(ctx));
+            return solutions;
+        }
+
+        std::vector<ConvSolution> GetAllSolutions(const ConvolutionContext& ctx) const override
+        {
+            return GetAllSolutions(ctx,
+                                   std::integral_constant<bool, TunableSolver::Is>(),
+                                   std::integral_constant<bool, LegacySolver::Is>());
+        }
+
         AnySolver_tmpl(T obj) : value(std::move(obj)){};
         bool IsApplicable(const ConvolutionContext& ctx) const override
         {
@@ -166,12 +233,46 @@ struct AnySolver
         bool IsTunable() const override { return TunableSolver::Is; }
         bool IsDynamic() const override { return value.IsDynamic(); }
         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
         {
             return miopen::solver::FindSolution(value, ctx, db, invoke_ctx);
         };
+
+        std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db, std::true_type) const
+        {
+            using PerformanceConfig = decltype(value.GetDefaultPerformanceConfig(ctx));
+            PerformanceConfig config{};
+            if(db.Load(ctx, value.SolverDbId(), config))
+            {
+                MIOPEN_LOG_I2("Perf Db: Record Loaded: " << value.SolverDbId());
+                if(value.IsValidPerformanceConfig(ctx, config))
+                {
+                    std::ostringstream ss;
+                    config.Serialize(ss);
+                    return ss.str();
+                }
+                MIOPEN_LOG_I2("Perf Db: Invalid Config: " << value.SolverDbId());
+            }
+            MIOPEN_LOG_I2("Perf Db: Failed Loading: " << value.SolverDbId());
+            return "";
+        }
+        std::string
+        GetPerfCfgParams(const ConvolutionContext& ctx, const Db& db, std::false_type) const
+        {
+            MIOPEN_LOG_I2("Perf Db: No Config: " << value.SolverDbId());
+            std::ignore = ctx;
+            std::ignore = db;
+            return "";
+        }
+
+        std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db) const override
+        {
+            return GetPerfCfgParams(ctx, db, std::integral_constant<bool, TunableSolver::Is>());
+        }
+
         size_t GetWorkspaceSize(const ConvolutionContext& ctx) const override
         {
             return value.GetWorkspaceSize(ctx);

src/include/miopen/execution_context.hpp

@@ -67,6 +67,7 @@ struct ExecutionContext
 {
     // Operation modes & environment
     bool do_search               = false;
+    bool db_update               = false;
     bool save_srch_req           = false;
     bool use_asm_kernels         = false;
     bool use_hip_kernels         = true;

src/include/miopen/find_solution.hpp

@@ -63,7 +63,8 @@ auto FindSolutionImpl(
     }
     else
     {
-        if((context.do_search || enforce.IsSearch(context)) && enforce.IsDbUpdate(context))
+        if((context.do_search || enforce.IsSearch(context)) &&
+           (context.db_update || enforce.IsDbUpdate(context)))
         {
             MIOPEN_LOG_W("Perf Db: load skipped: " << s.SolverDbId() << ", enforce: " << enforce);
         }

src/include/miopen/generic_search.hpp

@@ -296,6 +296,43 @@ using RunAndMeasure_t =
                                                           std::declval<ConvSolution>(),
                                                           std::declval<float&>()));
 
+template <class Solver, class Context>
+auto GetAllConfigs(const Solver s, const Context& context)
+    -> ComputedContainer<decltype(s.GetDefaultPerformanceConfig(context)), Context>
+{
+    using PerformanceConfig = decltype(s.GetDefaultPerformanceConfig(context));
+
+    ComputedContainer<PerformanceConfig, Context> primary(context);
+    const int primary_size = std::distance(primary.begin(), primary.end());
+    ComputedContainer<PerformanceConfig, Context> spare(context, true);
+    const int spare_size = std::distance(spare.begin(), spare.end());
+    const bool useSpare  = (primary_size == 0);
+
+    ComputedContainer<PerformanceConfig, Context> all_configs = useSpare ? spare : primary;
+    const int n_runs_total = useSpare ? spare_size : primary_size;
+    MIOPEN_LOG_W(s.SolverDbId() << ": Searching the best solution among " << n_runs_total
+                                << (useSpare ? " (spare)" : "") << "...");
+
+    return all_configs;
+}
+
+template <class Solver, class Context>
+std::vector<ConvSolution> GetAllSolutions(const Solver s, const Context& context_)
+{
+    auto context                  = context_;
+    context.is_for_generic_search = true;
+
+    auto all_configs = GetAllConfigs(s, context);
+
+    std::vector<ConvSolution> solutions;
+    for(const auto& current_config : all_configs)
+    {
+        ConvSolution current_solution = s.GetSolution(context, current_config);
+        solutions.push_back(current_solution);
+    }
+    return solutions;
+}
+
 template <class Solver, class Context>
 auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParams& invoke_ctx_)
     -> decltype(s.GetDefaultPerformanceConfig(context_))
@@ -320,16 +357,8 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam
     auto& profile_h = context.GetStream();
     AutoEnableProfiling enableProfiling{profile_h};
 
-    const ComputedContainer<PerformanceConfig, Context> main(context);
-    const int main_size = std::distance(main.begin(), main.end());
-    const ComputedContainer<PerformanceConfig, Context> spare(context, true);
-    const int spare_size = std::distance(spare.begin(), spare.end());
-    const bool useSpare  = (main_size == 0);
-
-    const ComputedContainer<PerformanceConfig, Context> all_configs = useSpare ? spare : main;
-    const int n_runs_total = useSpare ? spare_size : main_size;
-    MIOPEN_LOG_W(s.SolverDbId() << ": Searching the best solution among " << n_runs_total
-                                << (useSpare ? " (spare)" : "") << "...");
+    auto all_configs       = GetAllConfigs(s, context);
+    const int n_runs_total = std::distance(all_configs.begin(), all_configs.end());
 
     bool is_passed  = false; // left false only if all iterations failed.
     float best_time = std::numeric_limits<float>::max();

src/include/miopen/solver.hpp

@@ -2436,6 +2436,12 @@ struct ConvOclBwdWrW2 : ConvTunableSolver<PerformanceConfigConvOclBwdWrw2<N_BATC
 #pragma clang diagnostic ignored "-Wweak-template-vtables"
 #endif
 
+extern template struct PerformanceConfigConvOclBwdWrw2<1>;
+extern template struct PerformanceConfigConvOclBwdWrw2<2>;
+extern template struct PerformanceConfigConvOclBwdWrw2<4>;
+extern template struct PerformanceConfigConvOclBwdWrw2<8>;
+extern template struct PerformanceConfigConvOclBwdWrw2<16>;
+
 extern template struct ConvOclBwdWrW2<1>;
 extern template struct ConvOclBwdWrW2<2>;
 extern template struct ConvOclBwdWrW2<4>;
@@ -2463,6 +2469,7 @@ struct ConvOclBwdWrW2NonTunable final : ConvOclBwdWrW2<1>
 private:
     // This function dervied from ConvOclBwdWrW2 is declared private
     // so that this solver is not marked searchable/tunable.
+    using ConvOclBwdWrW2<1>::GetDefaultPerformanceConfig;
     using ConvOclBwdWrW2<1>::GetSolution;
 };
 

src/solver/conv_asm_1x1u.cpp

@@ -225,8 +225,7 @@ PerformanceConfigConvAsm1x1U::PerformanceConfigConvAsm1x1U(int read_size_,
 {
 }
 
-inline bool
-PerformanceConfigConvAsm1x1U::operator==(const PerformanceConfigConvAsm1x1U& other) const
+bool PerformanceConfigConvAsm1x1U::operator==(const PerformanceConfigConvAsm1x1U& other) const
 {
     // clang-format off
     return read_size == other.read_size

src/solver/conv_asm_1x1u_bias_activ.cpp

@@ -53,7 +53,7 @@ bool PerformanceConfigConvBiasActivAsm1x1U::IsValid(const ConvolutionContext& co
     return PerformanceConfigConvAsm1x1U::IsValid(config);
 }
 
-inline bool PerformanceConfigConvBiasActivAsm1x1U::operator==(
+bool PerformanceConfigConvBiasActivAsm1x1U::operator==(
     const PerformanceConfigConvBiasActivAsm1x1U& other) const
 {
     // clang-format off

src/solver/conv_asm_1x1u_stride2.cpp

@@ -287,8 +287,7 @@ PerformanceConfigConvAsm1x1UV2::PerformanceConfigConvAsm1x1UV2(int chunk_size_,
 {
 }
 
-inline bool
-PerformanceConfigConvAsm1x1UV2::operator==(const PerformanceConfigConvAsm1x1UV2& other) const
+bool PerformanceConfigConvAsm1x1UV2::operator==(const PerformanceConfigConvAsm1x1UV2& other) const
 {
     // clang-format off
     return chunk_size == other.chunk_size

src/solver/conv_asm_3x3u.cpp

@@ -70,8 +70,7 @@ PerformanceConfigConvAsm3x3U::PerformanceConfigConvAsm3x3U(int lwc, int fpw, int
 {
 }
 
-inline bool
-PerformanceConfigConvAsm3x3U::operator==(const PerformanceConfigConvAsm3x3U& other) const
+bool PerformanceConfigConvAsm3x3U::operator==(const PerformanceConfigConvAsm3x3U& other) const
 {
     return PerfFieldRules().Compare(*this, other);
 }

src/solver/conv_asm_dir_BwdWrW1x1.cpp

@@ -274,8 +274,8 @@ PerformanceConfigConvAsmBwdWrW1x1::PerformanceConfigConvAsmBwdWrW1x1(int chunk_s
 {
 }
 
-inline bool
-PerformanceConfigConvAsmBwdWrW1x1::operator==(const PerformanceConfigConvAsmBwdWrW1x1& other) const
+bool PerformanceConfigConvAsmBwdWrW1x1::operator==(
+    const PerformanceConfigConvAsmBwdWrW1x1& other) const
 {
     // clang-format off
     return chunk_size == other.chunk_size

src/solver/conv_asm_dir_BwdWrW3x3.cpp

@@ -111,8 +111,8 @@ PerformanceConfigAsmDirect3x3WrW::PerformanceConfigAsmDirect3x3WrW(
 {
 }
 
-inline bool
-PerformanceConfigAsmDirect3x3WrW::operator==(const PerformanceConfigAsmDirect3x3WrW& other) const
+bool PerformanceConfigAsmDirect3x3WrW::operator==(
+    const PerformanceConfigAsmDirect3x3WrW& other) const
 {
     // clang-format off
     return limit_wave_cnt == other.limit_wave_cnt

src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp

@@ -217,6 +217,8 @@ PerformanceImplicitGemmBwdDataV4R1::CalculateGemmABlockCopyPerformanceParameters
 
         // decide threadwise copy lengths
         const auto a_data_per_thread_copy_gemmm = SrcDataPerRead_GemmM;
+        if(a_data_per_thread_copy_gemmm == 0)
+            MIOPEN_THROW("DIV/0 with a_data_per_thread_copy_gemmm");
         const auto a_data_per_thread_copy_gemmk =
             a_data_per_thread_copy / a_data_per_thread_copy_gemmm;
 

src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp

@@ -315,6 +315,8 @@ PerformanceImplicitGemmV4R4Fwd::CalculateGemmBBlockCopyPerformanceParameters(
         // GemmBBlockCopyDstDataPerWrite_GemmN also bounded by size of threadwise copy
         DstDataPerWrite_GemmN = gcd(DstDataPerWrite_GemmN, b_data_per_thread_copy_gemmn);
 
+        if(b_data_per_thread_copy_gemmk == 0)
+            MIOPEN_THROW("DIV/0 with b_data_per_thread_copy_gemmk");
         // calculate blockwise copy thread cluster lengths
         ClusterLengths_GemmK = GemmKPerBlock / b_data_per_thread_copy_gemmk;
         ClusterLengths_GemmN = GemmNPerBlock / b_data_per_thread_copy_gemmn;

src/solver/conv_ocl_dir2D_bwdWrW_2.cpp

@@ -156,7 +156,7 @@ ConvSolution ConvOclBwdWrW2NonTunable::GetSolution(const ConvolutionContext& par
 }
 
 template <int N_BATCH_LOOPS>
-inline bool PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>::operator==(
+bool PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>::operator==(
     const PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>& other) const
 {
     // clang-format off
@@ -745,6 +745,12 @@ ConvOclBwdWrW2<N_BATCH_LOOPS>::Search(const ConvolutionContext& context,
 /// We need to instantiate required classes implicitly.
 /// The reason is that we do not define the whole template class
 /// in the header, only declaring it there.
+template struct PerformanceConfigConvOclBwdWrw2<1>;
+template struct PerformanceConfigConvOclBwdWrw2<2>;
+template struct PerformanceConfigConvOclBwdWrw2<4>;
+template struct PerformanceConfigConvOclBwdWrw2<8>;
+template struct PerformanceConfigConvOclBwdWrw2<16>;
+
 template struct ConvOclBwdWrW2<1>;
 template struct ConvOclBwdWrW2<2>;
 template struct ConvOclBwdWrW2<4>;