From 6115f17b4d99753911a67124b9520164fc1992d1 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 4 Sep 2025 14:42:09 -0500 Subject: [PATCH 01/26] add cutoff time for GenericSearch relative to worst case timing from other solvers --- projects/miopen/include/miopen/miopen.h | 8 ++-- .../miopen/src/include/miopen/any_solver.hpp | 15 +++---- .../src/include/miopen/execution_context.hpp | 8 ++-- .../src/include/miopen/find_controls.hpp | 7 ++++ .../src/include/miopen/find_solution.hpp | 3 ++ .../src/include/miopen/generic_search.hpp | 41 ++++++++++++------- 6 files changed, 55 insertions(+), 27 deletions(-) diff --git a/projects/miopen/include/miopen/miopen.h b/projects/miopen/include/miopen/miopen.h index 676c1bb086b..bd741fdd3cb 100644 --- a/projects/miopen/include/miopen/miopen.h +++ b/projects/miopen/include/miopen/miopen.h @@ -8523,7 +8523,8 @@ MIOPEN_EXPORT miopenStatus_t miopenMultiMarginLossForward(miopenHandle_t handle, * 2. DbUpdate: Tune and update the database. * 3. Search: Search the database first; if no record is found, tune but do not update the database. * 4. SearchDbUpdate: Combination of Search and DbUpdate. - * 5. DbClean: Remove existing entry, do not tune. + * 5. SearchCutoffDbUpdate: Search with tuning perf cutoff and DbUpdate. + * 6. DbClean: Remove existing entry, do not tune. * Note: MIOpenFindEnforce has additional features that are not supported by TuningPolicy. * Note: TuningPolicy has higher priority over MIOPEN_FIND_ENFORCE. */ @@ -8533,8 +8534,9 @@ typedef enum miopenTuningPolicyDbUpdate = 2, /* tune and update the db */ miopenTuningPolicySearch = 3, /* search db first, if record not found tune but do not update the db */ - miopenTuningPolicySearchDbUpdate = 4, /* combination of Search and DbUpdate */ - miopenTuningPolicyDbClean = 5, /* remove existing entry, do not tune */ + miopenTuningPolicySearchDbUpdate = 4, /* combination of Search and DbUpdate */ + miopenTuningPolicySearchCutoffDbUpdate = 5, + miopenTuningPolicyDbClean = 6, /* remove existing entry, do not tune */ } miopenTuningPolicy_t; /*! @ingroup handle diff --git a/projects/miopen/src/include/miopen/any_solver.hpp b/projects/miopen/src/include/miopen/any_solver.hpp index 18cdeba3bda..5de823cc133 100644 --- a/projects/miopen/src/include/miopen/any_solver.hpp +++ b/projects/miopen/src/include/miopen/any_solver.hpp @@ -132,10 +132,10 @@ struct AnySolver std::string GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, - std::vector* perf_sols = nullptr) const + std::vector* perf_solsp = nullptr) const { assert(ptr_value != nullptr); - return ptr_value->GenericSearch(ctx, problem, invoke_ctx, perf_sols); + return ptr_value->GenericSearch(ctx, problem, invoke_ctx, perf_solsp); } InvokerFactory GetInvokeFactory(const ExecutionContext& ctx, @@ -213,7 +213,7 @@ struct AnySolver GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, - std::vector* perf_sols) const = 0; + std::vector* perf_solsp) const = 0; virtual InvokerFactory GetInvokeFactory(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const std::string& perf_cfg) const = 0; @@ -457,11 +457,12 @@ struct AnySolver std::string GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, - std::vector* perf_sols, + std::vector* perf_solsp, std::true_type, std::false_type) const { - auto config = miopen::solver::GenericSearch(value, ctx, problem, invoke_ctx, perf_sols); + auto config = + miopen::solver::GenericSearch(value, ctx, problem, invoke_ctx, perf_solsp); return config.ToString(); } @@ -490,12 +491,12 @@ struct AnySolver GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, - std::vector* perf_sols) const override + std::vector* perf_solsp) const override { return GenericSearch(ctx, problem, invoke_ctx, - perf_sols, + perf_solsp, std::integral_constant(), std::integral_constant()); } diff --git a/projects/miopen/src/include/miopen/execution_context.hpp b/projects/miopen/src/include/miopen/execution_context.hpp index 81576c363e2..80c6d0829a6 100644 --- a/projects/miopen/src/include/miopen/execution_context.hpp +++ b/projects/miopen/src/include/miopen/execution_context.hpp @@ -84,6 +84,7 @@ struct MIOPEN_INTERNALS_EXPORT ExecutionContext // Operation modes & environment bool do_search = false; + bool search_cutoff = false; bool db_update = false; bool use_asm_kernels = false; bool use_hip_kernels = true; @@ -94,9 +95,10 @@ struct MIOPEN_INTERNALS_EXPORT ExecutionContext // to optimize the getWorkspaceSize() calls for speed. This specific optimization is correct // because Solvers shall be written so that the required workspace size does not depend on the // performance config. - bool disable_perfdb_access = false; - bool use_dynamic_solutions_only = false; - bool is_for_generic_search = false; + bool disable_perfdb_access = false; + bool use_dynamic_solutions_only = false; + bool is_for_generic_search = false; + float generic_search_cutoff_time = std::numeric_limits::max(); inline const Handle& GetStream() const { return *stream; } inline void SetStream(const Handle* stream_) { stream = stream_; } diff --git a/projects/miopen/src/include/miopen/find_controls.hpp b/projects/miopen/src/include/miopen/find_controls.hpp index 64dc340a10c..4f27bb60f66 100644 --- a/projects/miopen/src/include/miopen/find_controls.hpp +++ b/projects/miopen/src/include/miopen/find_controls.hpp @@ -54,6 +54,7 @@ enum class FindEnforceAction DbUpdate, Search, SearchDbUpdate, + SearchCutoffDbUpdate, DbClean, Last_ = DbClean, Default_ = None, @@ -90,6 +91,12 @@ class MIOPEN_INTERNALS_EXPORT FindEnforce action == FindEnforceAction::DbUpdate); } + template + bool IsSearchCutoff(const Context& context) const + { + return IsEnabled(context) && (action == FindEnforceAction::SearchCutoffDbUpdate); + } + template bool IsDbUpdate(const Context& context) const { diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index f964ea1dd54..4fa58687527 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -137,6 +137,9 @@ auto FindSolutionImpl(rank<1>, if(context.do_search || enforce.IsSearch(context)) // TODO: Make it a customization point { + if(enforce.IsSearchCutoff(context)) + context.search_cutoff = true; + MIOPEN_LOG_I("Starting search: " << s.SolverDbId() << ", enforce: " << enforce); try { diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 35426fd35fc..0fec650fdd0 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -426,7 +426,7 @@ auto GenericSearch(const Solver s, const Context& context_, const Problem& problem, const AnyInvokeParams& invoke_ctx_, - std::vector* perf_sols = nullptr) + std::vector* perf_solsp = nullptr) -> decltype(s.GetDefaultPerformanceConfig(context_, problem)) { auto context = context_; @@ -443,10 +443,7 @@ auto GenericSearch(const Solver s, }(); // list of sampled solutions - if(perf_sols) - { - perf_sols->erase(perf_sols->begin(), perf_sols->end()); - } + std::vector perf_sols; auto& profile_h = context.GetStream(); const AutoEnableProfiling enableProfiling{profile_h}; @@ -595,6 +592,16 @@ auto GenericSearch(const Solver s, if(ret == 0) { + // If 1st probe of 1st successful config is worse than the cutoff time abort the + // search + if(perf_sols.empty() && elapsed_time > context.generic_search_cutoff_time) + { + MIOPEN_LOG_E("Measured time: " << elapsed_time << " was greater than cutoff: " + << context.generic_search_cutoff_time + << " aborting search..."); + return current_config; + } + // Smooth the jitter of measurements: // If the 1st probe is NOT too bad (measured time <= 1.10 * worst sample of the best // config), then gather 9 more samples, and remove positive z-score outliers. Use @@ -646,13 +653,9 @@ auto GenericSearch(const Solver s, MIOPEN_LOG_I2("Mean is not better: " << elapsed_time << " >= " << best_time); } - } } - if(perf_sols) - { - perf_sols->push_back({current_config.ToString(), elapsed_time}); - } + perf_sols.push_back({current_config.ToString(), elapsed_time}); } // Banchmarked kernels will not be used anymore. @@ -692,10 +695,20 @@ auto GenericSearch(const Solver s, if(!is_passed) MIOPEN_THROW("Search failed"); - if(perf_sols) - std::sort(perf_sols->begin(), perf_sols->end(), [](SolutionPerf a, SolutionPerf b) { - return a.time < b.time; - }); + std::sort(perf_sols.begin(), perf_sols.end(), [](SolutionPerf a, SolutionPerf b) { + return a.time < b.time; + }); + + // if using cutoff time for search and new cutoff is shorter, update + if(context.search_cutoff) + { + float new_cutoff = perf_sols.end().time * 2; + if(new_cutoff < context.generic_search_cutoff_time) + context_.generic_search_cutoff_time = new_cutoff; + } + + if(perf_solsp) + *perf_solsp = std::move(perf_sols); // Run once with the default config and show score. const auto& invoker = profile_h.PrepareInvoker(*default_solution.invoker_factory, From 02a92b938d230196d2931d6d3ffe75ada4a61e5e Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 5 Sep 2025 12:29:26 -0500 Subject: [PATCH 02/26] remove const from execution context to allow message passing between solvers in generic search --- projects/miopen/src/conv/solver_finders.cpp | 12 +-- .../miopen/src/include/miopen/any_solver.hpp | 14 ++-- .../src/include/miopen/batchnorm/solvers.hpp | 4 +- .../include/miopen/conv/solver_finders.hpp | 8 +- .../src/include/miopen/conv/solvers.hpp | 82 +++++++++---------- .../src/include/miopen/find_solution.hpp | 2 +- .../src/include/miopen/fusion/solvers.hpp | 12 +-- projects/miopen/src/include/miopen/solver.hpp | 2 +- projects/miopen/src/mlo_dir_conv.cpp | 16 ++-- projects/miopen/src/solver.cpp | 2 +- .../src/solver/batchnorm/backward_spatial.cpp | 2 +- .../src/solver/batchnorm/forward_spatial.cpp | 2 +- .../miopen/src/solver/conv/conv_asm_1x1u.cpp | 2 +- .../src/solver/conv/conv_asm_1x1u_stride2.cpp | 2 +- .../miopen/src/solver/conv/conv_asm_3x3u.cpp | 2 +- .../solver/conv/conv_asm_dir_BwdWrW1x1.cpp | 2 +- .../solver/conv/conv_asm_dir_BwdWrW3x3.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 2 +- .../conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_bwd_xdlops.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_fwd_xdlops.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_wrw_xdlops.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_data_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_bwd_v1r1.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_bwd_v4r1.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_fwd_v4r1.cpp | 4 +- .../conv/conv_hip_implicit_gemm_fwd_v4r4.cpp | 2 +- ...conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp | 2 +- ...licit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp | 2 +- ...conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp | 2 +- .../conv_hip_implicit_gemm_fwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_bwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_fwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_wrw_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_wrw_v4r4.cpp | 2 +- ...conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp | 2 +- ...licit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_bwd.cpp | 2 +- .../conv/conv_mlir_igemm_bwd_xdlops.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_fwd.cpp | 2 +- .../conv/conv_mlir_igemm_fwd_xdlops.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_wrw.cpp | 2 +- .../conv/conv_mlir_igemm_wrw_xdlops.cpp | 2 +- .../solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp | 2 +- .../conv_ocl_dir2Dfwd_exhaustive_search.cpp | 2 +- .../miopen/src/solver/conv/conv_winoRxS.cpp | 2 +- .../solver/conv_asm_1x1u_bias_activ_fused.cpp | 2 +- .../conv_ck_igemm_fwd_bias_activ_fused.cpp | 2 +- ..._ck_igemm_fwd_bias_res_add_activ_fused.cpp | 2 +- .../conv_ck_igemm_grp_fwd_activ_fused.cpp | 2 +- ...conv_ck_igemm_grp_fwd_bias_activ_fused.cpp | 2 +- .../src/solver/conv_ocl_dir2Dfwd_fused.cpp | 2 +- 57 files changed, 125 insertions(+), 125 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 1b7a39e0e98..2638d09d48f 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -66,7 +66,7 @@ class DirectSolverFinder : public SolversFinderMixin FindImpl(const ExecutionContext& ctx, + std::vector FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -95,7 +95,7 @@ class ImplicitGemmSolverFinder : public SolversFinderMixin FindImpl(const ExecutionContext& ctx, + std::vector FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -126,7 +126,7 @@ class FftSolverFinder : public SolversFinderMixin FindImpl(const ExecutionContext& ctx, + std::vector FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -153,7 +153,7 @@ class GemmSolverFinder : public SolversFinderMixin FindImpl(const ExecutionContext& ctx, + std::vector FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -180,7 +180,7 @@ class WinogradSolverFinder : public SolversFinderMixin FindImpl(const ExecutionContext& ctx, + std::vector FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters& parameters, @@ -334,7 +334,7 @@ std::vector EvaluateInvokers(const Handle& handle, } FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, - const ExecutionContext& ctx, + ExecutionContext& ctx, const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters, const std::vector>& finders, diff --git a/projects/miopen/src/include/miopen/any_solver.hpp b/projects/miopen/src/include/miopen/any_solver.hpp index 5de823cc133..7e610719dae 100644 --- a/projects/miopen/src/include/miopen/any_solver.hpp +++ b/projects/miopen/src/include/miopen/any_solver.hpp @@ -129,7 +129,7 @@ struct AnySolver return ptr_value->FindSolution(ctx, problem, db_getter, invoke_ctx, perf_cfg); } - std::string GenericSearch(const ExecutionContext& ctx, + std::string GenericSearch(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp = nullptr) const @@ -210,7 +210,7 @@ struct AnySolver const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; virtual std::string - GenericSearch(const ExecutionContext& ctx, + GenericSearch(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp) const = 0; @@ -443,7 +443,7 @@ struct AnySolver } // tunable legacy solver - std::string GenericSearch(const ExecutionContext&, + std::string GenericSearch(ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -454,7 +454,7 @@ struct AnySolver } // tunable solver - std::string GenericSearch(const ExecutionContext& ctx, + std::string GenericSearch(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp, @@ -467,7 +467,7 @@ struct AnySolver } // non-tunable solver has no search - std::string GenericSearch(const ExecutionContext&, + std::string GenericSearch(ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -477,7 +477,7 @@ struct AnySolver MIOPEN_THROW("No GenericSearch for non-tunable Solvers."); } - std::string GenericSearch(const ExecutionContext&, + std::string GenericSearch(ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -488,7 +488,7 @@ struct AnySolver } std::string - GenericSearch(const ExecutionContext& ctx, + GenericSearch(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp) const override diff --git a/projects/miopen/src/include/miopen/batchnorm/solvers.hpp b/projects/miopen/src/include/miopen/batchnorm/solvers.hpp index 07c358efba3..e1d6b065429 100644 --- a/projects/miopen/src/include/miopen/batchnorm/solvers.hpp +++ b/projects/miopen/src/include/miopen/batchnorm/solvers.hpp @@ -130,7 +130,7 @@ struct BnBwdTrainingSpatial final : BatchNormTunableSolver - Find(const ExecutionContext& ctx, + Find(ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -93,7 +93,7 @@ class ISolversFinder const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters) const = 0; [[nodiscard]] virtual std::vector - FindImpl(const ExecutionContext& ctx, + FindImpl(ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -113,7 +113,7 @@ class SolversFinderMixin : public ISolversFinder } [[nodiscard]] std::vector - FindImpl(const ExecutionContext& ctx, + FindImpl(ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -140,7 +140,7 @@ class SolversFinderMixin : public ISolversFinder GetAlgorithmName(const ProblemDescription& problem) const = 0; [[nodiscard]] virtual std::vector - FindImpl(const ExecutionContext& ctx, + FindImpl(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const FindParameters& parameters, diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index 5fa56cb8da1..52677c4e83d 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/projects/miopen/src/include/miopen/conv/solvers.hpp @@ -123,7 +123,7 @@ struct ConvAsm3x3U final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm3x3U&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm3x3U - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -229,7 +229,7 @@ struct ConvAsm1x1U final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm1x1U&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm1x1U - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -320,7 +320,7 @@ struct ConvAsm1x1UV2 final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm1x1UV2&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm1x1UV2 - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -822,7 +822,7 @@ struct ConvHipImplicitGemmV4R1Fwd final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1001,7 +1001,7 @@ struct ConvMlirIgemmFwdXdlops final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1078,7 +1078,7 @@ struct ConvMlirIgemmWrWXdlops final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1516,7 +1516,7 @@ struct ConvMlirIgemmBwdXdlops final : ConvTunableSolver&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvOclBwdWrw2 - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -2664,7 +2664,7 @@ struct ConvHipImplicitGemmWrwV4R4Xdlops final const miopen::conv::ProblemDescription&, const PerformanceImplicitGemmWrwV4R4Xdlops&) const override; MIOPEN_INTERNALS_EXPORT PerformanceImplicitGemmWrwV4R4Xdlops - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; }; @@ -2761,7 +2761,7 @@ struct ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm final const miopen::conv::ProblemDescription&, const PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm - Search(const ExecutionContext&, + Search(ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; }; @@ -2814,7 +2814,7 @@ struct ConvCkIgemmFwdV6r1DlopsNchw final : ConvTunableSolver std::vector - SearchForAllSolutions(const Context& ctx, + SearchForAllSolutions(Context& ctx, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, diff --git a/projects/miopen/src/include/miopen/fusion/solvers.hpp b/projects/miopen/src/include/miopen/fusion/solvers.hpp index ec472b455a9..0363befe93e 100644 --- a/projects/miopen/src/include/miopen/fusion/solvers.hpp +++ b/projects/miopen/src/include/miopen/fusion/solvers.hpp @@ -76,7 +76,7 @@ struct ConvBiasActivAsm1x1U : FusionTunableSolver, TunableSolv /// Search virtual PerformanceConfig - Search(const Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; + Search(Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; /// Tunable solvers provide a GetSolution that takes a Context and PerformanceConfig virtual ConvSolution GetSolution(const Context& ctx, diff --git a/projects/miopen/src/mlo_dir_conv.cpp b/projects/miopen/src/mlo_dir_conv.cpp index 7ca8c342d2c..08e02e90b01 100644 --- a/projects/miopen/src/mlo_dir_conv.cpp +++ b/projects/miopen/src/mlo_dir_conv.cpp @@ -220,7 +220,7 @@ auto miopen::MakeConvDbGetter(const ExecutionContext& ctx) -> DbGetter } std::vector -FindAllGemmSolutions(const miopen::ExecutionContext& ctx, +FindAllGemmSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -236,7 +236,7 @@ AllGemmWorkspaceSize(const miopen::ExecutionContext& ctx, } std::vector -FindAllDirectSolutions(const miopen::ExecutionContext& ctx, +FindAllDirectSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -252,7 +252,7 @@ AllDirectForwardBackwardDataWorkspaceSize(const miopen::ExecutionContext& ctx, } std::vector> -FindAllWinogradWorkspaceSizes(const miopen::ExecutionContext& ctx, +FindAllWinogradWorkspaceSizes(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem) { return GetWindogradSolvers().GetWorkspaceSizes(ctx, problem); @@ -266,14 +266,14 @@ FindWinogradWrWWorkspaceSizes(const miopen::ExecutionContext& ctx, } std::vector> -FindAllImplicitGemmWorkspaceSizes(const miopen::ExecutionContext& ctx, +FindAllImplicitGemmWorkspaceSizes(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem) { return GetImplicitGemmSolvers().GetWorkspaceSizes(ctx, problem); } std::vector -FindAllImplicitGemmSolutions(const miopen::ExecutionContext& ctx, +FindAllImplicitGemmSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -282,7 +282,7 @@ FindAllImplicitGemmSolutions(const miopen::ExecutionContext& ctx, } std::vector -FindAllWinogradSolutions(const miopen::ExecutionContext& ctx, +FindAllWinogradSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -323,7 +323,7 @@ FindImplicitGemmWrWAllSolutions(const miopen::ExecutionContext& ctx, } std::vector -FindAllBwdWrW2DSolutions(const miopen::ExecutionContext& ctx, +FindAllBwdWrW2DSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -332,7 +332,7 @@ FindAllBwdWrW2DSolutions(const miopen::ExecutionContext& ctx, } std::vector -FindAllFFTSolutions(const miopen::ExecutionContext& ctx, +FindAllFFTSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { diff --git a/projects/miopen/src/solver.cpp b/projects/miopen/src/solver.cpp index 699844ddbf9..f99b143e2fc 100644 --- a/projects/miopen/src/solver.cpp +++ b/projects/miopen/src/solver.cpp @@ -724,7 +724,7 @@ inline SolverRegistrar::SolverRegistrar(IdRegistryData& registry) Primitive::Fusion, fusion::ConvCKIgemmGrpFwdActivFused{}.SolverDbId(), miopenConvolutionAlgoImplicitGEMM); - + Register(registry, ++id, Primitive::Normalization, layernorm::LayernormBackward().SolverDbId()); // IMPORTANT: New solvers should be added to the end of the function, and don't leave a white // space between this comment and the newly registered solver(s)! diff --git a/projects/miopen/src/solver/batchnorm/backward_spatial.cpp b/projects/miopen/src/solver/batchnorm/backward_spatial.cpp index 26a386add1f..1eba75aca70 100644 --- a/projects/miopen/src/solver/batchnorm/backward_spatial.cpp +++ b/projects/miopen/src/solver/batchnorm/backward_spatial.cpp @@ -184,7 +184,7 @@ bool BnBwdTrainingSpatial::IsValidPerformanceConfig( } PerformanceConfigBnBwdBackward -BnBwdTrainingSpatial::Search(const ExecutionContext& ctx, +BnBwdTrainingSpatial::Search(ExecutionContext& ctx, const miopen::batchnorm::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/batchnorm/forward_spatial.cpp b/projects/miopen/src/solver/batchnorm/forward_spatial.cpp index 5f7dbc57e50..0fe64589b78 100644 --- a/projects/miopen/src/solver/batchnorm/forward_spatial.cpp +++ b/projects/miopen/src/solver/batchnorm/forward_spatial.cpp @@ -173,7 +173,7 @@ bool BnFwdTrainingSpatial::IsValidPerformanceConfig( } PerformanceConfigBnFwdTraining -BnFwdTrainingSpatial::Search(const ExecutionContext& ctx, +BnFwdTrainingSpatial::Search(ExecutionContext& ctx, const miopen::batchnorm::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp b/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp index 6d7e6761492..a31c83afd88 100644 --- a/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp @@ -911,7 +911,7 @@ ConvSolution ConvAsm1x1U::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm1x1U ConvAsm1x1U::Search(const ExecutionContext& ctx, +PerformanceConfigConvAsm1x1U ConvAsm1x1U::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp b/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp index 420d37d2174..0b899e6831d 100644 --- a/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp @@ -753,7 +753,7 @@ ConvSolution ConvAsm1x1UV2::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm1x1UV2 ConvAsm1x1UV2::Search(const ExecutionContext& ctx, +PerformanceConfigConvAsm1x1UV2 ConvAsm1x1UV2::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp b/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp index 839517e3c7c..33805c38558 100644 --- a/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp @@ -320,7 +320,7 @@ ConvSolution ConvAsm3x3U::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm3x3U ConvAsm3x3U::Search(const ExecutionContext& ctx, +PerformanceConfigConvAsm3x3U ConvAsm3x3U::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp index 53486560097..8272537a4cd 100644 --- a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp @@ -853,7 +853,7 @@ ConvSolution ConvAsmBwdWrW1x1::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsmBwdWrW1x1 ConvAsmBwdWrW1x1::Search(const ExecutionContext& ctx, +PerformanceConfigConvAsmBwdWrW1x1 ConvAsmBwdWrW1x1::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp index 6d49a26fdf4..8db941cd038 100644 --- a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp @@ -597,7 +597,7 @@ ConvSolution ConvAsmBwdWrW3x3::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigAsmDirect3x3WrW ConvAsmBwdWrW3x3::Search(const ExecutionContext& ctx, +PerformanceConfigAsmDirect3x3WrW ConvAsmBwdWrW3x3::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index 117fba50c2f..95f831fee88 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -951,7 +951,7 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(const ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp index 041e69b9261..48870e3478d 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp @@ -541,7 +541,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCFwdDlopsNCHWC -ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::Search(const ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index 93d2e256685..8ebdaa2e8e7 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -804,7 +804,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCFwdXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::Search(const ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 1779f124cbb..6f1b622085c 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -850,7 +850,7 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCWrwXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(const ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp b/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp index 61dba30c7f6..72dff87f0fe 100644 --- a/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp +++ b/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp @@ -268,7 +268,7 @@ std::size_t ConvCkIgemmFwdV6r1DlopsNchw::GetWorkspaceSize(const ExecutionContext } PerformanceConvCkIgemmFwdV6r1DlopsNchw -ConvCkIgemmFwdV6r1DlopsNchw::Search(const ExecutionContext& ctx, +ConvCkIgemmFwdV6r1DlopsNchw::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp index 78a09c392b6..2c7dc6bfa2d 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp @@ -498,7 +498,7 @@ ConvHipImplicitGemm3DGroupBwdXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupBwdXdlops -ConvHipImplicitGemm3DGroupBwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupBwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp index 751bd964c9d..f7907c93d46 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp @@ -561,7 +561,7 @@ ConvHipImplicitGemm3DGroupFwdXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupFwdXdlops -ConvHipImplicitGemm3DGroupFwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupFwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp index 0c15d775091..078c95f6440 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp @@ -473,7 +473,7 @@ ConvHipImplicitGemm3DGroupWrwXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupWrwXdlops -ConvHipImplicitGemm3DGroupWrwXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupWrwXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp index 9f7bc5f0876..1c6a697990a 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp @@ -261,7 +261,7 @@ bool ConvHipImplicitGemmBwdXdlops::IsValidPerformanceConfig( } PerformanceConfigHipImplicitGemmBwdXdlops -ConvHipImplicitGemmBwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmBwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp index 47a8a6c0ca3..5f69c52a3fd 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -709,7 +709,7 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV1R1 -ConvHipImplicitGemmBwdDataV1R1::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV1R1::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp index d29f6b3f49d..0502ca0eaac 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp @@ -831,7 +831,7 @@ bool ConvHipImplicitGemmBwdDataV1R1Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmBwdV1R1Xdlops -ConvHipImplicitGemmBwdDataV1R1Xdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV1R1Xdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp index a8d38c7dac3..4f00fc40e71 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp @@ -819,7 +819,7 @@ bool ConvHipImplicitGemmBwdDataV4R1::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV4R1 -ConvHipImplicitGemmBwdDataV4R1::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV4R1::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp index 63677941afe..23723f2beeb 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp @@ -902,7 +902,7 @@ bool ConvHipImplicitGemmBwdDataV4R1Xdlops::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV4R1Xdlops -ConvHipImplicitGemmBwdDataV4R1Xdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV4R1Xdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp index 15dffe49f46..acca114e953 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp @@ -197,14 +197,14 @@ bool ConvHipImplicitGemmV4R1WrW::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R1 -ConvHipImplicitGemmV4R1Fwd::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmV4R1Fwd::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { return GenericSearch(*this, ctx, problem, invoke_ctx); } PerformanceImplicitGemmV4R1 -ConvHipImplicitGemmV4R1WrW::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmV4R1WrW::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp index 597553bfee3..f0fc97cb2f4 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp @@ -639,7 +639,7 @@ bool ConvHipImplicitGemmV4R4Fwd::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R4Fwd -ConvHipImplicitGemmV4R4Fwd::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmV4R4Fwd::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp index cbe3edc4a90..b360a903a1c 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp @@ -1053,7 +1053,7 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmForwardV4R4Xdlops -ConvHipImplicitGemmForwardV4R4Xdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R4Xdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp index db410f306dc..4f09e27d56c 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp @@ -1117,7 +1117,7 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::IsApplicable( } PerformanceImplicitGemmForwardV4R4Xdlops_Padded_Gemm -ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp index 145e53b4dab..1aa7100653e 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp @@ -1090,7 +1090,7 @@ bool ConvHipImplicitGemmForwardV4R5Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmForwardV4R5Xdlops -ConvHipImplicitGemmForwardV4R5Xdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R5Xdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp index ece0e1b1197..caad7df232b 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp @@ -263,7 +263,7 @@ bool ConvHipImplicitGemmFwdXdlops::IsValidPerformanceConfig( } PerformanceConfigHipImplicitGemmFwdXdlops -ConvHipImplicitGemmFwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmFwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp index 4533508e15b..e876ef1b7fc 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp @@ -534,7 +534,7 @@ size_t ConvHipImplicitGemmGroupBwdXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupBwdXdlops -ConvHipImplicitGemmGroupBwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmGroupBwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp index 173ad0c74be..140a43e8424 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp @@ -495,7 +495,7 @@ size_t ConvHipImplicitGemmGroupFwdXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupFwdXdlops -ConvHipImplicitGemmGroupFwdXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmGroupFwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp index f6a38370b43..3442366f2a7 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp @@ -574,7 +574,7 @@ size_t ConvHipImplicitGemmGroupWrwXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupWrwXdlops -ConvHipImplicitGemmGroupWrwXdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmGroupWrwXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp index e99b8c225f3..3225ca550a7 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp @@ -642,7 +642,7 @@ bool ConvHipImplicitGemmV4R4WrW::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R4WrW -ConvHipImplicitGemmV4R4WrW::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmV4R4WrW::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp index faa244742b1..47787c94eed 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp @@ -1119,7 +1119,7 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops::IsApplicable(const ExecutionContext& ctx, } PerformanceImplicitGemmWrwV4R4Xdlops -ConvHipImplicitGemmWrwV4R4Xdlops::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmWrwV4R4Xdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp index 4d807c4c7b7..bfad02e6e3d 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp @@ -1198,7 +1198,7 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::IsApplicable( } PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm -ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::Search(const ExecutionContext& ctx, +ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp index 03357d02517..7fa4cfc8b6b 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp @@ -91,7 +91,7 @@ bool ConvMlirIgemmBwd::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmBwd::Search(const ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmBwd::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp index 79d6ff3694e..068f4f585fb 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp @@ -87,7 +87,7 @@ bool ConvMlirIgemmBwdXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmBwdXdlops::Search(const ExecutionContext& ctx, +ConvMlirIgemmBwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp index 7d03507d392..4f6f58fcf01 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp @@ -153,7 +153,7 @@ bool ConvMlirIgemmFwd::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmFwd::Search(const ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmFwd::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp index 02865f6e9ab..01c91b8ecb6 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp @@ -212,7 +212,7 @@ bool ConvMlirIgemmFwdXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmFwdXdlops::Search(const ExecutionContext& ctx, +ConvMlirIgemmFwdXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp index 64b19192071..fcc9091d9b8 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp @@ -92,7 +92,7 @@ bool ConvMlirIgemmWrW::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmWrW::Search(const ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmWrW::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp index 15ef9363ab9..791908acb35 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp @@ -88,7 +88,7 @@ bool ConvMlirIgemmWrWXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmWrWXdlops::Search(const ExecutionContext& ctx, +ConvMlirIgemmWrWXdlops::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp index b6a177d31b5..d99403d4389 100644 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp +++ b/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp @@ -770,7 +770,7 @@ ConvSolution ConvOclBwdWrW2::GetSolution( template PerformanceConfigConvOclBwdWrw2 -ConvOclBwdWrW2::Search(const ExecutionContext& ctx, +ConvOclBwdWrW2::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp index 213c989f376..f7c786d9166 100644 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp +++ b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp @@ -213,7 +213,7 @@ static int MeasurePerfConfig(const Handle& handle, } LegacyPerformanceConfig -ConvOclDirectFwdLegacyExhaustiveSearch::Search(const ExecutionContext& ctx, +ConvOclDirectFwdLegacyExhaustiveSearch::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_winoRxS.cpp b/projects/miopen/src/solver/conv/conv_winoRxS.cpp index 2ec293fef73..3b8f0861cd9 100644 --- a/projects/miopen/src/solver/conv/conv_winoRxS.cpp +++ b/projects/miopen/src/solver/conv/conv_winoRxS.cpp @@ -429,7 +429,7 @@ bool ConvBinWinoRxS::IsValidPerformanceConfig( template PerformanceConfigConvBinWinogradRxS -ConvBinWinoRxS::Search(const ExecutionContext& ctx, +ConvBinWinoRxS::Search(ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp index 5b1e4b0a006..fc7ac3c0fa0 100644 --- a/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp @@ -85,7 +85,7 @@ bool ConvBiasActivAsm1x1U::IsValidPerformanceConfig( } PerformanceConfigConvBiasActivAsm1x1U -ConvBiasActivAsm1x1U::Search(const FusionContext& context, +ConvBiasActivAsm1x1U::Search(FusionContext& context, const FusionDescription& problem, const AnyInvokeParams& invoke_params) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp index 0d3a82fab7a..cfaa85479e5 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp @@ -387,7 +387,7 @@ bool ConvCKIgemmFwdBiasActivFused::IsValidPerformanceConfig( } PerformanceConfigConvCKIgemmFwdBiasActivFused -ConvCKIgemmFwdBiasActivFused::Search(const FusionContext& ctx, +ConvCKIgemmFwdBiasActivFused::Search(FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp index 0495d4598bd..547d24b2b80 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp @@ -370,7 +370,7 @@ bool ConvCKIgemmFwdBiasResAddActivFused::IsValidPerformanceConfig( } PerfConfigConvCKIgemmFwdBiasResAddActivFused -ConvCKIgemmFwdBiasResAddActivFused::Search(const FusionContext& ctx, +ConvCKIgemmFwdBiasResAddActivFused::Search(FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp index 03c4f35874e..a088f4a7005 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp @@ -581,7 +581,7 @@ size_t ConvCKIgemmGrpFwdActivFused::GetWorkspaceSize(const FusionContext&, } PerformanceConfigConvCKIgemmGrpFwdActivFused -ConvCKIgemmGrpFwdActivFused::Search(const FusionContext& ctx, +ConvCKIgemmGrpFwdActivFused::Search(FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp index fed835c8b60..1e4679eb700 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp @@ -592,7 +592,7 @@ ConvCKIgemmGrpFwdBiasActivFused::GetWorkspaceSize(const FusionContext&, } PerformanceConfigConvCKIgemmGrpFwdBiasActivFused -ConvCKIgemmGrpFwdBiasActivFused::Search(const FusionContext& ctx, +ConvCKIgemmGrpFwdBiasActivFused::Search(FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp index 6a1a9f35877..d3d2e1e6b5f 100644 --- a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp +++ b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp @@ -43,7 +43,7 @@ namespace solver { namespace fusion { PerformanceConfigConvOclDirectFwdFused -ConvOclDirectFwdFused::Search(const FusionContext& context, +ConvOclDirectFwdFused::Search(FusionContext& context, const FusionDescription& problem, const AnyInvokeParams& invoke_params) const { From ca3384b4fa831c98b1df3fabe447e5c657a4f257 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 5 Sep 2025 22:44:39 +0000 Subject: [PATCH 03/26] removing const modifiers --- projects/miopen/src/adam.cpp | 2 +- projects/miopen/src/addlayernorm.cpp | 2 +- projects/miopen/src/cat.cpp | 2 +- projects/miopen/src/find_controls.cpp | 7 ++++ projects/miopen/src/fusion.cpp | 8 ++--- projects/miopen/src/getitem.cpp | 2 +- projects/miopen/src/glu.cpp | 4 +-- projects/miopen/src/groupnorm.cpp | 2 +- projects/miopen/src/include/miopen/activ.hpp | 4 +-- .../miopen/src/include/miopen/any_solver.hpp | 18 +++++----- .../src/include/miopen/find_solution.hpp | 12 +++---- .../src/include/miopen/generic_search.hpp | 4 +-- .../src/include/miopen/mlo_internal.hpp | 16 ++++----- projects/miopen/src/include/miopen/rnn.hpp | 30 ++++++++-------- projects/miopen/src/include/miopen/solver.hpp | 4 +-- projects/miopen/src/kthvalue.cpp | 2 +- projects/miopen/src/layernorm.cpp | 4 +-- projects/miopen/src/mlo_dir_conv.cpp | 4 +-- projects/miopen/src/multimarginloss.cpp | 2 +- projects/miopen/src/ocl/activ_ocl.cpp | 4 +-- projects/miopen/src/ocl/convolutionocl.cpp | 34 +++++++++---------- projects/miopen/src/ocl/pooling_ocl.cpp | 4 +-- projects/miopen/src/ocl/rnnocl.cpp | 18 +++++----- projects/miopen/src/prelu.cpp | 2 +- projects/miopen/src/reducecalculation.cpp | 2 +- projects/miopen/src/reduceextreme.cpp | 2 +- projects/miopen/src/rnn.cpp | 14 ++++---- .../src/rnn/Solutions/rnn_transformer.cpp | 6 ++-- projects/miopen/src/rope.cpp | 4 +-- projects/miopen/src/softmarginloss.cpp | 4 +-- .../conv/conv_MP_bidirectional_winograd.cpp | 4 +-- .../src/solver/conv_ocl_dir2Dfwd_fused.cpp | 2 +- projects/miopen/src/t5layernorm.cpp | 4 +-- projects/miopen/src/transformers_adam_w.cpp | 2 +- 34 files changed, 120 insertions(+), 115 deletions(-) diff --git a/projects/miopen/src/adam.cpp b/projects/miopen/src/adam.cpp index 16e7b23d7b9..fedc32791d3 100644 --- a/projects/miopen/src/adam.cpp +++ b/projects/miopen/src/adam.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t Adam(const Handle& handle, +miopenStatus_t Adam(Handle& handle, const TensorDescriptor& paramInDesc, ConstData_t paramIn, const TensorDescriptor& paramOutDesc, diff --git a/projects/miopen/src/addlayernorm.cpp b/projects/miopen/src/addlayernorm.cpp index e0ab9e1ade5..271c81ca60f 100644 --- a/projects/miopen/src/addlayernorm.cpp +++ b/projects/miopen/src/addlayernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t AddLayerNormForward(const Handle& handle, +miopenStatus_t AddLayerNormForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& x2Desc, diff --git a/projects/miopen/src/cat.cpp b/projects/miopen/src/cat.cpp index e4e35a98711..85e8e89ea73 100644 --- a/projects/miopen/src/cat.cpp +++ b/projects/miopen/src/cat.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t CatForward(const Handle& handle, +miopenStatus_t CatForward(Handle& handle, int32_t xCount, const TensorDescriptor* const* xDescs, ConstData_t* xs, diff --git a/projects/miopen/src/find_controls.cpp b/projects/miopen/src/find_controls.cpp index 4e87575cd89..6eb332fc69d 100644 --- a/projects/miopen/src/find_controls.cpp +++ b/projects/miopen/src/find_controls.cpp @@ -62,6 +62,8 @@ static_assert(FindEnforceAction::Search == static_cast(miopenTuningPolicySearch)); static_assert(FindEnforceAction::SearchDbUpdate == static_cast(miopenTuningPolicySearchDbUpdate)); +static_assert(FindEnforceAction::SearchCutoffDbUpdate == + static_cast(miopenTuningPolicySearchCutoffDbUpdate)); static_assert(FindEnforceAction::DbClean == static_cast(miopenTuningPolicyDbClean)); @@ -75,6 +77,7 @@ const char* ToCString(const FindEnforceAction mode) case FindEnforceAction::DbUpdate: return "DB_UPDATE"; case FindEnforceAction::Search: return "SEARCH"; case FindEnforceAction::SearchDbUpdate: return "SEARCH_DB_UPDATE"; + case FindEnforceAction::SearchCutoffDbUpdate: return "SEARCH_CUTOFF_DB_UPDATE"; case FindEnforceAction::DbClean: return "CLEAN"; } return ""; @@ -103,6 +106,10 @@ FindEnforceAction GetFindEnforceActionImpl() { return FindEnforceAction::SearchDbUpdate; } + else if(str == "SEARCH_CUTOFF_DB_UPDATE") + { + return FindEnforceAction::SearchCutoffDbUpdate; + } else if(str == "DB_CLEAN") { return FindEnforceAction::DbClean; diff --git a/projects/miopen/src/fusion.cpp b/projects/miopen/src/fusion.cpp index 7694e2e59fc..284580e747e 100644 --- a/projects/miopen/src/fusion.cpp +++ b/projects/miopen/src/fusion.cpp @@ -787,7 +787,7 @@ static auto GetAllFusionSolvers() GetFusedWinogradSolvers(); } -solver::ConvSolution MakeFusedSolution(const FusionContext& ctx, +solver::ConvSolution MakeFusedSolution(FusionContext& ctx, solver::Id id, const std::optional& perf_cfg_override, const FusionDescription& problem, @@ -829,13 +829,13 @@ class FusionSolverFinder : public SolversFinderMixin - FindImpl(const ExecutionContext& ctx, + FindImpl(ExecutionContext& ctx, const FusionDescription& problem, const AnyInvokeParams& invoke_ctx, const FusionFindParameters&, const std::optional& options) const override { - const auto fusion_ctx = FusionContext(ctx); + auto fusion_ctx = FusionContext(ctx); return solvers.SearchForAllSolutions(fusion_ctx, problem, MakeConvDbGetter(ctx), @@ -1046,7 +1046,7 @@ miopenStatus_t FusionPlanDescriptor::Compile(const Handle& handle) const auto id = solver::Id{sol->solution_id}; GetAllFusionSolvers().FindById(id, [&](auto solver) { - const auto ctx = FusionContext{handle}; + auto ctx = FusionContext{handle}; auto db_getter = MakeConvDbGetter(ctx); const auto solution = solver::FindSolution( solver, ctx, fusion_problem, db_getter, {}); // auto tune is not expected here diff --git a/projects/miopen/src/getitem.cpp b/projects/miopen/src/getitem.cpp index 2920708932d..e06e6299b15 100644 --- a/projects/miopen/src/getitem.cpp +++ b/projects/miopen/src/getitem.cpp @@ -50,7 +50,7 @@ std::size_t GetGetitemWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t GetitemBackward(const Handle& handle, +miopenStatus_t GetitemBackward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/glu.cpp b/projects/miopen/src/glu.cpp index 8198549bd04..9d1181d32e0 100644 --- a/projects/miopen/src/glu.cpp +++ b/projects/miopen/src/glu.cpp @@ -37,7 +37,7 @@ namespace miopen { namespace glu { -miopenStatus_t GLUForward(const Handle& handle, +miopenStatus_t GLUForward(Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputDesc, @@ -65,7 +65,7 @@ miopenStatus_t GLUForward(const Handle& handle, return miopenStatusSuccess; } -miopenStatus_t GLUBackward(const Handle& handle, +miopenStatus_t GLUBackward(Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputGradDesc, diff --git a/projects/miopen/src/groupnorm.cpp b/projects/miopen/src/groupnorm.cpp index dd12a5f593a..2cf21102af2 100644 --- a/projects/miopen/src/groupnorm.cpp +++ b/projects/miopen/src/groupnorm.cpp @@ -33,7 +33,7 @@ namespace miopen { -miopenStatus_t GroupNormForward(const Handle& handle, +miopenStatus_t GroupNormForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, diff --git a/projects/miopen/src/include/miopen/activ.hpp b/projects/miopen/src/include/miopen/activ.hpp index b7db13b4782..26ce91112f4 100644 --- a/projects/miopen/src/include/miopen/activ.hpp +++ b/projects/miopen/src/include/miopen/activ.hpp @@ -50,7 +50,7 @@ struct MIOPEN_INTERNALS_EXPORT ActivationDescriptor : miopenActivationDescriptor double GetBeta() const; double GetGamma() const; - miopenStatus_t Forward(const Handle& handle, + miopenStatus_t Forward(Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -60,7 +60,7 @@ struct MIOPEN_INTERNALS_EXPORT ActivationDescriptor : miopenActivationDescriptor size_t xOffset = 0, size_t yOffset = 0) const; - miopenStatus_t Backward(const Handle& handle, + miopenStatus_t Backward(Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t y, diff --git a/projects/miopen/src/include/miopen/any_solver.hpp b/projects/miopen/src/include/miopen/any_solver.hpp index 7e610719dae..636fa9e9adc 100644 --- a/projects/miopen/src/include/miopen/any_solver.hpp +++ b/projects/miopen/src/include/miopen/any_solver.hpp @@ -98,7 +98,7 @@ struct AnySolver bool IsEmpty() const { return ptr_value == nullptr; }; #if defined(FIN_ANY_SOLVER_FIND_SOLUTION_COMPAT) - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, @@ -109,7 +109,7 @@ struct AnySolver } #endif - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -119,7 +119,7 @@ struct AnySolver return ptr_value->FindSolution(ctx, problem, db_getter, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -194,17 +194,17 @@ struct AnySolver const miopen::conv::ProblemDescription& problem) const = 0; virtual const std::type_info& Type() const = 0; virtual std::string GetSolverDbId() const = 0; - virtual ConvSolution FindSolution(const ExecutionContext& ctx, + virtual ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; - virtual ConvSolution FindSolution(const ExecutionContext& ctx, + virtual ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; - virtual ConvSolution FindSolution(const ExecutionContext& ctx, + virtual ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -413,7 +413,7 @@ struct AnySolver return value.GetWti(ctx, problem); } - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, @@ -422,7 +422,7 @@ struct AnySolver return miopen::solver::FindSolution(value, ctx, problem, db, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -432,7 +432,7 @@ struct AnySolver value, ctx, problem, db_getter, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(const ExecutionContext& ctx, + ConvSolution FindSolution(ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index e0506f804a7..0b70b3a29b4 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -52,7 +52,7 @@ namespace solver { template auto FindSolutionImpl(rank<1>, const Solver& s, - const Context& context, + Context& context, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, @@ -161,7 +161,7 @@ auto FindSolutionImpl(rank<1>, template auto FindSolutionImpl(rank<0>, const Solver& s, - const Context& context, + Context& context, const Problem& problem, Db&&, const AnyInvokeParams&, @@ -213,7 +213,7 @@ auto GetInvokeFactoryImpl( /// May read/write perfDb. template ConvSolution FindSolution(const Solver& s, - const Context& context, + Context& context, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, @@ -343,7 +343,7 @@ struct SolverContainer // Search for all applicable solutions among many solvers template std::vector - SearchForSolutions(const ExecutionContext& ctx, + SearchForSolutions(ExecutionContext& ctx, const Problem& problem, std::size_t limit = std::numeric_limits::max(), const AnyInvokeParams& invoke_params = {}) const @@ -451,7 +451,7 @@ struct SolverContainer } template - void ExecutePrimitive(const ExecutionContext& ctx, + void ExecutePrimitive(ExecutionContext& ctx, const Problem& problem, const AlgorithmName& algo, const AnyInvokeParams& invoke_params) const @@ -480,7 +480,7 @@ struct SolverContainer } template - void ExecutePrimitive(const Handle& handle, + void ExecutePrimitive(Handle& handle, const Problem& problem, const AlgorithmName& algo, const AnyInvokeParams& invoke_params) const diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 0fec650fdd0..fde433c0afa 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -423,7 +423,7 @@ struct SolutionPerf template auto GenericSearch(const Solver s, - const Context& context_, + Context& context_, const Problem& problem, const AnyInvokeParams& invoke_ctx_, std::vector* perf_solsp = nullptr) @@ -702,7 +702,7 @@ auto GenericSearch(const Solver s, // if using cutoff time for search and new cutoff is shorter, update if(context.search_cutoff) { - float new_cutoff = perf_sols.end().time * 2; + float new_cutoff = perf_sols.end()->time * 2; if(new_cutoff < context.generic_search_cutoff_time) context_.generic_search_cutoff_time = new_cutoff; } diff --git a/projects/miopen/src/include/miopen/mlo_internal.hpp b/projects/miopen/src/include/miopen/mlo_internal.hpp index fcf64a00764..63e2f3d4e31 100644 --- a/projects/miopen/src/include/miopen/mlo_internal.hpp +++ b/projects/miopen/src/include/miopen/mlo_internal.hpp @@ -174,7 +174,7 @@ auto mloConstruct(T& x) -> decltype(x.mloConstruct(), void()) } MIOPEN_INTERNALS_EXPORT std::vector -FindAllGemmSolutions(const miopen::ExecutionContext& ctx, +FindAllGemmSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); @@ -211,37 +211,37 @@ AllFFTForwardBackwardDataWorkspaceSize(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem); MIOPEN_INTERNALS_EXPORT std::vector -FindAllDirectSolutions(const miopen::ExecutionContext& ctx, +FindAllDirectSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllImplicitGemmSolutions(const miopen::ExecutionContext& ctx, +FindAllImplicitGemmSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllWinogradSolutions(const miopen::ExecutionContext& ctx, +FindAllWinogradSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindWinogradWrWAllSolutions(const miopen::ExecutionContext& ctx, +FindWinogradWrWAllSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindImplicitGemmWrWAllSolutions(const miopen::ExecutionContext& ctx, +FindImplicitGemmWrWAllSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllBwdWrW2DSolutions(const miopen::ExecutionContext& ctx, +FindAllBwdWrW2DSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllFFTSolutions(const miopen::ExecutionContext& ctx, +FindAllFFTSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); diff --git a/projects/miopen/src/include/miopen/rnn.hpp b/projects/miopen/src/include/miopen/rnn.hpp index 49e31910eea..ccbd96563b4 100644 --- a/projects/miopen/src/include/miopen/rnn.hpp +++ b/projects/miopen/src/include/miopen/rnn.hpp @@ -255,7 +255,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardTraining(const Handle& handle, + void RNNForwardTraining(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -276,7 +276,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardInference(const Handle& handle, + void RNNForwardInference(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -295,7 +295,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t workSpace, size_t workSpaceSize) const; - void RNNBackwardData(const Handle& handle, + void RNNBackwardData(Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t y, ConstData_t dy, @@ -316,7 +316,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardData(const Handle& handle, + void RNNBackwardData(Handle& handle, int seqLen, c_array_view yDesc, ConstData_t y, @@ -343,7 +343,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeights(const Handle& handle, + void RNNBackwardWeights(Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -357,7 +357,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeights(const Handle& handle, + void RNNBackwardWeights(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -430,7 +430,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t /*reserveSpaceSize*/) const; - void RNNTransformerForward(const Handle& handle, + void RNNTransformerForward(Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -448,7 +448,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNTransformerBackwardData(const Handle& handle, + void RNNTransformerBackwardData(Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -467,7 +467,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNTransformerBackwardWeights(const Handle& handle, + void RNNTransformerBackwardWeights(Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -479,7 +479,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNVanillaForward(const Handle& handle, + void RNNVanillaForward(Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -528,7 +528,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardTrainingPackedTensors(const Handle& handle, + void RNNForwardTrainingPackedTensors(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -549,7 +549,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardMS(const Handle& handle, + void RNNForwardMS(Handle& handle, std::vector& seq_array, const TensorDescriptor& xDesc, ConstData_t x, @@ -566,7 +566,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor size_t extra_space_size, miopenRNNFWDMode_t fwd_mode) const; - void RNNForwardInferencePacked(const Handle& handle, + void RNNForwardInferencePacked(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -585,7 +585,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t workSpace, size_t workSpaceSize) const; - void RNNBackwardDataPackedTensors(const Handle& handle, + void RNNBackwardDataPackedTensors(Handle& handle, int seqLen, c_array_view dyDesc, ConstData_t dy, @@ -605,7 +605,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeightsPackedTensors(const Handle& handle, + void RNNBackwardWeightsPackedTensors(Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, diff --git a/projects/miopen/src/include/miopen/solver.hpp b/projects/miopen/src/include/miopen/solver.hpp index 772fe3161cb..ffa65a08d3b 100644 --- a/projects/miopen/src/include/miopen/solver.hpp +++ b/projects/miopen/src/include/miopen/solver.hpp @@ -162,7 +162,7 @@ struct SolverInterfaceTunable : SolverInterface { /// This function is a simplified version of FindSolution(), it does not obey search parameters /// from the Context and does not use the database. Intended to be used in unit tests. - virtual ConvSolution FindSolutionSimple(const Context& ctx, + virtual ConvSolution FindSolutionSimple(Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; }; @@ -213,7 +213,7 @@ struct SolverBaseTunable : SolverInterfaceTunable, TunableSolv const Problem& problem, const PerformanceConfig& config) const = 0; - ConvSolution FindSolutionSimple(const Context& ctx, + ConvSolution FindSolutionSimple(Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const final { diff --git a/projects/miopen/src/kthvalue.cpp b/projects/miopen/src/kthvalue.cpp index 02dd9f96185..a9f4e730678 100644 --- a/projects/miopen/src/kthvalue.cpp +++ b/projects/miopen/src/kthvalue.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t KthvalueForward(const Handle& handle, +miopenStatus_t KthvalueForward(Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputDesc, diff --git a/projects/miopen/src/layernorm.cpp b/projects/miopen/src/layernorm.cpp index 7529d1ee7f6..b84824e52fd 100644 --- a/projects/miopen/src/layernorm.cpp +++ b/projects/miopen/src/layernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t LayerNormForward(const Handle& handle, +miopenStatus_t LayerNormForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, @@ -105,7 +105,7 @@ std::size_t GetLayerNormBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t LayerNormBackward(const Handle& handle, +miopenStatus_t LayerNormBackward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/mlo_dir_conv.cpp b/projects/miopen/src/mlo_dir_conv.cpp index 08e02e90b01..78028e75b4a 100644 --- a/projects/miopen/src/mlo_dir_conv.cpp +++ b/projects/miopen/src/mlo_dir_conv.cpp @@ -291,7 +291,7 @@ FindAllWinogradSolutions(miopen::ExecutionContext& ctx, } std::vector -FindWinogradWrWAllSolutions(const miopen::ExecutionContext& ctx, +FindWinogradWrWAllSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -314,7 +314,7 @@ FindImplicitGemmWrWWorkspaceSizes(const miopen::ExecutionContext& ctx, } std::vector -FindImplicitGemmWrWAllSolutions(const miopen::ExecutionContext& ctx, +FindImplicitGemmWrWAllSolutions(miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { diff --git a/projects/miopen/src/multimarginloss.cpp b/projects/miopen/src/multimarginloss.cpp index e3daa240120..5638ef6f704 100644 --- a/projects/miopen/src/multimarginloss.cpp +++ b/projects/miopen/src/multimarginloss.cpp @@ -55,7 +55,7 @@ std::size_t GetMultiMarginLossForwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t MultiMarginLossForward(const Handle& handle, +miopenStatus_t MultiMarginLossForward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& iDesc, diff --git a/projects/miopen/src/ocl/activ_ocl.cpp b/projects/miopen/src/ocl/activ_ocl.cpp index 185cfd1d3de..fc301acaddf 100644 --- a/projects/miopen/src/ocl/activ_ocl.cpp +++ b/projects/miopen/src/ocl/activ_ocl.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t ActivationDescriptor::Forward(const Handle& handle, +miopenStatus_t ActivationDescriptor::Forward(Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -74,7 +74,7 @@ miopenStatus_t ActivationDescriptor::Forward(const Handle& handle, return miopenStatusSuccess; } -miopenStatus_t ActivationDescriptor::Backward(const Handle& handle, +miopenStatus_t ActivationDescriptor::Backward(Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t y, diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index 7a0fca77444..fe3598ee94e 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -255,7 +255,7 @@ static void ShrinkToFind10Results(std::vector& found) MIOPEN_INTERNALS_EXPORT std::vector -GetConvSolutions(const ExecutionContext& ctx, +GetConvSolutions(ExecutionContext& ctx, const conv::ProblemDescription& problem, const std::vector solutions) { @@ -352,7 +352,7 @@ bool HasGoodSolution(const std::vector solutions, return good_entry; } -std::vector VerifiedFDBSolution(const ExecutionContext& ctx, +std::vector VerifiedFDBSolution(ExecutionContext& ctx, const conv::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, bool force_attach_binary, @@ -362,10 +362,9 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, const auto& conv = problem.GetConv(); const auto& findMode = conv.findMode; auto results = UserFindDbRecord::TryLoad(ctx.GetStream(), problem, [&]() { - auto ctx_copy = ctx; - ctx_copy.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); + ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); const auto params = - conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx_copy, problem)}; + conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx, problem)}; auto conv_sols = GetConvSolutions(ctx, problem, solutions); auto eval_sols = EvaluateConvSolutions(ctx, problem, invoke_ctx, conv_sols, model_result); @@ -391,11 +390,11 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, { // entry considered bad, trigger tuning MIOPEN_LOG_I2("TrustVerify: Regenerate entry for user db"); - ctx_copy.do_search = true; - ctx_copy.db_update = true; + ctx.do_search = true; + ctx.db_update = true; return FindCore(invoke_ctx, - ctx_copy, + ctx, problem, params, conv::GetConvSolverFinders(), @@ -407,7 +406,7 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, return results; } -std::vector FindConvolution(const ExecutionContext& ctx, +std::vector FindConvolution(ExecutionContext& ctx, const conv::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, int requestAlgoCount, @@ -476,20 +475,19 @@ std::vector FindConvolution(const ExecutionContext& ctx, else { results = UserFindDbRecord::TryLoad(ctx.GetStream(), problem, [&]() { - auto ctx_copy = ctx; - ctx_copy.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); + ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); const auto params = - conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx_copy, problem)}; + conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx, problem)}; if(findMode.IsTrustVerify(ctx)) { MIOPEN_LOG_I2("TrustVerify: Generate entry for user db"); - ctx_copy.do_search = true; - ctx_copy.db_update = true; + ctx.do_search = true; + ctx.db_update = true; } return FindCore(invoke_ctx, - ctx_copy, + ctx, problem, params, conv::GetConvSolverFinders(), @@ -564,7 +562,7 @@ void ConvolutionDescriptor::FindConvFwdAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription(xDesc, wDesc, yDesc, *this, conv::Direction::Forward); - const auto ctx = [&] { + auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; @@ -1103,7 +1101,7 @@ void ConvolutionDescriptor::FindConvBwdDataAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription{dyDesc, wDesc, dxDesc, *this, conv::Direction::BackwardData}; - const auto ctx = [&] { + auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; @@ -1310,7 +1308,7 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription{dyDesc, dwDesc, xDesc, *this, conv::Direction::BackwardWeights}; - const auto ctx = [&] { + auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; diff --git a/projects/miopen/src/ocl/pooling_ocl.cpp b/projects/miopen/src/ocl/pooling_ocl.cpp index ab3a63520d8..9881c1596f4 100644 --- a/projects/miopen/src/ocl/pooling_ocl.cpp +++ b/projects/miopen/src/ocl/pooling_ocl.cpp @@ -54,7 +54,7 @@ static auto PoolingBackwardSolvers() solver::pooling::TransposedPoolingBwdNd>{}; } -miopenStatus_t PoolingDescriptor::Forward(const Handle& handle, +miopenStatus_t PoolingDescriptor::Forward(Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -139,7 +139,7 @@ miopenStatus_t PoolingDescriptor::Forward(const Handle& handle, return miopenStatusSuccess; } -miopenStatus_t PoolingDescriptor::Backward(const Handle& handle, +miopenStatus_t PoolingDescriptor::Backward(Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t /*y*/, diff --git a/projects/miopen/src/ocl/rnnocl.cpp b/projects/miopen/src/ocl/rnnocl.cpp index 5562f7dd389..8e940505820 100644 --- a/projects/miopen/src/ocl/rnnocl.cpp +++ b/projects/miopen/src/ocl/rnnocl.cpp @@ -257,7 +257,7 @@ miopenStatus_t ReducAddBias(const miopen::Handle& handle, } // namespace -void RNNDescriptor::RNNForwardMS(const Handle& handle, +void RNNDescriptor::RNNForwardMS(Handle& handle, std::vector& seq_array, const TensorDescriptor& xDesc, ConstData_t x, @@ -1177,7 +1177,7 @@ void RNNDescriptor::RNNForwardMS(const Handle& handle, } // Assuming sequence length is set to > 0 otherwise throw exception. -void RNNDescriptor::RNNForwardInference(const Handle& handle, +void RNNDescriptor::RNNForwardInference(Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -1310,7 +1310,7 @@ void RNNDescriptor::RNNForwardInference(const Handle& handle, } #endif } -void RNNDescriptor::RNNForwardInferencePacked(const Handle& handle, +void RNNDescriptor::RNNForwardInferencePacked(Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -2563,7 +2563,7 @@ void RNNDescriptor::RNNForwardInferencePacked(const Handle& handle, #endif } -void RNNDescriptor::RNNForwardTraining(const Handle& handle, +void RNNDescriptor::RNNForwardTraining(Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -2708,7 +2708,7 @@ void RNNDescriptor::RNNForwardTraining(const Handle& handle, }; void RNNDescriptor::RNNForwardTrainingPackedTensors( - const Handle& handle, + Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -4080,7 +4080,7 @@ void RNNDescriptor::RNNForwardTrainingPackedTensors( #endif }; -void RNNDescriptor::RNNBackwardData(const Handle& handle, +void RNNDescriptor::RNNBackwardData(Handle& handle, const int seqLen, c_array_view yDesc, ConstData_t y, @@ -4223,7 +4223,7 @@ void RNNDescriptor::RNNBackwardData(const Handle& handle, } void RNNDescriptor::RNNBackwardDataPackedTensors( - const Handle& handle, + Handle& handle, const int seqLen, c_array_view dyDesc, ConstData_t dy, @@ -5810,7 +5810,7 @@ void RNNDescriptor::RNNBackwardDataPackedTensors( #endif }; -void RNNDescriptor::RNNBackwardWeights(const Handle& handle, +void RNNDescriptor::RNNBackwardWeights(Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -5910,7 +5910,7 @@ void RNNDescriptor::RNNBackwardWeights(const Handle& handle, } void RNNDescriptor::RNNBackwardWeightsPackedTensors( - const Handle& handle, + Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, diff --git a/projects/miopen/src/prelu.cpp b/projects/miopen/src/prelu.cpp index 34ec9d59ab9..b43e5ea84bd 100644 --- a/projects/miopen/src/prelu.cpp +++ b/projects/miopen/src/prelu.cpp @@ -48,7 +48,7 @@ size_t GetPReLUBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t PReLUBackward(const Handle& handle, +miopenStatus_t PReLUBackward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& inputDesc, diff --git a/projects/miopen/src/reducecalculation.cpp b/projects/miopen/src/reducecalculation.cpp index d7bc58f89d5..ea4fda314d7 100644 --- a/projects/miopen/src/reducecalculation.cpp +++ b/projects/miopen/src/reducecalculation.cpp @@ -71,7 +71,7 @@ std::size_t GetReduceCalculationWorkspaceSize(const Handle& handle, return static_cast(-1); } -miopenStatus_t ReduceCalculationForward(const Handle& handle, +miopenStatus_t ReduceCalculationForward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& xDesc, diff --git a/projects/miopen/src/reduceextreme.cpp b/projects/miopen/src/reduceextreme.cpp index 37fbfd82dbe..a35df325986 100644 --- a/projects/miopen/src/reduceextreme.cpp +++ b/projects/miopen/src/reduceextreme.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t ReduceExtremeForward(const Handle& handle, +miopenStatus_t ReduceExtremeForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& indiceDesc, diff --git a/projects/miopen/src/rnn.cpp b/projects/miopen/src/rnn.cpp index 83fe4e5abdd..ce7a834d856 100644 --- a/projects/miopen/src/rnn.cpp +++ b/projects/miopen/src/rnn.cpp @@ -44,7 +44,7 @@ namespace miopen { -void profileRNNkernels(const Handle& handle, unsigned char select, float& ctime) +void profileRNNkernels(Handle& handle, unsigned char select, float& ctime) { float ktime = 0.; @@ -1207,7 +1207,7 @@ void RNNDescriptor::SeqTensorToTensorDescArray(const SeqTensorDescriptor& desc, }); } -void RNNDescriptor::RNNVanillaForward(const Handle& handle, +void RNNDescriptor::RNNVanillaForward(Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -1284,7 +1284,7 @@ void RNNDescriptor::RNNVanillaForward(const Handle& handle, } } -void RNNDescriptor::RNNVanillaBackwardData(const Handle& handle, +void RNNDescriptor::RNNVanillaBackwardData(Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -1335,7 +1335,7 @@ void RNNDescriptor::RNNVanillaBackwardData(const Handle& handle, reserveSpaceSize); } -void RNNDescriptor::RNNVanillaBackwardWeights(const Handle& handle, +void RNNDescriptor::RNNVanillaBackwardWeights(Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -1374,7 +1374,7 @@ void RNNDescriptor::RNNVanillaBackwardWeights(const Handle& handle, reserveSpaceSize); } -void RNNDescriptor::RNNForward(const Handle& handle, +void RNNDescriptor::RNNForward(Handle& handle, miopenRNNFWDMode_t fwdMode, const SeqTensorDescriptor& xDesc, ConstData_t x, @@ -1476,7 +1476,7 @@ void RNNDescriptor::RNNForward(const Handle& handle, #endif } -void RNNDescriptor::RNNBackwardData(const Handle& handle, +void RNNDescriptor::RNNBackwardData(Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t, ConstData_t dy, @@ -1584,7 +1584,7 @@ void RNNDescriptor::RNNBackwardData(const Handle& handle, #endif } -void RNNDescriptor::RNNBackwardWeights(const Handle& handle, +void RNNDescriptor::RNNBackwardWeights(Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, diff --git a/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp b/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp index 57c6866167d..72ebb560ee6 100644 --- a/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp +++ b/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp @@ -120,7 +120,7 @@ size_t RNNDescriptor::RNNTransformerWorkspaceSize(const SeqTensorDescriptor& xDe typeSize; } -void RNNDescriptor::RNNTransformerForward(const Handle& handle, +void RNNDescriptor::RNNTransformerForward(Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -278,7 +278,7 @@ void RNNDescriptor::RNNTransformerForward(const Handle& handle, } } -void RNNDescriptor::RNNTransformerBackwardData(const Handle& handle, +void RNNDescriptor::RNNTransformerBackwardData(Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -446,7 +446,7 @@ void RNNDescriptor::RNNTransformerBackwardData(const Handle& handle, } } -void RNNDescriptor::RNNTransformerBackwardWeights(const Handle& handle, +void RNNDescriptor::RNNTransformerBackwardWeights(Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, diff --git a/projects/miopen/src/rope.cpp b/projects/miopen/src/rope.cpp index d54f68bbd8a..9dca9b145e9 100644 --- a/projects/miopen/src/rope.cpp +++ b/projects/miopen/src/rope.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t RoPEForward(const Handle& handle, +miopenStatus_t RoPEForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& cosDesc, @@ -69,7 +69,7 @@ miopenStatus_t RoPEForward(const Handle& handle, return miopenStatusSuccess; } -miopenStatus_t RoPEBackward(const Handle& handle, +miopenStatus_t RoPEBackward(Handle& handle, const TensorDescriptor& dyDesc, ConstData_t dy, const TensorDescriptor& cosDesc, diff --git a/projects/miopen/src/softmarginloss.cpp b/projects/miopen/src/softmarginloss.cpp index a0b8072f76a..4ce38baa591 100644 --- a/projects/miopen/src/softmarginloss.cpp +++ b/projects/miopen/src/softmarginloss.cpp @@ -50,7 +50,7 @@ std::size_t GetSoftMarginLossForwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t SoftMarginLossForward(const Handle& handle, +miopenStatus_t SoftMarginLossForward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& iDesc, @@ -84,7 +84,7 @@ miopenStatus_t SoftMarginLossForward(const Handle& handle, return miopenStatusSuccess; } -miopenStatus_t SoftMarginLossBackward(const Handle& handle, +miopenStatus_t SoftMarginLossBackward(Handle& handle, const TensorDescriptor& iDesc, ConstData_t i, const TensorDescriptor& tDesc, diff --git a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp index 9b934ec041a..90067c24639 100644 --- a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp +++ b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp @@ -940,7 +940,7 @@ ConvMPBidirectWinograd_xdlops::G template PerformanceImplicitGemmForwardV4R4Xdlops ConvMPBidirectWinograd_xdlops::Search( - const ExecutionContext& ctx, + ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { @@ -948,7 +948,7 @@ ConvMPBidirectWinograd_xdlops::S GetTransformedInvokeContext(problem, invoke_ctx); const auto xdlops_problem = GetTransformedProblem(problem); - const auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); + auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); return ConvHipImplicitGemmForwardV4R4Xdlops().Search( xdlops_ctx, xdlops_problem, xdlops_invoke_ctx); diff --git a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp index d3d2e1e6b5f..e480db107a8 100644 --- a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp +++ b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp @@ -48,7 +48,7 @@ ConvOclDirectFwdFused::Search(FusionContext& context, const AnyInvokeParams& invoke_params) const { const auto conv_problem = problem.GetConvProblem(0, miopen::conv::Direction::Forward); - const auto conv_ctx = context.GetConvContext(conv_problem); + auto conv_ctx = context.GetConvContext(conv_problem); const auto legacy = conv::ConvOclDirectFwd{}; const auto& fusion_invoke_params = invoke_params.CastTo(); const auto wei_ocl_ptr = dynamic_cast( diff --git a/projects/miopen/src/t5layernorm.cpp b/projects/miopen/src/t5layernorm.cpp index 08b2600c53a..dacd51a208f 100644 --- a/projects/miopen/src/t5layernorm.cpp +++ b/projects/miopen/src/t5layernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t T5LayerNormForward(const Handle& handle, +miopenStatus_t T5LayerNormForward(Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, @@ -91,7 +91,7 @@ std::size_t GetT5LayerNormBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t T5LayerNormBackward(const Handle& handle, +miopenStatus_t T5LayerNormBackward(Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/transformers_adam_w.cpp b/projects/miopen/src/transformers_adam_w.cpp index 3b5253c9403..0e644d8344e 100644 --- a/projects/miopen/src/transformers_adam_w.cpp +++ b/projects/miopen/src/transformers_adam_w.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t TransformersAdamW(const Handle& handle, +miopenStatus_t TransformersAdamW(Handle& handle, const TensorDescriptor& paramInDesc, ConstData_t paramIn, const TensorDescriptor& paramOutDesc, From a303d234ddadf2e36dae43ed8aacc778d13f47a3 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 8 Sep 2025 18:09:53 +0000 Subject: [PATCH 04/26] use context copy for ExecutePrimitive --- projects/miopen/src/adam.cpp | 2 +- projects/miopen/src/addlayernorm.cpp | 2 +- projects/miopen/src/cat.cpp | 2 +- projects/miopen/src/getitem.cpp | 2 +- projects/miopen/src/glu.cpp | 4 +-- projects/miopen/src/groupnorm.cpp | 2 +- projects/miopen/src/include/miopen/activ.hpp | 4 +-- .../src/include/miopen/find_solution.hpp | 7 +++-- projects/miopen/src/include/miopen/rnn.hpp | 30 +++++++++---------- projects/miopen/src/kthvalue.cpp | 2 +- projects/miopen/src/layernorm.cpp | 4 +-- projects/miopen/src/multimarginloss.cpp | 2 +- projects/miopen/src/ocl/activ_ocl.cpp | 4 +-- projects/miopen/src/ocl/pooling_ocl.cpp | 4 +-- projects/miopen/src/ocl/rnnocl.cpp | 18 +++++------ projects/miopen/src/prelu.cpp | 2 +- projects/miopen/src/reducecalculation.cpp | 2 +- projects/miopen/src/reduceextreme.cpp | 2 +- projects/miopen/src/rnn.cpp | 14 ++++----- .../src/rnn/Solutions/rnn_transformer.cpp | 6 ++-- projects/miopen/src/rope.cpp | 4 +-- projects/miopen/src/softmarginloss.cpp | 4 +-- projects/miopen/src/t5layernorm.cpp | 4 +-- projects/miopen/src/transformers_adam_w.cpp | 2 +- 24 files changed, 65 insertions(+), 64 deletions(-) diff --git a/projects/miopen/src/adam.cpp b/projects/miopen/src/adam.cpp index fedc32791d3..16e7b23d7b9 100644 --- a/projects/miopen/src/adam.cpp +++ b/projects/miopen/src/adam.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t Adam(Handle& handle, +miopenStatus_t Adam(const Handle& handle, const TensorDescriptor& paramInDesc, ConstData_t paramIn, const TensorDescriptor& paramOutDesc, diff --git a/projects/miopen/src/addlayernorm.cpp b/projects/miopen/src/addlayernorm.cpp index 271c81ca60f..e0ab9e1ade5 100644 --- a/projects/miopen/src/addlayernorm.cpp +++ b/projects/miopen/src/addlayernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t AddLayerNormForward(Handle& handle, +miopenStatus_t AddLayerNormForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& x2Desc, diff --git a/projects/miopen/src/cat.cpp b/projects/miopen/src/cat.cpp index 85e8e89ea73..e4e35a98711 100644 --- a/projects/miopen/src/cat.cpp +++ b/projects/miopen/src/cat.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t CatForward(Handle& handle, +miopenStatus_t CatForward(const Handle& handle, int32_t xCount, const TensorDescriptor* const* xDescs, ConstData_t* xs, diff --git a/projects/miopen/src/getitem.cpp b/projects/miopen/src/getitem.cpp index e06e6299b15..2920708932d 100644 --- a/projects/miopen/src/getitem.cpp +++ b/projects/miopen/src/getitem.cpp @@ -50,7 +50,7 @@ std::size_t GetGetitemWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t GetitemBackward(Handle& handle, +miopenStatus_t GetitemBackward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/glu.cpp b/projects/miopen/src/glu.cpp index 9d1181d32e0..8198549bd04 100644 --- a/projects/miopen/src/glu.cpp +++ b/projects/miopen/src/glu.cpp @@ -37,7 +37,7 @@ namespace miopen { namespace glu { -miopenStatus_t GLUForward(Handle& handle, +miopenStatus_t GLUForward(const Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputDesc, @@ -65,7 +65,7 @@ miopenStatus_t GLUForward(Handle& handle, return miopenStatusSuccess; } -miopenStatus_t GLUBackward(Handle& handle, +miopenStatus_t GLUBackward(const Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputGradDesc, diff --git a/projects/miopen/src/groupnorm.cpp b/projects/miopen/src/groupnorm.cpp index 2cf21102af2..dd12a5f593a 100644 --- a/projects/miopen/src/groupnorm.cpp +++ b/projects/miopen/src/groupnorm.cpp @@ -33,7 +33,7 @@ namespace miopen { -miopenStatus_t GroupNormForward(Handle& handle, +miopenStatus_t GroupNormForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, diff --git a/projects/miopen/src/include/miopen/activ.hpp b/projects/miopen/src/include/miopen/activ.hpp index 26ce91112f4..b7db13b4782 100644 --- a/projects/miopen/src/include/miopen/activ.hpp +++ b/projects/miopen/src/include/miopen/activ.hpp @@ -50,7 +50,7 @@ struct MIOPEN_INTERNALS_EXPORT ActivationDescriptor : miopenActivationDescriptor double GetBeta() const; double GetGamma() const; - miopenStatus_t Forward(Handle& handle, + miopenStatus_t Forward(const Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -60,7 +60,7 @@ struct MIOPEN_INTERNALS_EXPORT ActivationDescriptor : miopenActivationDescriptor size_t xOffset = 0, size_t yOffset = 0) const; - miopenStatus_t Backward(Handle& handle, + miopenStatus_t Backward(const Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t y, diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index 0b70b3a29b4..f088f9680b7 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -451,7 +451,7 @@ struct SolverContainer } template - void ExecutePrimitive(ExecutionContext& ctx, + void ExecutePrimitive(const ExecutionContext& ctx, const Problem& problem, const AlgorithmName& algo, const AnyInvokeParams& invoke_params) const @@ -465,7 +465,8 @@ struct SolverContainer return; } - const auto slns = SearchForSolutions(ctx, problem, 1, invoke_params); + auto ctx_cpy = ctx; + const auto slns = SearchForSolutions(ctx_cpy, problem, 1, invoke_params); if(slns.empty()) MIOPEN_THROW(miopenStatusNotImplemented, "No solver found."); @@ -480,7 +481,7 @@ struct SolverContainer } template - void ExecutePrimitive(Handle& handle, + void ExecutePrimitive(const Handle& handle, const Problem& problem, const AlgorithmName& algo, const AnyInvokeParams& invoke_params) const diff --git a/projects/miopen/src/include/miopen/rnn.hpp b/projects/miopen/src/include/miopen/rnn.hpp index ccbd96563b4..49e31910eea 100644 --- a/projects/miopen/src/include/miopen/rnn.hpp +++ b/projects/miopen/src/include/miopen/rnn.hpp @@ -255,7 +255,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardTraining(Handle& handle, + void RNNForwardTraining(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -276,7 +276,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardInference(Handle& handle, + void RNNForwardInference(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -295,7 +295,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t workSpace, size_t workSpaceSize) const; - void RNNBackwardData(Handle& handle, + void RNNBackwardData(const Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t y, ConstData_t dy, @@ -316,7 +316,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardData(Handle& handle, + void RNNBackwardData(const Handle& handle, int seqLen, c_array_view yDesc, ConstData_t y, @@ -343,7 +343,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeights(Handle& handle, + void RNNBackwardWeights(const Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -357,7 +357,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeights(Handle& handle, + void RNNBackwardWeights(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -430,7 +430,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t /*reserveSpaceSize*/) const; - void RNNTransformerForward(Handle& handle, + void RNNTransformerForward(const Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -448,7 +448,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNTransformerBackwardData(Handle& handle, + void RNNTransformerBackwardData(const Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -467,7 +467,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNTransformerBackwardWeights(Handle& handle, + void RNNTransformerBackwardWeights(const Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -479,7 +479,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNVanillaForward(Handle& handle, + void RNNVanillaForward(const Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -528,7 +528,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor ConstData_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardTrainingPackedTensors(Handle& handle, + void RNNForwardTrainingPackedTensors(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -549,7 +549,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNForwardMS(Handle& handle, + void RNNForwardMS(const Handle& handle, std::vector& seq_array, const TensorDescriptor& xDesc, ConstData_t x, @@ -566,7 +566,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor size_t extra_space_size, miopenRNNFWDMode_t fwd_mode) const; - void RNNForwardInferencePacked(Handle& handle, + void RNNForwardInferencePacked(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, @@ -585,7 +585,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t workSpace, size_t workSpaceSize) const; - void RNNBackwardDataPackedTensors(Handle& handle, + void RNNBackwardDataPackedTensors(const Handle& handle, int seqLen, c_array_view dyDesc, ConstData_t dy, @@ -605,7 +605,7 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor Data_t reserveSpace, size_t reserveSpaceSize) const; - void RNNBackwardWeightsPackedTensors(Handle& handle, + void RNNBackwardWeightsPackedTensors(const Handle& handle, int seqLen, c_array_view xDesc, ConstData_t x, diff --git a/projects/miopen/src/kthvalue.cpp b/projects/miopen/src/kthvalue.cpp index a9f4e730678..02dd9f96185 100644 --- a/projects/miopen/src/kthvalue.cpp +++ b/projects/miopen/src/kthvalue.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t KthvalueForward(Handle& handle, +miopenStatus_t KthvalueForward(const Handle& handle, const TensorDescriptor& inputDesc, ConstData_t input, const TensorDescriptor& outputDesc, diff --git a/projects/miopen/src/layernorm.cpp b/projects/miopen/src/layernorm.cpp index b84824e52fd..7529d1ee7f6 100644 --- a/projects/miopen/src/layernorm.cpp +++ b/projects/miopen/src/layernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t LayerNormForward(Handle& handle, +miopenStatus_t LayerNormForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, @@ -105,7 +105,7 @@ std::size_t GetLayerNormBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t LayerNormBackward(Handle& handle, +miopenStatus_t LayerNormBackward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/multimarginloss.cpp b/projects/miopen/src/multimarginloss.cpp index 5638ef6f704..e3daa240120 100644 --- a/projects/miopen/src/multimarginloss.cpp +++ b/projects/miopen/src/multimarginloss.cpp @@ -55,7 +55,7 @@ std::size_t GetMultiMarginLossForwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t MultiMarginLossForward(Handle& handle, +miopenStatus_t MultiMarginLossForward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& iDesc, diff --git a/projects/miopen/src/ocl/activ_ocl.cpp b/projects/miopen/src/ocl/activ_ocl.cpp index fc301acaddf..185cfd1d3de 100644 --- a/projects/miopen/src/ocl/activ_ocl.cpp +++ b/projects/miopen/src/ocl/activ_ocl.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t ActivationDescriptor::Forward(Handle& handle, +miopenStatus_t ActivationDescriptor::Forward(const Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -74,7 +74,7 @@ miopenStatus_t ActivationDescriptor::Forward(Handle& handle, return miopenStatusSuccess; } -miopenStatus_t ActivationDescriptor::Backward(Handle& handle, +miopenStatus_t ActivationDescriptor::Backward(const Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t y, diff --git a/projects/miopen/src/ocl/pooling_ocl.cpp b/projects/miopen/src/ocl/pooling_ocl.cpp index 9881c1596f4..ab3a63520d8 100644 --- a/projects/miopen/src/ocl/pooling_ocl.cpp +++ b/projects/miopen/src/ocl/pooling_ocl.cpp @@ -54,7 +54,7 @@ static auto PoolingBackwardSolvers() solver::pooling::TransposedPoolingBwdNd>{}; } -miopenStatus_t PoolingDescriptor::Forward(Handle& handle, +miopenStatus_t PoolingDescriptor::Forward(const Handle& handle, const void* alpha, const TensorDescriptor& xDesc, ConstData_t x, @@ -139,7 +139,7 @@ miopenStatus_t PoolingDescriptor::Forward(Handle& handle, return miopenStatusSuccess; } -miopenStatus_t PoolingDescriptor::Backward(Handle& handle, +miopenStatus_t PoolingDescriptor::Backward(const Handle& handle, const void* alpha, const TensorDescriptor& yDesc, ConstData_t /*y*/, diff --git a/projects/miopen/src/ocl/rnnocl.cpp b/projects/miopen/src/ocl/rnnocl.cpp index 8e940505820..5562f7dd389 100644 --- a/projects/miopen/src/ocl/rnnocl.cpp +++ b/projects/miopen/src/ocl/rnnocl.cpp @@ -257,7 +257,7 @@ miopenStatus_t ReducAddBias(const miopen::Handle& handle, } // namespace -void RNNDescriptor::RNNForwardMS(Handle& handle, +void RNNDescriptor::RNNForwardMS(const Handle& handle, std::vector& seq_array, const TensorDescriptor& xDesc, ConstData_t x, @@ -1177,7 +1177,7 @@ void RNNDescriptor::RNNForwardMS(Handle& handle, } // Assuming sequence length is set to > 0 otherwise throw exception. -void RNNDescriptor::RNNForwardInference(Handle& handle, +void RNNDescriptor::RNNForwardInference(const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -1310,7 +1310,7 @@ void RNNDescriptor::RNNForwardInference(Handle& handle, } #endif } -void RNNDescriptor::RNNForwardInferencePacked(Handle& handle, +void RNNDescriptor::RNNForwardInferencePacked(const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -2563,7 +2563,7 @@ void RNNDescriptor::RNNForwardInferencePacked(Handle& handle, #endif } -void RNNDescriptor::RNNForwardTraining(Handle& handle, +void RNNDescriptor::RNNForwardTraining(const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -2708,7 +2708,7 @@ void RNNDescriptor::RNNForwardTraining(Handle& handle, }; void RNNDescriptor::RNNForwardTrainingPackedTensors( - Handle& handle, + const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -4080,7 +4080,7 @@ void RNNDescriptor::RNNForwardTrainingPackedTensors( #endif }; -void RNNDescriptor::RNNBackwardData(Handle& handle, +void RNNDescriptor::RNNBackwardData(const Handle& handle, const int seqLen, c_array_view yDesc, ConstData_t y, @@ -4223,7 +4223,7 @@ void RNNDescriptor::RNNBackwardData(Handle& handle, } void RNNDescriptor::RNNBackwardDataPackedTensors( - Handle& handle, + const Handle& handle, const int seqLen, c_array_view dyDesc, ConstData_t dy, @@ -5810,7 +5810,7 @@ void RNNDescriptor::RNNBackwardDataPackedTensors( #endif }; -void RNNDescriptor::RNNBackwardWeights(Handle& handle, +void RNNDescriptor::RNNBackwardWeights(const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, @@ -5910,7 +5910,7 @@ void RNNDescriptor::RNNBackwardWeights(Handle& handle, } void RNNDescriptor::RNNBackwardWeightsPackedTensors( - Handle& handle, + const Handle& handle, const int seqLen, c_array_view xDesc, ConstData_t x, diff --git a/projects/miopen/src/prelu.cpp b/projects/miopen/src/prelu.cpp index b43e5ea84bd..34ec9d59ab9 100644 --- a/projects/miopen/src/prelu.cpp +++ b/projects/miopen/src/prelu.cpp @@ -48,7 +48,7 @@ size_t GetPReLUBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t PReLUBackward(Handle& handle, +miopenStatus_t PReLUBackward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& inputDesc, diff --git a/projects/miopen/src/reducecalculation.cpp b/projects/miopen/src/reducecalculation.cpp index ea4fda314d7..d7bc58f89d5 100644 --- a/projects/miopen/src/reducecalculation.cpp +++ b/projects/miopen/src/reducecalculation.cpp @@ -71,7 +71,7 @@ std::size_t GetReduceCalculationWorkspaceSize(const Handle& handle, return static_cast(-1); } -miopenStatus_t ReduceCalculationForward(Handle& handle, +miopenStatus_t ReduceCalculationForward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& xDesc, diff --git a/projects/miopen/src/reduceextreme.cpp b/projects/miopen/src/reduceextreme.cpp index a35df325986..37fbfd82dbe 100644 --- a/projects/miopen/src/reduceextreme.cpp +++ b/projects/miopen/src/reduceextreme.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t ReduceExtremeForward(Handle& handle, +miopenStatus_t ReduceExtremeForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& indiceDesc, diff --git a/projects/miopen/src/rnn.cpp b/projects/miopen/src/rnn.cpp index ce7a834d856..83fe4e5abdd 100644 --- a/projects/miopen/src/rnn.cpp +++ b/projects/miopen/src/rnn.cpp @@ -44,7 +44,7 @@ namespace miopen { -void profileRNNkernels(Handle& handle, unsigned char select, float& ctime) +void profileRNNkernels(const Handle& handle, unsigned char select, float& ctime) { float ktime = 0.; @@ -1207,7 +1207,7 @@ void RNNDescriptor::SeqTensorToTensorDescArray(const SeqTensorDescriptor& desc, }); } -void RNNDescriptor::RNNVanillaForward(Handle& handle, +void RNNDescriptor::RNNVanillaForward(const Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -1284,7 +1284,7 @@ void RNNDescriptor::RNNVanillaForward(Handle& handle, } } -void RNNDescriptor::RNNVanillaBackwardData(Handle& handle, +void RNNDescriptor::RNNVanillaBackwardData(const Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -1335,7 +1335,7 @@ void RNNDescriptor::RNNVanillaBackwardData(Handle& handle, reserveSpaceSize); } -void RNNDescriptor::RNNVanillaBackwardWeights(Handle& handle, +void RNNDescriptor::RNNVanillaBackwardWeights(const Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, @@ -1374,7 +1374,7 @@ void RNNDescriptor::RNNVanillaBackwardWeights(Handle& handle, reserveSpaceSize); } -void RNNDescriptor::RNNForward(Handle& handle, +void RNNDescriptor::RNNForward(const Handle& handle, miopenRNNFWDMode_t fwdMode, const SeqTensorDescriptor& xDesc, ConstData_t x, @@ -1476,7 +1476,7 @@ void RNNDescriptor::RNNForward(Handle& handle, #endif } -void RNNDescriptor::RNNBackwardData(Handle& handle, +void RNNDescriptor::RNNBackwardData(const Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t, ConstData_t dy, @@ -1584,7 +1584,7 @@ void RNNDescriptor::RNNBackwardData(Handle& handle, #endif } -void RNNDescriptor::RNNBackwardWeights(Handle& handle, +void RNNDescriptor::RNNBackwardWeights(const Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, diff --git a/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp b/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp index 72ebb560ee6..57c6866167d 100644 --- a/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp +++ b/projects/miopen/src/rnn/Solutions/rnn_transformer.cpp @@ -120,7 +120,7 @@ size_t RNNDescriptor::RNNTransformerWorkspaceSize(const SeqTensorDescriptor& xDe typeSize; } -void RNNDescriptor::RNNTransformerForward(Handle& handle, +void RNNDescriptor::RNNTransformerForward(const Handle& handle, miopenRNNFWDMode_t fwdMode, ConstData_t w, const SeqTensorDescriptor& xDesc, @@ -278,7 +278,7 @@ void RNNDescriptor::RNNTransformerForward(Handle& handle, } } -void RNNDescriptor::RNNTransformerBackwardData(Handle& handle, +void RNNDescriptor::RNNTransformerBackwardData(const Handle& handle, const SeqTensorDescriptor& yDesc, ConstData_t dy, const TensorDescriptor& hDesc, @@ -446,7 +446,7 @@ void RNNDescriptor::RNNTransformerBackwardData(Handle& handle, } } -void RNNDescriptor::RNNTransformerBackwardWeights(Handle& handle, +void RNNDescriptor::RNNTransformerBackwardWeights(const Handle& handle, const SeqTensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& hDesc, diff --git a/projects/miopen/src/rope.cpp b/projects/miopen/src/rope.cpp index 9dca9b145e9..d54f68bbd8a 100644 --- a/projects/miopen/src/rope.cpp +++ b/projects/miopen/src/rope.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t RoPEForward(Handle& handle, +miopenStatus_t RoPEForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& cosDesc, @@ -69,7 +69,7 @@ miopenStatus_t RoPEForward(Handle& handle, return miopenStatusSuccess; } -miopenStatus_t RoPEBackward(Handle& handle, +miopenStatus_t RoPEBackward(const Handle& handle, const TensorDescriptor& dyDesc, ConstData_t dy, const TensorDescriptor& cosDesc, diff --git a/projects/miopen/src/softmarginloss.cpp b/projects/miopen/src/softmarginloss.cpp index 4ce38baa591..a0b8072f76a 100644 --- a/projects/miopen/src/softmarginloss.cpp +++ b/projects/miopen/src/softmarginloss.cpp @@ -50,7 +50,7 @@ std::size_t GetSoftMarginLossForwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(-1) : pair_size_vector.front().second; } -miopenStatus_t SoftMarginLossForward(Handle& handle, +miopenStatus_t SoftMarginLossForward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& iDesc, @@ -84,7 +84,7 @@ miopenStatus_t SoftMarginLossForward(Handle& handle, return miopenStatusSuccess; } -miopenStatus_t SoftMarginLossBackward(Handle& handle, +miopenStatus_t SoftMarginLossBackward(const Handle& handle, const TensorDescriptor& iDesc, ConstData_t i, const TensorDescriptor& tDesc, diff --git a/projects/miopen/src/t5layernorm.cpp b/projects/miopen/src/t5layernorm.cpp index dacd51a208f..08b2600c53a 100644 --- a/projects/miopen/src/t5layernorm.cpp +++ b/projects/miopen/src/t5layernorm.cpp @@ -34,7 +34,7 @@ namespace miopen { -miopenStatus_t T5LayerNormForward(Handle& handle, +miopenStatus_t T5LayerNormForward(const Handle& handle, const TensorDescriptor& xDesc, ConstData_t x, const TensorDescriptor& weightDesc, @@ -91,7 +91,7 @@ std::size_t GetT5LayerNormBackwardWorkspaceSize(const Handle& handle, return pair_size_vector.empty() ? static_cast(0) : pair_size_vector.front().second; } -miopenStatus_t T5LayerNormBackward(Handle& handle, +miopenStatus_t T5LayerNormBackward(const Handle& handle, Data_t workspace, size_t workspaceSizeInBytes, const TensorDescriptor& dyDesc, diff --git a/projects/miopen/src/transformers_adam_w.cpp b/projects/miopen/src/transformers_adam_w.cpp index 0e644d8344e..3b5253c9403 100644 --- a/projects/miopen/src/transformers_adam_w.cpp +++ b/projects/miopen/src/transformers_adam_w.cpp @@ -35,7 +35,7 @@ namespace miopen { -miopenStatus_t TransformersAdamW(Handle& handle, +miopenStatus_t TransformersAdamW(const Handle& handle, const TensorDescriptor& paramInDesc, ConstData_t paramIn, const TensorDescriptor& paramOutDesc, From f8e3c3094df40a44186e9ee21d6e9f4caa6ea108 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 9 Sep 2025 01:03:18 +0000 Subject: [PATCH 05/26] whitespace --- projects/miopen/src/include/miopen/find_solution.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index f088f9680b7..ca914cf1c5d 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -465,7 +465,7 @@ struct SolverContainer return; } - auto ctx_cpy = ctx; + auto ctx_cpy = ctx; const auto slns = SearchForSolutions(ctx_cpy, problem, 1, invoke_params); if(slns.empty()) From bbba8f6bc7d74bc3dcb7c8c041447cd54498e29d Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 9 Sep 2025 20:18:05 +0000 Subject: [PATCH 06/26] adjust function signatures, corrections for cutoff updates --- projects/miopen/src/fusion.cpp | 2 +- .../include/miopen/conv/solver_finders.hpp | 2 +- .../src/include/miopen/find_controls.hpp | 6 ++-- .../src/include/miopen/generic_search.hpp | 28 ++++++++++++------- projects/miopen/src/mlo_dir_conv.cpp | 4 +-- 5 files changed, 26 insertions(+), 16 deletions(-) diff --git a/projects/miopen/src/fusion.cpp b/projects/miopen/src/fusion.cpp index 284580e747e..fdc4da6f308 100644 --- a/projects/miopen/src/fusion.cpp +++ b/projects/miopen/src/fusion.cpp @@ -867,7 +867,7 @@ static const std::vector>& GetFusionSolverFinder } static std::vector -FindFusion(const ExecutionContext& ctx, +FindFusion(ExecutionContext& ctx, const FusionDescription& fusion_problem, const std::function& invoke_params, const std::optional& options = std::nullopt) diff --git a/projects/miopen/src/include/miopen/conv/solver_finders.hpp b/projects/miopen/src/include/miopen/conv/solver_finders.hpp index 8600cfbc60f..0945b18954a 100644 --- a/projects/miopen/src/include/miopen/conv/solver_finders.hpp +++ b/projects/miopen/src/include/miopen/conv/solver_finders.hpp @@ -172,7 +172,7 @@ std::vector EvaluateInvokers(const Handle& handle, bool force_attach_binary); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, - const ExecutionContext& ctx, + ExecutionContext& ctx, const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters, const std::vector>& finders, diff --git a/projects/miopen/src/include/miopen/find_controls.hpp b/projects/miopen/src/include/miopen/find_controls.hpp index 4f27bb60f66..dcffea0feca 100644 --- a/projects/miopen/src/include/miopen/find_controls.hpp +++ b/projects/miopen/src/include/miopen/find_controls.hpp @@ -88,7 +88,8 @@ class MIOPEN_INTERNALS_EXPORT FindEnforce { return IsEnabled(context) && (action == FindEnforceAction::Search || action == FindEnforceAction::SearchDbUpdate || - action == FindEnforceAction::DbUpdate); + action == FindEnforceAction::DbUpdate || + action == FindEnforceAction::SearchCutoffDbUpdate); } template @@ -101,7 +102,8 @@ class MIOPEN_INTERNALS_EXPORT FindEnforce bool IsDbUpdate(const Context& context) const { return IsEnabled(context) && (action == FindEnforceAction::DbUpdate || - action == FindEnforceAction::SearchDbUpdate); + action == FindEnforceAction::SearchDbUpdate || + action == FindEnforceAction::SearchCutoffDbUpdate); } template diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index fde433c0afa..d43e93b9963 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -482,6 +482,7 @@ auto GenericSearch(const Solver s, float worst_time = std::numeric_limits::max(); size_t n_failed = 0; size_t n_best = 0; + bool perf_cutoff = false; HeartBeat heartbeat; heartbeat.Start(); @@ -592,14 +593,17 @@ auto GenericSearch(const Solver s, if(ret == 0) { - // If 1st probe of 1st successful config is worse than the cutoff time abort the - // search - if(perf_sols.empty() && elapsed_time > context.generic_search_cutoff_time) + // If config is worse than the cutoff time abort the search + if(elapsed_time > context.generic_search_cutoff_time) { - MIOPEN_LOG_E("Measured time: " << elapsed_time << " was greater than cutoff: " - << context.generic_search_cutoff_time - << " aborting search..."); - return current_config; + MIOPEN_LOG_I2("Ending Search, measured time: " << elapsed_time << " was greater than cutoff: " + << context.generic_search_cutoff_time); + perf_cutoff = true; + for(const auto& kernelInfo : current_solution.construction_params) + profile_h.ClearProgram(kernelInfo.kernel_file, kernelInfo.comp_options); + if(perf_sols.empty()) + best_config = current_config; + break; } // Smooth the jitter of measurements: @@ -700,11 +704,15 @@ auto GenericSearch(const Solver s, }); // if using cutoff time for search and new cutoff is shorter, update - if(context.search_cutoff) + if(context.search_cutoff && !perf_cutoff) { - float new_cutoff = perf_sols.end()->time * 2; - if(new_cutoff < context.generic_search_cutoff_time) + // cutoff based on worst performing config + float new_cutoff = (perf_sols.end()-1)->time * 1.1f; + if(new_cutoff > 0.0f && new_cutoff < context.generic_search_cutoff_time) + { context_.generic_search_cutoff_time = new_cutoff; + MIOPEN_LOG_I2("Cutoff time updated: " << new_cutoff); + } } if(perf_solsp) diff --git a/projects/miopen/src/mlo_dir_conv.cpp b/projects/miopen/src/mlo_dir_conv.cpp index 78028e75b4a..9843b4ab4ff 100644 --- a/projects/miopen/src/mlo_dir_conv.cpp +++ b/projects/miopen/src/mlo_dir_conv.cpp @@ -252,7 +252,7 @@ AllDirectForwardBackwardDataWorkspaceSize(const miopen::ExecutionContext& ctx, } std::vector> -FindAllWinogradWorkspaceSizes(miopen::ExecutionContext& ctx, +FindAllWinogradWorkspaceSizes(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem) { return GetWindogradSolvers().GetWorkspaceSizes(ctx, problem); @@ -266,7 +266,7 @@ FindWinogradWrWWorkspaceSizes(const miopen::ExecutionContext& ctx, } std::vector> -FindAllImplicitGemmWorkspaceSizes(miopen::ExecutionContext& ctx, +FindAllImplicitGemmWorkspaceSizes(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem) { return GetImplicitGemmSolvers().GetWorkspaceSizes(ctx, problem); From 233cdcc84584634967536afad8dccc7650c654e5 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 9 Sep 2025 20:23:21 +0000 Subject: [PATCH 07/26] clang format --- projects/miopen/src/conv/solver_finders.cpp | 4 ++-- projects/miopen/src/fusion.cpp | 2 +- projects/miopen/src/include/miopen/find_solution.hpp | 2 +- projects/miopen/src/include/miopen/generic_search.hpp | 11 ++++++----- .../solver/conv/conv_MP_bidirectional_winograd.cpp | 2 +- 5 files changed, 11 insertions(+), 10 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 2638d09d48f..87d5aeb8c45 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -202,10 +202,10 @@ const std::vector>& GetConvSolverFinders() { static const auto finders = []() { auto tmp = std::vector>{}; - tmp.emplace_back(std::make_unique()); - tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); + tmp.emplace_back(std::make_unique()); + tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); return tmp; }(); diff --git a/projects/miopen/src/fusion.cpp b/projects/miopen/src/fusion.cpp index fdc4da6f308..b1539e50bbd 100644 --- a/projects/miopen/src/fusion.cpp +++ b/projects/miopen/src/fusion.cpp @@ -1046,7 +1046,7 @@ miopenStatus_t FusionPlanDescriptor::Compile(const Handle& handle) const auto id = solver::Id{sol->solution_id}; GetAllFusionSolvers().FindById(id, [&](auto solver) { - auto ctx = FusionContext{handle}; + auto ctx = FusionContext{handle}; auto db_getter = MakeConvDbGetter(ctx); const auto solution = solver::FindSolution( solver, ctx, fusion_problem, db_getter, {}); // auto tune is not expected here diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index ca914cf1c5d..c4b89d8745e 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -465,7 +465,7 @@ struct SolverContainer return; } - auto ctx_cpy = ctx; + auto ctx_cpy = ctx; const auto slns = SearchForSolutions(ctx_cpy, problem, 1, invoke_params); if(slns.empty()) diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index d43e93b9963..2a075dbef96 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -596,8 +596,9 @@ auto GenericSearch(const Solver s, // If config is worse than the cutoff time abort the search if(elapsed_time > context.generic_search_cutoff_time) { - MIOPEN_LOG_I2("Ending Search, measured time: " << elapsed_time << " was greater than cutoff: " - << context.generic_search_cutoff_time); + MIOPEN_LOG_I2("Ending Search, measured time: " + << elapsed_time << " was greater than cutoff: " + << context.generic_search_cutoff_time); perf_cutoff = true; for(const auto& kernelInfo : current_solution.construction_params) profile_h.ClearProgram(kernelInfo.kernel_file, kernelInfo.comp_options); @@ -707,12 +708,12 @@ auto GenericSearch(const Solver s, if(context.search_cutoff && !perf_cutoff) { // cutoff based on worst performing config - float new_cutoff = (perf_sols.end()-1)->time * 1.1f; + float new_cutoff = (perf_sols.end() - 1)->time * 1.1f; if(new_cutoff > 0.0f && new_cutoff < context.generic_search_cutoff_time) - { + { context_.generic_search_cutoff_time = new_cutoff; MIOPEN_LOG_I2("Cutoff time updated: " << new_cutoff); - } + } } if(perf_solsp) diff --git a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp index 90067c24639..3cab1911f40 100644 --- a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp +++ b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp @@ -948,7 +948,7 @@ ConvMPBidirectWinograd_xdlops::S GetTransformedInvokeContext(problem, invoke_ctx); const auto xdlops_problem = GetTransformedProblem(problem); - auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); + auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); return ConvHipImplicitGemmForwardV4R4Xdlops().Search( xdlops_ctx, xdlops_problem, xdlops_invoke_ctx); From 13523647d3945cd27c508ecccc9b52d566db955c Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 10 Sep 2025 18:06:48 +0000 Subject: [PATCH 08/26] split cutoff time to best_time and worst_time --- .../src/include/miopen/execution_context.hpp | 3 +- .../src/include/miopen/generic_search.hpp | 32 +++++++++---------- 2 files changed, 18 insertions(+), 17 deletions(-) diff --git a/projects/miopen/src/include/miopen/execution_context.hpp b/projects/miopen/src/include/miopen/execution_context.hpp index 80c6d0829a6..b6a9c41849d 100644 --- a/projects/miopen/src/include/miopen/execution_context.hpp +++ b/projects/miopen/src/include/miopen/execution_context.hpp @@ -98,7 +98,8 @@ struct MIOPEN_INTERNALS_EXPORT ExecutionContext bool disable_perfdb_access = false; bool use_dynamic_solutions_only = false; bool is_for_generic_search = false; - float generic_search_cutoff_time = std::numeric_limits::max(); + float generic_search_worst_time = std::numeric_limits::max(); + float generic_search_best_time = std::numeric_limits::max(); inline const Handle& GetStream() const { return *stream; } inline void SetStream(const Handle* stream_) { stream = stream_; } diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 2a075dbef96..fa8d76d8704 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -482,7 +482,13 @@ auto GenericSearch(const Solver s, float worst_time = std::numeric_limits::max(); size_t n_failed = 0; size_t n_best = 0; - bool perf_cutoff = false; + float cutoff_time = context.generic_search_worst_time; + if(cutoff_time < std::numeric_limits::max()) + cutoff_time *= 10; + float overall_best_time = context.generic_search_best_time; + if(overall_best_time < std::numeric_limits::max()) + overall_best_time *= 1.30f; + HeartBeat heartbeat; heartbeat.Start(); @@ -594,16 +600,13 @@ auto GenericSearch(const Solver s, if(ret == 0) { // If config is worse than the cutoff time abort the search - if(elapsed_time > context.generic_search_cutoff_time) + if(elapsed_time > cutoff_time) { MIOPEN_LOG_I2("Ending Search, measured time: " << elapsed_time << " was greater than cutoff: " - << context.generic_search_cutoff_time); - perf_cutoff = true; + << cutoff_time); for(const auto& kernelInfo : current_solution.construction_params) profile_h.ClearProgram(kernelInfo.kernel_file, kernelInfo.comp_options); - if(perf_sols.empty()) - best_config = current_config; break; } @@ -613,7 +616,7 @@ auto GenericSearch(const Solver s, // the mean value with outliers removed for calculating best config. constexpr int N_RUNS = 10; last_imprv++; - if(elapsed_time / worst_time < 1.10f) + if(elapsed_time < worst_time * 1.10f && elapsed_time < overall_best_time) { MIOPEN_LOG_I2("Finding average for: " << elapsed_time << " / " << best_time << " = " << (elapsed_time / best_time)); @@ -704,16 +707,13 @@ auto GenericSearch(const Solver s, return a.time < b.time; }); - // if using cutoff time for search and new cutoff is shorter, update - if(context.search_cutoff && !perf_cutoff) + // if using cutoff for search update timing + if(context.search_cutoff && best_time < context.generic_search_best_time) { - // cutoff based on worst performing config - float new_cutoff = (perf_sols.end() - 1)->time * 1.1f; - if(new_cutoff > 0.0f && new_cutoff < context.generic_search_cutoff_time) - { - context_.generic_search_cutoff_time = new_cutoff; - MIOPEN_LOG_I2("Cutoff time updated: " << new_cutoff); - } + float new_worst = (perf_sols.end() - 1)->time; + context_.generic_search_best_time = best_time; + context_.generic_search_worst_time = new_worst; + MIOPEN_LOG_I2("Times updated, best: " << best_time << " worst: " << new_worst); } if(perf_solsp) From 2362fd67d2efa00f77ea36db81c7837e328b97e5 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 11 Sep 2025 22:14:21 +0000 Subject: [PATCH 09/26] restore const to ExecutionContext, add mutable to search time fields --- projects/miopen/src/conv/solver_finders.cpp | 12 +-- projects/miopen/src/fusion.cpp | 10 +-- .../miopen/src/include/miopen/any_solver.hpp | 32 ++++---- .../src/include/miopen/batchnorm/solvers.hpp | 4 +- .../include/miopen/conv/solver_finders.hpp | 10 +-- .../src/include/miopen/conv/solvers.hpp | 82 +++++++++---------- .../src/include/miopen/execution_context.hpp | 6 +- .../src/include/miopen/find_solution.hpp | 13 ++- .../src/include/miopen/fusion/solvers.hpp | 12 +-- .../src/include/miopen/generic_search.hpp | 2 +- .../src/include/miopen/mlo_internal.hpp | 16 ++-- projects/miopen/src/include/miopen/solver.hpp | 6 +- projects/miopen/src/mlo_dir_conv.cpp | 16 ++-- projects/miopen/src/ocl/convolutionocl.cpp | 44 ++++++---- .../src/solver/batchnorm/backward_spatial.cpp | 2 +- .../src/solver/batchnorm/forward_spatial.cpp | 2 +- .../conv/conv_MP_bidirectional_winograd.cpp | 4 +- .../miopen/src/solver/conv/conv_asm_1x1u.cpp | 2 +- .../src/solver/conv/conv_asm_1x1u_stride2.cpp | 2 +- .../miopen/src/solver/conv/conv_asm_3x3u.cpp | 2 +- .../solver/conv/conv_asm_dir_BwdWrW1x1.cpp | 2 +- .../solver/conv/conv_asm_dir_BwdWrW3x3.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp | 2 +- .../conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 2 +- .../conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_bwd_xdlops.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_fwd_xdlops.cpp | 2 +- ...ip_implicit_gemm_3d_grouped_wrw_xdlops.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_data_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_bwd_v1r1.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_bwd_v4r1.cpp | 2 +- ...conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_fwd_v4r1.cpp | 4 +- .../conv/conv_hip_implicit_gemm_fwd_v4r4.cpp | 2 +- ...conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp | 2 +- ...licit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp | 2 +- ...conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp | 2 +- .../conv_hip_implicit_gemm_fwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_bwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_fwd_xdlops.cpp | 2 +- ...v_hip_implicit_gemm_grouped_wrw_xdlops.cpp | 2 +- .../conv/conv_hip_implicit_gemm_wrw_v4r4.cpp | 2 +- ...conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp | 2 +- ...licit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_bwd.cpp | 2 +- .../conv/conv_mlir_igemm_bwd_xdlops.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_fwd.cpp | 2 +- .../conv/conv_mlir_igemm_fwd_xdlops.cpp | 2 +- .../src/solver/conv/conv_mlir_igemm_wrw.cpp | 2 +- .../conv/conv_mlir_igemm_wrw_xdlops.cpp | 2 +- .../solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp | 2 +- .../conv_ocl_dir2Dfwd_exhaustive_search.cpp | 2 +- .../miopen/src/solver/conv/conv_winoRxS.cpp | 2 +- .../solver/conv_asm_1x1u_bias_activ_fused.cpp | 2 +- .../conv_ck_igemm_fwd_bias_activ_fused.cpp | 2 +- ..._ck_igemm_fwd_bias_res_add_activ_fused.cpp | 2 +- .../conv_ck_igemm_grp_fwd_activ_fused.cpp | 2 +- ...conv_ck_igemm_grp_fwd_bias_activ_fused.cpp | 2 +- .../src/solver/conv_ocl_dir2Dfwd_fused.cpp | 4 +- 62 files changed, 187 insertions(+), 180 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 87d5aeb8c45..2c008d4e8f8 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -66,7 +66,7 @@ class DirectSolverFinder : public SolversFinderMixin FindImpl(ExecutionContext& ctx, + std::vector FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -95,7 +95,7 @@ class ImplicitGemmSolverFinder : public SolversFinderMixin FindImpl(ExecutionContext& ctx, + std::vector FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -126,7 +126,7 @@ class FftSolverFinder : public SolversFinderMixin FindImpl(ExecutionContext& ctx, + std::vector FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -153,7 +153,7 @@ class GemmSolverFinder : public SolversFinderMixin FindImpl(ExecutionContext& ctx, + std::vector FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters&, @@ -180,7 +180,7 @@ class WinogradSolverFinder : public SolversFinderMixin FindImpl(ExecutionContext& ctx, + std::vector FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const ConvFindParameters& parameters, @@ -334,7 +334,7 @@ std::vector EvaluateInvokers(const Handle& handle, } FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, - ExecutionContext& ctx, + const ExecutionContext& ctx, const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters, const std::vector>& finders, diff --git a/projects/miopen/src/fusion.cpp b/projects/miopen/src/fusion.cpp index b1539e50bbd..7694e2e59fc 100644 --- a/projects/miopen/src/fusion.cpp +++ b/projects/miopen/src/fusion.cpp @@ -787,7 +787,7 @@ static auto GetAllFusionSolvers() GetFusedWinogradSolvers(); } -solver::ConvSolution MakeFusedSolution(FusionContext& ctx, +solver::ConvSolution MakeFusedSolution(const FusionContext& ctx, solver::Id id, const std::optional& perf_cfg_override, const FusionDescription& problem, @@ -829,13 +829,13 @@ class FusionSolverFinder : public SolversFinderMixin - FindImpl(ExecutionContext& ctx, + FindImpl(const ExecutionContext& ctx, const FusionDescription& problem, const AnyInvokeParams& invoke_ctx, const FusionFindParameters&, const std::optional& options) const override { - auto fusion_ctx = FusionContext(ctx); + const auto fusion_ctx = FusionContext(ctx); return solvers.SearchForAllSolutions(fusion_ctx, problem, MakeConvDbGetter(ctx), @@ -867,7 +867,7 @@ static const std::vector>& GetFusionSolverFinder } static std::vector -FindFusion(ExecutionContext& ctx, +FindFusion(const ExecutionContext& ctx, const FusionDescription& fusion_problem, const std::function& invoke_params, const std::optional& options = std::nullopt) @@ -1046,7 +1046,7 @@ miopenStatus_t FusionPlanDescriptor::Compile(const Handle& handle) const auto id = solver::Id{sol->solution_id}; GetAllFusionSolvers().FindById(id, [&](auto solver) { - auto ctx = FusionContext{handle}; + const auto ctx = FusionContext{handle}; auto db_getter = MakeConvDbGetter(ctx); const auto solution = solver::FindSolution( solver, ctx, fusion_problem, db_getter, {}); // auto tune is not expected here diff --git a/projects/miopen/src/include/miopen/any_solver.hpp b/projects/miopen/src/include/miopen/any_solver.hpp index 636fa9e9adc..5de823cc133 100644 --- a/projects/miopen/src/include/miopen/any_solver.hpp +++ b/projects/miopen/src/include/miopen/any_solver.hpp @@ -98,7 +98,7 @@ struct AnySolver bool IsEmpty() const { return ptr_value == nullptr; }; #if defined(FIN_ANY_SOLVER_FIND_SOLUTION_COMPAT) - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, @@ -109,7 +109,7 @@ struct AnySolver } #endif - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -119,7 +119,7 @@ struct AnySolver return ptr_value->FindSolution(ctx, problem, db_getter, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -129,7 +129,7 @@ struct AnySolver return ptr_value->FindSolution(ctx, problem, db_getter, invoke_ctx, perf_cfg); } - std::string GenericSearch(ExecutionContext& ctx, + std::string GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp = nullptr) const @@ -194,23 +194,23 @@ struct AnySolver const miopen::conv::ProblemDescription& problem) const = 0; virtual const std::type_info& Type() const = 0; virtual std::string GetSolverDbId() const = 0; - virtual ConvSolution FindSolution(ExecutionContext& ctx, + virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; - virtual ConvSolution FindSolution(ExecutionContext& ctx, + virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; - virtual ConvSolution FindSolution(ExecutionContext& ctx, + virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, const std::string& perf_cfg) const = 0; virtual std::string - GenericSearch(ExecutionContext& ctx, + GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp) const = 0; @@ -413,7 +413,7 @@ struct AnySolver return value.GetWti(ctx, problem); } - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, @@ -422,7 +422,7 @@ struct AnySolver return miopen::solver::FindSolution(value, ctx, problem, db, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -432,7 +432,7 @@ struct AnySolver value, ctx, problem, db_getter, invoke_ctx, perf_cfg); } - ConvSolution FindSolution(ExecutionContext& ctx, + ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, @@ -443,7 +443,7 @@ struct AnySolver } // tunable legacy solver - std::string GenericSearch(ExecutionContext&, + std::string GenericSearch(const ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -454,7 +454,7 @@ struct AnySolver } // tunable solver - std::string GenericSearch(ExecutionContext& ctx, + std::string GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp, @@ -467,7 +467,7 @@ struct AnySolver } // non-tunable solver has no search - std::string GenericSearch(ExecutionContext&, + std::string GenericSearch(const ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -477,7 +477,7 @@ struct AnySolver MIOPEN_THROW("No GenericSearch for non-tunable Solvers."); } - std::string GenericSearch(ExecutionContext&, + std::string GenericSearch(const ExecutionContext&, const miopen::conv::ProblemDescription&, const miopen::AnyInvokeParams&, std::vector*, @@ -488,7 +488,7 @@ struct AnySolver } std::string - GenericSearch(ExecutionContext& ctx, + GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx, std::vector* perf_solsp) const override diff --git a/projects/miopen/src/include/miopen/batchnorm/solvers.hpp b/projects/miopen/src/include/miopen/batchnorm/solvers.hpp index e1d6b065429..07c358efba3 100644 --- a/projects/miopen/src/include/miopen/batchnorm/solvers.hpp +++ b/projects/miopen/src/include/miopen/batchnorm/solvers.hpp @@ -130,7 +130,7 @@ struct BnBwdTrainingSpatial final : BatchNormTunableSolver - Find(ExecutionContext& ctx, + Find(const ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -93,7 +93,7 @@ class ISolversFinder const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters) const = 0; [[nodiscard]] virtual std::vector - FindImpl(ExecutionContext& ctx, + FindImpl(const ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -113,7 +113,7 @@ class SolversFinderMixin : public ISolversFinder } [[nodiscard]] std::vector - FindImpl(ExecutionContext& ctx, + FindImpl(const ExecutionContext& ctx, const ProblemDescriptionBase& problem, const AnyInvokeParams& invoke_ctx, const PrimitiveFindParameters& parameters, @@ -140,7 +140,7 @@ class SolversFinderMixin : public ISolversFinder GetAlgorithmName(const ProblemDescription& problem) const = 0; [[nodiscard]] virtual std::vector - FindImpl(ExecutionContext& ctx, + FindImpl(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, const FindParameters& parameters, @@ -172,7 +172,7 @@ std::vector EvaluateInvokers(const Handle& handle, bool force_attach_binary); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, - ExecutionContext& ctx, + const ExecutionContext& ctx, const ProblemDescriptionBase& problem, const PrimitiveFindParameters& parameters, const std::vector>& finders, diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index 52677c4e83d..5fa56cb8da1 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/projects/miopen/src/include/miopen/conv/solvers.hpp @@ -123,7 +123,7 @@ struct ConvAsm3x3U final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm3x3U&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm3x3U - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -229,7 +229,7 @@ struct ConvAsm1x1U final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm1x1U&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm1x1U - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -320,7 +320,7 @@ struct ConvAsm1x1UV2 final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConfigConvAsm1x1UV2&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvAsm1x1UV2 - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -822,7 +822,7 @@ struct ConvHipImplicitGemmV4R1Fwd final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1001,7 +1001,7 @@ struct ConvMlirIgemmFwdXdlops final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1078,7 +1078,7 @@ struct ConvMlirIgemmWrWXdlops final : ConvTunableSolver const miopen::conv::ProblemDescription&, const PerformanceConvMlirIgemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConvMlirIgemm - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT ConvSolution @@ -1516,7 +1516,7 @@ struct ConvMlirIgemmBwdXdlops final : ConvTunableSolver&) const override; MIOPEN_INTERNALS_EXPORT PerformanceConfigConvOclBwdWrw2 - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; MIOPEN_INTERNALS_EXPORT bool @@ -2664,7 +2664,7 @@ struct ConvHipImplicitGemmWrwV4R4Xdlops final const miopen::conv::ProblemDescription&, const PerformanceImplicitGemmWrwV4R4Xdlops&) const override; MIOPEN_INTERNALS_EXPORT PerformanceImplicitGemmWrwV4R4Xdlops - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; }; @@ -2761,7 +2761,7 @@ struct ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm final const miopen::conv::ProblemDescription&, const PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm&) const override; MIOPEN_INTERNALS_EXPORT PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm - Search(ExecutionContext&, + Search(const ExecutionContext&, const miopen::conv::ProblemDescription&, const AnyInvokeParams& invoke_ctx) const override; }; @@ -2814,7 +2814,7 @@ struct ConvCkIgemmFwdV6r1DlopsNchw final : ConvTunableSolver::max(); - float generic_search_best_time = std::numeric_limits::max(); + mutable bool search_cutoff = false; + mutable float generic_search_worst_time = std::numeric_limits::max(); + mutable float generic_search_best_time = std::numeric_limits::max(); inline const Handle& GetStream() const { return *stream; } inline void SetStream(const Handle* stream_) { stream = stream_; } diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index c4b89d8745e..4fa58687527 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -52,7 +52,7 @@ namespace solver { template auto FindSolutionImpl(rank<1>, const Solver& s, - Context& context, + const Context& context, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, @@ -161,7 +161,7 @@ auto FindSolutionImpl(rank<1>, template auto FindSolutionImpl(rank<0>, const Solver& s, - Context& context, + const Context& context, const Problem& problem, Db&&, const AnyInvokeParams&, @@ -213,7 +213,7 @@ auto GetInvokeFactoryImpl( /// May read/write perfDb. template ConvSolution FindSolution(const Solver& s, - Context& context, + const Context& context, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, @@ -285,7 +285,7 @@ struct SolverContainer // Search for all applicable solutions among many solvers template std::vector - SearchForAllSolutions(Context& ctx, + SearchForAllSolutions(const Context& ctx, const Problem& problem, Db&& db, const AnyInvokeParams& invoke_ctx, @@ -343,7 +343,7 @@ struct SolverContainer // Search for all applicable solutions among many solvers template std::vector - SearchForSolutions(ExecutionContext& ctx, + SearchForSolutions(const ExecutionContext& ctx, const Problem& problem, std::size_t limit = std::numeric_limits::max(), const AnyInvokeParams& invoke_params = {}) const @@ -465,8 +465,7 @@ struct SolverContainer return; } - auto ctx_cpy = ctx; - const auto slns = SearchForSolutions(ctx_cpy, problem, 1, invoke_params); + const auto slns = SearchForSolutions(ctx, problem, 1, invoke_params); if(slns.empty()) MIOPEN_THROW(miopenStatusNotImplemented, "No solver found."); diff --git a/projects/miopen/src/include/miopen/fusion/solvers.hpp b/projects/miopen/src/include/miopen/fusion/solvers.hpp index 0363befe93e..ec472b455a9 100644 --- a/projects/miopen/src/include/miopen/fusion/solvers.hpp +++ b/projects/miopen/src/include/miopen/fusion/solvers.hpp @@ -76,7 +76,7 @@ struct ConvBiasActivAsm1x1U : FusionTunableSolver auto GenericSearch(const Solver s, - Context& context_, + const Context& context_, const Problem& problem, const AnyInvokeParams& invoke_ctx_, std::vector* perf_solsp = nullptr) diff --git a/projects/miopen/src/include/miopen/mlo_internal.hpp b/projects/miopen/src/include/miopen/mlo_internal.hpp index 63e2f3d4e31..fcf64a00764 100644 --- a/projects/miopen/src/include/miopen/mlo_internal.hpp +++ b/projects/miopen/src/include/miopen/mlo_internal.hpp @@ -174,7 +174,7 @@ auto mloConstruct(T& x) -> decltype(x.mloConstruct(), void()) } MIOPEN_INTERNALS_EXPORT std::vector -FindAllGemmSolutions(miopen::ExecutionContext& ctx, +FindAllGemmSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); @@ -211,37 +211,37 @@ AllFFTForwardBackwardDataWorkspaceSize(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem); MIOPEN_INTERNALS_EXPORT std::vector -FindAllDirectSolutions(miopen::ExecutionContext& ctx, +FindAllDirectSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllImplicitGemmSolutions(miopen::ExecutionContext& ctx, +FindAllImplicitGemmSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllWinogradSolutions(miopen::ExecutionContext& ctx, +FindAllWinogradSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindWinogradWrWAllSolutions(miopen::ExecutionContext& ctx, +FindWinogradWrWAllSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindImplicitGemmWrWAllSolutions(miopen::ExecutionContext& ctx, +FindImplicitGemmWrWAllSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllBwdWrW2DSolutions(miopen::ExecutionContext& ctx, +FindAllBwdWrW2DSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); MIOPEN_INTERNALS_EXPORT std::vector -FindAllFFTSolutions(miopen::ExecutionContext& ctx, +FindAllFFTSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx); diff --git a/projects/miopen/src/include/miopen/solver.hpp b/projects/miopen/src/include/miopen/solver.hpp index ffa65a08d3b..dd812b39946 100644 --- a/projects/miopen/src/include/miopen/solver.hpp +++ b/projects/miopen/src/include/miopen/solver.hpp @@ -162,7 +162,7 @@ struct SolverInterfaceTunable : SolverInterface { /// This function is a simplified version of FindSolution(), it does not obey search parameters /// from the Context and does not use the database. Intended to be used in unit tests. - virtual ConvSolution FindSolutionSimple(Context& ctx, + virtual ConvSolution FindSolutionSimple(const Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; }; @@ -206,14 +206,14 @@ struct SolverBaseTunable : SolverInterfaceTunable, TunableSolv /// Search virtual PerformanceConfig - Search(Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; + Search(const Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const = 0; /// Tunable solvers provide a GetSolution that takes a Context and PerformanceConfig virtual ConvSolution GetSolution(const Context& ctx, const Problem& problem, const PerformanceConfig& config) const = 0; - ConvSolution FindSolutionSimple(Context& ctx, + ConvSolution FindSolutionSimple(const Context& ctx, const Problem& problem, const AnyInvokeParams& invoke_ctx) const final { diff --git a/projects/miopen/src/mlo_dir_conv.cpp b/projects/miopen/src/mlo_dir_conv.cpp index 9843b4ab4ff..7ca8c342d2c 100644 --- a/projects/miopen/src/mlo_dir_conv.cpp +++ b/projects/miopen/src/mlo_dir_conv.cpp @@ -220,7 +220,7 @@ auto miopen::MakeConvDbGetter(const ExecutionContext& ctx) -> DbGetter } std::vector -FindAllGemmSolutions(miopen::ExecutionContext& ctx, +FindAllGemmSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -236,7 +236,7 @@ AllGemmWorkspaceSize(const miopen::ExecutionContext& ctx, } std::vector -FindAllDirectSolutions(miopen::ExecutionContext& ctx, +FindAllDirectSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -273,7 +273,7 @@ FindAllImplicitGemmWorkspaceSizes(const miopen::ExecutionContext& ctx, } std::vector -FindAllImplicitGemmSolutions(miopen::ExecutionContext& ctx, +FindAllImplicitGemmSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -282,7 +282,7 @@ FindAllImplicitGemmSolutions(miopen::ExecutionContext& ctx, } std::vector -FindAllWinogradSolutions(miopen::ExecutionContext& ctx, +FindAllWinogradSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -291,7 +291,7 @@ FindAllWinogradSolutions(miopen::ExecutionContext& ctx, } std::vector -FindWinogradWrWAllSolutions(miopen::ExecutionContext& ctx, +FindWinogradWrWAllSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -314,7 +314,7 @@ FindImplicitGemmWrWWorkspaceSizes(const miopen::ExecutionContext& ctx, } std::vector -FindImplicitGemmWrWAllSolutions(miopen::ExecutionContext& ctx, +FindImplicitGemmWrWAllSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -323,7 +323,7 @@ FindImplicitGemmWrWAllSolutions(miopen::ExecutionContext& ctx, } std::vector -FindAllBwdWrW2DSolutions(miopen::ExecutionContext& ctx, +FindAllBwdWrW2DSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { @@ -332,7 +332,7 @@ FindAllBwdWrW2DSolutions(miopen::ExecutionContext& ctx, } std::vector -FindAllFFTSolutions(miopen::ExecutionContext& ctx, +FindAllFFTSolutions(const miopen::ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, const miopen::AnyInvokeParams& invoke_ctx) { diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index fe3598ee94e..51b9375b8b3 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -255,7 +255,7 @@ static void ShrinkToFind10Results(std::vector& found) MIOPEN_INTERNALS_EXPORT std::vector -GetConvSolutions(ExecutionContext& ctx, +GetConvSolutions(const ExecutionContext& ctx, const conv::ProblemDescription& problem, const std::vector solutions) { @@ -352,7 +352,7 @@ bool HasGoodSolution(const std::vector solutions, return good_entry; } -std::vector VerifiedFDBSolution(ExecutionContext& ctx, +std::vector VerifiedFDBSolution(const ExecutionContext& ctx, const conv::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, bool force_attach_binary, @@ -362,9 +362,10 @@ std::vector VerifiedFDBSolution(ExecutionContext& ctx, const auto& conv = problem.GetConv(); const auto& findMode = conv.findMode; auto results = UserFindDbRecord::TryLoad(ctx.GetStream(), problem, [&]() { - ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); + auto ctx_copy = ctx; + ctx_copy.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); const auto params = - conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx, problem)}; + conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx_copy, problem)}; auto conv_sols = GetConvSolutions(ctx, problem, solutions); auto eval_sols = EvaluateConvSolutions(ctx, problem, invoke_ctx, conv_sols, model_result); @@ -390,23 +391,26 @@ std::vector VerifiedFDBSolution(ExecutionContext& ctx, { // entry considered bad, trigger tuning MIOPEN_LOG_I2("TrustVerify: Regenerate entry for user db"); - ctx.do_search = true; - ctx.db_update = true; + ctx_copy.do_search = true; + ctx_copy.db_update = true; - return FindCore(invoke_ctx, - ctx, + auto ret = FindCore(invoke_ctx, + ctx_copy, problem, params, conv::GetConvSolverFinders(), std::nullopt, force_attach_binary); + ctx.generic_search_worst_time = ctx_copy.generic_search_worst_time; + ctx.generic_search_best_time = ctx_copy.generic_search_best_time; + return ret; } }); return results; } -std::vector FindConvolution(ExecutionContext& ctx, +std::vector FindConvolution(const ExecutionContext& ctx, const conv::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx, int requestAlgoCount, @@ -475,24 +479,28 @@ std::vector FindConvolution(ExecutionContext& ctx, else { results = UserFindDbRecord::TryLoad(ctx.GetStream(), problem, [&]() { - ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); + auto ctx_copy = ctx; + ctx_copy.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); const auto params = - conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx, problem)}; + conv::ConvFindParameters{conv.IsWinograd3x3SupportedAndFast(ctx_copy, problem)}; if(findMode.IsTrustVerify(ctx)) { MIOPEN_LOG_I2("TrustVerify: Generate entry for user db"); - ctx.do_search = true; - ctx.db_update = true; + ctx_copy.do_search = true; + ctx_copy.db_update = true; } - return FindCore(invoke_ctx, - ctx, + auto ret = FindCore(invoke_ctx, + ctx_copy, problem, params, conv::GetConvSolverFinders(), std::nullopt, force_attach_binary); + ctx.generic_search_worst_time = ctx_copy.generic_search_worst_time; + ctx.generic_search_best_time = ctx_copy.generic_search_best_time; + return ret; }); } @@ -562,7 +570,7 @@ void ConvolutionDescriptor::FindConvFwdAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription(xDesc, wDesc, yDesc, *this, conv::Direction::Forward); - auto ctx = [&] { + const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; @@ -1101,7 +1109,7 @@ void ConvolutionDescriptor::FindConvBwdDataAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription{dyDesc, wDesc, dxDesc, *this, conv::Direction::BackwardData}; - auto ctx = [&] { + const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; @@ -1308,7 +1316,7 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(const Handle& handle, const auto problem = conv::ProblemDescription{dyDesc, dwDesc, xDesc, *this, conv::Direction::BackwardWeights}; - auto ctx = [&] { + const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); tmp.do_search = exhaustiveSearch; diff --git a/projects/miopen/src/solver/batchnorm/backward_spatial.cpp b/projects/miopen/src/solver/batchnorm/backward_spatial.cpp index 1eba75aca70..26a386add1f 100644 --- a/projects/miopen/src/solver/batchnorm/backward_spatial.cpp +++ b/projects/miopen/src/solver/batchnorm/backward_spatial.cpp @@ -184,7 +184,7 @@ bool BnBwdTrainingSpatial::IsValidPerformanceConfig( } PerformanceConfigBnBwdBackward -BnBwdTrainingSpatial::Search(ExecutionContext& ctx, +BnBwdTrainingSpatial::Search(const ExecutionContext& ctx, const miopen::batchnorm::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/batchnorm/forward_spatial.cpp b/projects/miopen/src/solver/batchnorm/forward_spatial.cpp index 0fe64589b78..5f7dbc57e50 100644 --- a/projects/miopen/src/solver/batchnorm/forward_spatial.cpp +++ b/projects/miopen/src/solver/batchnorm/forward_spatial.cpp @@ -173,7 +173,7 @@ bool BnFwdTrainingSpatial::IsValidPerformanceConfig( } PerformanceConfigBnFwdTraining -BnFwdTrainingSpatial::Search(ExecutionContext& ctx, +BnFwdTrainingSpatial::Search(const ExecutionContext& ctx, const miopen::batchnorm::ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp index 3cab1911f40..9b934ec041a 100644 --- a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp +++ b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp @@ -940,7 +940,7 @@ ConvMPBidirectWinograd_xdlops::G template PerformanceImplicitGemmForwardV4R4Xdlops ConvMPBidirectWinograd_xdlops::Search( - ExecutionContext& ctx, + const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { @@ -948,7 +948,7 @@ ConvMPBidirectWinograd_xdlops::S GetTransformedInvokeContext(problem, invoke_ctx); const auto xdlops_problem = GetTransformedProblem(problem); - auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); + const auto xdlops_ctx = GetTransformedConvContext(ctx, xdlops_problem); return ConvHipImplicitGemmForwardV4R4Xdlops().Search( xdlops_ctx, xdlops_problem, xdlops_invoke_ctx); diff --git a/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp b/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp index a31c83afd88..6d7e6761492 100644 --- a/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_1x1u.cpp @@ -911,7 +911,7 @@ ConvSolution ConvAsm1x1U::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm1x1U ConvAsm1x1U::Search(ExecutionContext& ctx, +PerformanceConfigConvAsm1x1U ConvAsm1x1U::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp b/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp index 0b899e6831d..420d37d2174 100644 --- a/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_1x1u_stride2.cpp @@ -753,7 +753,7 @@ ConvSolution ConvAsm1x1UV2::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm1x1UV2 ConvAsm1x1UV2::Search(ExecutionContext& ctx, +PerformanceConfigConvAsm1x1UV2 ConvAsm1x1UV2::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp b/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp index 33805c38558..839517e3c7c 100644 --- a/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_3x3u.cpp @@ -320,7 +320,7 @@ ConvSolution ConvAsm3x3U::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsm3x3U ConvAsm3x3U::Search(ExecutionContext& ctx, +PerformanceConfigConvAsm3x3U ConvAsm3x3U::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp index 8272537a4cd..53486560097 100644 --- a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp @@ -853,7 +853,7 @@ ConvSolution ConvAsmBwdWrW1x1::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigConvAsmBwdWrW1x1 ConvAsmBwdWrW1x1::Search(ExecutionContext& ctx, +PerformanceConfigConvAsmBwdWrW1x1 ConvAsmBwdWrW1x1::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp index 8db941cd038..6d49a26fdf4 100644 --- a/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp @@ -597,7 +597,7 @@ ConvSolution ConvAsmBwdWrW3x3::GetSolution(const ExecutionContext& ctx, return result; } -PerformanceConfigAsmDirect3x3WrW ConvAsmBwdWrW3x3::Search(ExecutionContext& ctx, +PerformanceConfigAsmDirect3x3WrW ConvAsmBwdWrW3x3::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index 95f831fee88..117fba50c2f 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -951,7 +951,7 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp index 48870e3478d..041e69b9261 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp @@ -541,7 +541,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCFwdDlopsNCHWC -ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::Search(ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index 8ebdaa2e8e7..93d2e256685 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -804,7 +804,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCFwdXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::Search(ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 6f1b622085c..1779f124cbb 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -850,7 +850,7 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsValidPerformanceConfig( } PerformanceConfigAsmImplicitGemmGTCWrwXdlopsNHWC -ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(ExecutionContext& ctx, +ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp b/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp index 72dff87f0fe..61dba30c7f6 100644 --- a/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp +++ b/projects/miopen/src/solver/conv/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp @@ -268,7 +268,7 @@ std::size_t ConvCkIgemmFwdV6r1DlopsNchw::GetWorkspaceSize(const ExecutionContext } PerformanceConvCkIgemmFwdV6r1DlopsNchw -ConvCkIgemmFwdV6r1DlopsNchw::Search(ExecutionContext& ctx, +ConvCkIgemmFwdV6r1DlopsNchw::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp index 2c7dc6bfa2d..78a09c392b6 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp @@ -498,7 +498,7 @@ ConvHipImplicitGemm3DGroupBwdXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupBwdXdlops -ConvHipImplicitGemm3DGroupBwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupBwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp index f7907c93d46..751bd964c9d 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp @@ -561,7 +561,7 @@ ConvHipImplicitGemm3DGroupFwdXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupFwdXdlops -ConvHipImplicitGemm3DGroupFwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupFwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp index 078c95f6440..0c15d775091 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp @@ -473,7 +473,7 @@ ConvHipImplicitGemm3DGroupWrwXdlops::GetWorkspaceSize(const ExecutionContext&, } PerformanceConfigHipImplicitGemm3DGroupWrwXdlops -ConvHipImplicitGemm3DGroupWrwXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemm3DGroupWrwXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp index 1c6a697990a..9f7bc5f0876 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_data_xdlops.cpp @@ -261,7 +261,7 @@ bool ConvHipImplicitGemmBwdXdlops::IsValidPerformanceConfig( } PerformanceConfigHipImplicitGemmBwdXdlops -ConvHipImplicitGemmBwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmBwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp index e7e06663481..de2f54d1971 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -705,7 +705,7 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV1R1 -ConvHipImplicitGemmBwdDataV1R1::Search(ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV1R1::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp index 8a464b1d7c6..65e4042063c 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp @@ -827,7 +827,7 @@ bool ConvHipImplicitGemmBwdDataV1R1Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmBwdV1R1Xdlops -ConvHipImplicitGemmBwdDataV1R1Xdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV1R1Xdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp index 7956eb166a6..3fa7367f269 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp @@ -815,7 +815,7 @@ bool ConvHipImplicitGemmBwdDataV4R1::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV4R1 -ConvHipImplicitGemmBwdDataV4R1::Search(ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV4R1::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp index 11c01bf38c1..d63d23d1bd4 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp @@ -898,7 +898,7 @@ bool ConvHipImplicitGemmBwdDataV4R1Xdlops::IsValidPerformanceConfig( } PerformanceImplicitGemmBwdDataV4R1Xdlops -ConvHipImplicitGemmBwdDataV4R1Xdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmBwdDataV4R1Xdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp index 47188934411..e73581da159 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r1.cpp @@ -189,14 +189,14 @@ bool ConvHipImplicitGemmV4R1WrW::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R1 -ConvHipImplicitGemmV4R1Fwd::Search(ExecutionContext& ctx, +ConvHipImplicitGemmV4R1Fwd::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { return GenericSearch(*this, ctx, problem, invoke_ctx); } PerformanceImplicitGemmV4R1 -ConvHipImplicitGemmV4R1WrW::Search(ExecutionContext& ctx, +ConvHipImplicitGemmV4R1WrW::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp index ee62d628952..3d8f866e1aa 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4.cpp @@ -636,7 +636,7 @@ bool ConvHipImplicitGemmV4R4Fwd::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R4Fwd -ConvHipImplicitGemmV4R4Fwd::Search(ExecutionContext& ctx, +ConvHipImplicitGemmV4R4Fwd::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp index d4605dcff5c..55be2a71b33 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp @@ -1049,7 +1049,7 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmForwardV4R4Xdlops -ConvHipImplicitGemmForwardV4R4Xdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R4Xdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp index 55f8ad65b87..850edf8da70 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp @@ -1113,7 +1113,7 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::IsApplicable( } PerformanceImplicitGemmForwardV4R4Xdlops_Padded_Gemm -ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::Search(ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp index a7a27d1721a..cf3bc09021b 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp @@ -1086,7 +1086,7 @@ bool ConvHipImplicitGemmForwardV4R5Xdlops::IsApplicable(const ExecutionContext& } PerformanceImplicitGemmForwardV4R5Xdlops -ConvHipImplicitGemmForwardV4R5Xdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmForwardV4R5Xdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp index caad7df232b..ece0e1b1197 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_fwd_xdlops.cpp @@ -263,7 +263,7 @@ bool ConvHipImplicitGemmFwdXdlops::IsValidPerformanceConfig( } PerformanceConfigHipImplicitGemmFwdXdlops -ConvHipImplicitGemmFwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmFwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp index e876ef1b7fc..4533508e15b 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp @@ -534,7 +534,7 @@ size_t ConvHipImplicitGemmGroupBwdXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupBwdXdlops -ConvHipImplicitGemmGroupBwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmGroupBwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp index 140a43e8424..173ad0c74be 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp @@ -495,7 +495,7 @@ size_t ConvHipImplicitGemmGroupFwdXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupFwdXdlops -ConvHipImplicitGemmGroupFwdXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmGroupFwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp index 3442366f2a7..f6a38370b43 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp @@ -574,7 +574,7 @@ size_t ConvHipImplicitGemmGroupWrwXdlops::GetWorkspaceSize(const ExecutionContex } PerformanceConfigHipImplicitGemmGroupWrwXdlops -ConvHipImplicitGemmGroupWrwXdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmGroupWrwXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp index 8cf382f340b..113fb6af4c3 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4.cpp @@ -639,7 +639,7 @@ bool ConvHipImplicitGemmV4R4WrW::IsValidPerformanceConfig( } PerformanceImplicitGemmV4R4WrW -ConvHipImplicitGemmV4R4WrW::Search(ExecutionContext& ctx, +ConvHipImplicitGemmV4R4WrW::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp index df026aee619..8a0887a06c6 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp @@ -1115,7 +1115,7 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops::IsApplicable(const ExecutionContext& ctx, } PerformanceImplicitGemmWrwV4R4Xdlops -ConvHipImplicitGemmWrwV4R4Xdlops::Search(ExecutionContext& ctx, +ConvHipImplicitGemmWrwV4R4Xdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp index acad1009b79..dd5032a5c88 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp @@ -1194,7 +1194,7 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::IsApplicable( } PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm -ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::Search(ExecutionContext& ctx, +ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp index 7fa4cfc8b6b..03357d02517 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd.cpp @@ -91,7 +91,7 @@ bool ConvMlirIgemmBwd::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmBwd::Search(ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmBwd::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp index 068f4f585fb..79d6ff3694e 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_bwd_xdlops.cpp @@ -87,7 +87,7 @@ bool ConvMlirIgemmBwdXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmBwdXdlops::Search(ExecutionContext& ctx, +ConvMlirIgemmBwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp index 4f6f58fcf01..7d03507d392 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd.cpp @@ -153,7 +153,7 @@ bool ConvMlirIgemmFwd::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmFwd::Search(ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmFwd::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp index 01c91b8ecb6..02865f6e9ab 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_fwd_xdlops.cpp @@ -212,7 +212,7 @@ bool ConvMlirIgemmFwdXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmFwdXdlops::Search(ExecutionContext& ctx, +ConvMlirIgemmFwdXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp index fcc9091d9b8..64b19192071 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw.cpp @@ -92,7 +92,7 @@ bool ConvMlirIgemmWrW::IsValidPerformanceConfig(const ExecutionContext& ctx, return config.IsValid(ctx, problem); } -PerformanceConvMlirIgemm ConvMlirIgemmWrW::Search(ExecutionContext& ctx, +PerformanceConvMlirIgemm ConvMlirIgemmWrW::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp index 791908acb35..15ef9363ab9 100644 --- a/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_mlir_igemm_wrw_xdlops.cpp @@ -88,7 +88,7 @@ bool ConvMlirIgemmWrWXdlops::IsValidPerformanceConfig( } PerformanceConvMlirIgemmXdlops -ConvMlirIgemmWrWXdlops::Search(ExecutionContext& ctx, +ConvMlirIgemmWrWXdlops::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp index d99403d4389..b6a177d31b5 100644 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp +++ b/projects/miopen/src/solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp @@ -770,7 +770,7 @@ ConvSolution ConvOclBwdWrW2::GetSolution( template PerformanceConfigConvOclBwdWrw2 -ConvOclBwdWrW2::Search(ExecutionContext& ctx, +ConvOclBwdWrW2::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp index f7c786d9166..213c989f376 100644 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp +++ b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp @@ -213,7 +213,7 @@ static int MeasurePerfConfig(const Handle& handle, } LegacyPerformanceConfig -ConvOclDirectFwdLegacyExhaustiveSearch::Search(ExecutionContext& ctx, +ConvOclDirectFwdLegacyExhaustiveSearch::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv/conv_winoRxS.cpp b/projects/miopen/src/solver/conv/conv_winoRxS.cpp index 3b8f0861cd9..2ec293fef73 100644 --- a/projects/miopen/src/solver/conv/conv_winoRxS.cpp +++ b/projects/miopen/src/solver/conv/conv_winoRxS.cpp @@ -429,7 +429,7 @@ bool ConvBinWinoRxS::IsValidPerformanceConfig( template PerformanceConfigConvBinWinogradRxS -ConvBinWinoRxS::Search(ExecutionContext& ctx, +ConvBinWinoRxS::Search(const ExecutionContext& ctx, const ProblemDescription& problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp index fc7ac3c0fa0..5b1e4b0a006 100644 --- a/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_asm_1x1u_bias_activ_fused.cpp @@ -85,7 +85,7 @@ bool ConvBiasActivAsm1x1U::IsValidPerformanceConfig( } PerformanceConfigConvBiasActivAsm1x1U -ConvBiasActivAsm1x1U::Search(FusionContext& context, +ConvBiasActivAsm1x1U::Search(const FusionContext& context, const FusionDescription& problem, const AnyInvokeParams& invoke_params) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp index cfaa85479e5..0d3a82fab7a 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp @@ -387,7 +387,7 @@ bool ConvCKIgemmFwdBiasActivFused::IsValidPerformanceConfig( } PerformanceConfigConvCKIgemmFwdBiasActivFused -ConvCKIgemmFwdBiasActivFused::Search(FusionContext& ctx, +ConvCKIgemmFwdBiasActivFused::Search(const FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp index 547d24b2b80..0495d4598bd 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_fwd_bias_res_add_activ_fused.cpp @@ -370,7 +370,7 @@ bool ConvCKIgemmFwdBiasResAddActivFused::IsValidPerformanceConfig( } PerfConfigConvCKIgemmFwdBiasResAddActivFused -ConvCKIgemmFwdBiasResAddActivFused::Search(FusionContext& ctx, +ConvCKIgemmFwdBiasResAddActivFused::Search(const FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp index a088f4a7005..03c4f35874e 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp @@ -581,7 +581,7 @@ size_t ConvCKIgemmGrpFwdActivFused::GetWorkspaceSize(const FusionContext&, } PerformanceConfigConvCKIgemmGrpFwdActivFused -ConvCKIgemmGrpFwdActivFused::Search(FusionContext& ctx, +ConvCKIgemmGrpFwdActivFused::Search(const FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp index 1e4679eb700..fed835c8b60 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_bias_activ_fused.cpp @@ -592,7 +592,7 @@ ConvCKIgemmGrpFwdBiasActivFused::GetWorkspaceSize(const FusionContext&, } PerformanceConfigConvCKIgemmGrpFwdBiasActivFused -ConvCKIgemmGrpFwdBiasActivFused::Search(FusionContext& ctx, +ConvCKIgemmGrpFwdBiasActivFused::Search(const FusionContext& ctx, const FusionDescription& fdesc_problem, const AnyInvokeParams& invoke_ctx) const { diff --git a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp index e480db107a8..6a1a9f35877 100644 --- a/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp +++ b/projects/miopen/src/solver/conv_ocl_dir2Dfwd_fused.cpp @@ -43,12 +43,12 @@ namespace solver { namespace fusion { PerformanceConfigConvOclDirectFwdFused -ConvOclDirectFwdFused::Search(FusionContext& context, +ConvOclDirectFwdFused::Search(const FusionContext& context, const FusionDescription& problem, const AnyInvokeParams& invoke_params) const { const auto conv_problem = problem.GetConvProblem(0, miopen::conv::Direction::Forward); - auto conv_ctx = context.GetConvContext(conv_problem); + const auto conv_ctx = context.GetConvContext(conv_problem); const auto legacy = conv::ConvOclDirectFwd{}; const auto& fusion_invoke_params = invoke_params.CastTo(); const auto wei_ocl_ptr = dynamic_cast( From a6ea744de5f75ed662ed047786615606091c92ee Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 11 Sep 2025 23:23:51 +0000 Subject: [PATCH 10/26] add envs to modify GenericSearch skips --- .../miopen/src/include/miopen/generic_search.hpp | 12 +++++++----- .../src/include/miopen/generic_search_controls.hpp | 4 ++++ 2 files changed, 11 insertions(+), 5 deletions(-) diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 1aa2c77df5e..a526480d786 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -482,12 +482,14 @@ auto GenericSearch(const Solver s, float worst_time = std::numeric_limits::max(); size_t n_failed = 0; size_t n_best = 0; + // terminate search when perf is less than cutoff float cutoff_time = context.generic_search_worst_time; if(cutoff_time < std::numeric_limits::max()) - cutoff_time *= 10; - float overall_best_time = context.generic_search_best_time; - if(overall_best_time < std::numeric_limits::max()) - overall_best_time *= 1.30f; + cutoff_time *= env::value(MIOPEN_SEARCH_CUTOFF_MUL); + // skip detailed measurement for configs slower than skip_time + float skip_time = context.generic_search_best_time; + if(skip_time < std::numeric_limits::max()) + skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; HeartBeat heartbeat; heartbeat.Start(); @@ -616,7 +618,7 @@ auto GenericSearch(const Solver s, // the mean value with outliers removed for calculating best config. constexpr int N_RUNS = 10; last_imprv++; - if(elapsed_time < worst_time * 1.10f && elapsed_time < overall_best_time) + if(elapsed_time < worst_time * 1.10f && elapsed_time < skip_time) { MIOPEN_LOG_I2("Finding average for: " << elapsed_time << " / " << best_time << " = " << (elapsed_time / best_time)); diff --git a/projects/miopen/src/include/miopen/generic_search_controls.hpp b/projects/miopen/src/include/miopen/generic_search_controls.hpp index 10a1b1b1e97..a2ea1cded2b 100644 --- a/projects/miopen/src/include/miopen/generic_search_controls.hpp +++ b/projects/miopen/src/include/miopen/generic_search_controls.hpp @@ -52,3 +52,7 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY) const size_t MIOPEN_DEFAULT_VERIFY_TOLERANCE_PCT = 15; MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_VERIFY_TOLERANCE_PCT, MIOPEN_DEFAULT_VERIFY_TOLERANCE_PCT) + +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_CUTOFF_MUL, 10) + +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_SKIP_PCT, 130) From 2f548b24655a8084a780393333cd336ad9b795d3 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 14 Oct 2025 11:04:48 -0500 Subject: [PATCH 11/26] format --- .../src/include/miopen/execution_context.hpp | 12 ++++---- .../src/include/miopen/generic_search.hpp | 9 +++--- projects/miopen/src/ocl/convolutionocl.cpp | 28 +++++++++---------- 3 files changed, 24 insertions(+), 25 deletions(-) diff --git a/projects/miopen/src/include/miopen/execution_context.hpp b/projects/miopen/src/include/miopen/execution_context.hpp index e730845a4ba..23d78a67ac5 100644 --- a/projects/miopen/src/include/miopen/execution_context.hpp +++ b/projects/miopen/src/include/miopen/execution_context.hpp @@ -94,12 +94,12 @@ struct MIOPEN_INTERNALS_EXPORT ExecutionContext // to optimize the getWorkspaceSize() calls for speed. This specific optimization is correct // because Solvers shall be written so that the required workspace size does not depend on the // performance config. - bool disable_perfdb_access = false; - bool use_dynamic_solutions_only = false; - bool is_for_generic_search = false; - mutable bool search_cutoff = false; - mutable float generic_search_worst_time = std::numeric_limits::max(); - mutable float generic_search_best_time = std::numeric_limits::max(); + bool disable_perfdb_access = false; + bool use_dynamic_solutions_only = false; + bool is_for_generic_search = false; + mutable bool search_cutoff = false; + mutable float generic_search_worst_time = std::numeric_limits::max(); + mutable float generic_search_best_time = std::numeric_limits::max(); inline const Handle& GetStream() const { return *stream; } inline void SetStream(const Handle* stream_) { stream = stream_; } diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 3dbd32f293c..212a457d9e1 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -486,7 +486,7 @@ auto GenericSearch(const Solver s, float cutoff_time = context.generic_search_worst_time; if(cutoff_time < std::numeric_limits::max()) cutoff_time *= env::value(MIOPEN_SEARCH_CUTOFF_MUL); - // skip detailed measurement for configs slower than skip_time + // skip detailed measurement for configs slower than skip_time float skip_time = context.generic_search_best_time; if(skip_time < std::numeric_limits::max()) skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; @@ -593,8 +593,7 @@ auto GenericSearch(const Solver s, if(elapsed_time > cutoff_time) { MIOPEN_LOG_I2("Ending Search, measured time: " - << elapsed_time << " was greater than cutoff: " - << cutoff_time); + << elapsed_time << " was greater than cutoff: " << cutoff_time); for(const auto& kernelInfo : current_solution.construction_params) profile_h.ClearProgram(kernelInfo.kernel_file, kernelInfo.comp_options); break; @@ -700,8 +699,8 @@ auto GenericSearch(const Solver s, // if using cutoff for search update timing if(context.search_cutoff && best_time < context.generic_search_best_time) { - float new_worst = (perf_sols.end() - 1)->time; - context_.generic_search_best_time = best_time; + float new_worst = (perf_sols.end() - 1)->time; + context_.generic_search_best_time = best_time; context_.generic_search_worst_time = new_worst; MIOPEN_LOG_I2("Times updated, best: " << best_time << " worst: " << new_worst); } diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index 51b9375b8b3..01db3d73529 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -395,14 +395,14 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, ctx_copy.db_update = true; auto ret = FindCore(invoke_ctx, - ctx_copy, - problem, - params, - conv::GetConvSolverFinders(), - std::nullopt, - force_attach_binary); + ctx_copy, + problem, + params, + conv::GetConvSolverFinders(), + std::nullopt, + force_attach_binary); ctx.generic_search_worst_time = ctx_copy.generic_search_worst_time; - ctx.generic_search_best_time = ctx_copy.generic_search_best_time; + ctx.generic_search_best_time = ctx_copy.generic_search_best_time; return ret; } }); @@ -492,14 +492,14 @@ std::vector FindConvolution(const ExecutionContext& ctx, } auto ret = FindCore(invoke_ctx, - ctx_copy, - problem, - params, - conv::GetConvSolverFinders(), - std::nullopt, - force_attach_binary); + ctx_copy, + problem, + params, + conv::GetConvSolverFinders(), + std::nullopt, + force_attach_binary); ctx.generic_search_worst_time = ctx_copy.generic_search_worst_time; - ctx.generic_search_best_time = ctx_copy.generic_search_best_time; + ctx.generic_search_best_time = ctx_copy.generic_search_best_time; return ret; }); } From 9ef96643ea51e5011dbea92766cda8ce44a8f425 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 14 Oct 2025 12:24:59 -0500 Subject: [PATCH 12/26] format --- projects/miopen/src/ocl/convolutionocl.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index 01db3d73529..c3f9ca20f48 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -394,7 +394,7 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, ctx_copy.do_search = true; ctx_copy.db_update = true; - auto ret = FindCore(invoke_ctx, + auto ret = FindCore(invoke_ctx, ctx_copy, problem, params, @@ -491,7 +491,7 @@ std::vector FindConvolution(const ExecutionContext& ctx, ctx_copy.db_update = true; } - auto ret = FindCore(invoke_ctx, + auto ret = FindCore(invoke_ctx, ctx_copy, problem, params, From 179b0d6a8de57f0a6c7c50ab966f905d0e72a755 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 15 Oct 2025 15:16:50 -0500 Subject: [PATCH 13/26] adjust tuning policy smoke test --- projects/miopen/test/gtest/smoke_tuning_policy.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/projects/miopen/test/gtest/smoke_tuning_policy.cpp b/projects/miopen/test/gtest/smoke_tuning_policy.cpp index 3b2c7ace9cc..98e2bbc751a 100644 --- a/projects/miopen/test/gtest/smoke_tuning_policy.cpp +++ b/projects/miopen/test/gtest/smoke_tuning_policy.cpp @@ -29,6 +29,11 @@ TEST_F(CPU_TuningPolicy_NONE, TestTuningPolicyGetterAndSetter) EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchDbUpdate); + EXPECT_EQ(miopenSetTuningPolicy(&handle, miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate), + miopenStatusSuccess); + EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); + EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate); + EXPECT_EQ(miopenSetTuningPolicy(&handle, miopenTuningPolicy_t::miopenTuningPolicyNone), miopenStatusSuccess); EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); @@ -39,7 +44,7 @@ TEST_F(CPU_TuningPolicy_NONE, TestTuningPolicyGetterAndSetter) EXPECT_EQ(miopenSetTuningPolicy(&handle, static_cast(0)), miopenStatusBadParm); - EXPECT_EQ(miopenSetTuningPolicy(&handle, static_cast(6)), + EXPECT_EQ(miopenSetTuningPolicy(&handle, static_cast(7)), miopenStatusBadParm); EXPECT_EQ(miopenSetTuningPolicy(&handle, static_cast(1000)), From dc3b189cc71ed561a27b939e07ca9e3c594d54d0 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 17 Oct 2025 09:52:16 -0500 Subject: [PATCH 14/26] format --- projects/miopen/test/gtest/smoke_tuning_policy.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/miopen/test/gtest/smoke_tuning_policy.cpp b/projects/miopen/test/gtest/smoke_tuning_policy.cpp index 98e2bbc751a..3f093f983f8 100644 --- a/projects/miopen/test/gtest/smoke_tuning_policy.cpp +++ b/projects/miopen/test/gtest/smoke_tuning_policy.cpp @@ -29,7 +29,8 @@ TEST_F(CPU_TuningPolicy_NONE, TestTuningPolicyGetterAndSetter) EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchDbUpdate); - EXPECT_EQ(miopenSetTuningPolicy(&handle, miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate), + EXPECT_EQ(miopenSetTuningPolicy(&handle, + miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate), miopenStatusSuccess); EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate); From 8795a3a3ed4672fc7144b1b30328beed20c7e128 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 7 Nov 2025 19:01:41 +0000 Subject: [PATCH 15/26] switch to using env MIOPEN_SEARCH_CUTOFF --- projects/miopen/include/miopen/miopen.h | 6 ++---- projects/miopen/src/find_controls.cpp | 7 ------- .../src/include/miopen/execution_context.hpp | 1 - .../miopen/src/include/miopen/find_controls.hpp | 15 +++------------ .../miopen/src/include/miopen/find_solution.hpp | 3 --- .../miopen/src/include/miopen/generic_search.hpp | 4 +++- .../include/miopen/generic_search_controls.hpp | 2 ++ .../miopen/test/gtest/smoke_tuning_policy.cpp | 8 +------- 8 files changed, 11 insertions(+), 35 deletions(-) diff --git a/projects/miopen/include/miopen/miopen.h b/projects/miopen/include/miopen/miopen.h index 65a6e66ce7e..cdd1ace9c2c 100644 --- a/projects/miopen/include/miopen/miopen.h +++ b/projects/miopen/include/miopen/miopen.h @@ -8523,8 +8523,7 @@ MIOPEN_EXPORT miopenStatus_t miopenMultiMarginLossForward(miopenHandle_t handle, * 2. DbUpdate: Do not skip auto-tune even if PerfDb already contains optimized values. * 3. Search: Search the database first; if no record is found, tune and update the database. * 4. SearchDbUpdate: Combination of Search and DbUpdate. - * 5. SearchCutoffDbUpdate: Search with tuning perf cutoff and DbUpdate. - * 6. DbClean: Remove existing entry, do not tune. + * 5. DbClean: Remove existing entry, do not tune. * Note: TuningPolicy has higher priority over MIOPEN_FIND_ENFORCE. */ typedef enum @@ -8533,8 +8532,7 @@ typedef enum miopenTuningPolicyDbUpdate = 2, /*!< tune and update the db */ miopenTuningPolicySearch = 3, /*!< search db first, if record not found tune but do not update the db */ miopenTuningPolicySearchDbUpdate = 4, /*!< combination of Search and DbUpdate */ - miopenTuningPolicySearchCutoffDbUpdate = 5, /*!< combination of Search and DbUpdate with early exit strategy*/ - miopenTuningPolicyDbClean = 6, /*!< remove existing entry, do not tune */ + miopenTuningPolicyDbClean = 5, /*!< remove existing entry, do not tune */ } miopenTuningPolicy_t; /*! @ingroup handle diff --git a/projects/miopen/src/find_controls.cpp b/projects/miopen/src/find_controls.cpp index 6eb332fc69d..4e87575cd89 100644 --- a/projects/miopen/src/find_controls.cpp +++ b/projects/miopen/src/find_controls.cpp @@ -62,8 +62,6 @@ static_assert(FindEnforceAction::Search == static_cast(miopenTuningPolicySearch)); static_assert(FindEnforceAction::SearchDbUpdate == static_cast(miopenTuningPolicySearchDbUpdate)); -static_assert(FindEnforceAction::SearchCutoffDbUpdate == - static_cast(miopenTuningPolicySearchCutoffDbUpdate)); static_assert(FindEnforceAction::DbClean == static_cast(miopenTuningPolicyDbClean)); @@ -77,7 +75,6 @@ const char* ToCString(const FindEnforceAction mode) case FindEnforceAction::DbUpdate: return "DB_UPDATE"; case FindEnforceAction::Search: return "SEARCH"; case FindEnforceAction::SearchDbUpdate: return "SEARCH_DB_UPDATE"; - case FindEnforceAction::SearchCutoffDbUpdate: return "SEARCH_CUTOFF_DB_UPDATE"; case FindEnforceAction::DbClean: return "CLEAN"; } return ""; @@ -106,10 +103,6 @@ FindEnforceAction GetFindEnforceActionImpl() { return FindEnforceAction::SearchDbUpdate; } - else if(str == "SEARCH_CUTOFF_DB_UPDATE") - { - return FindEnforceAction::SearchCutoffDbUpdate; - } else if(str == "DB_CLEAN") { return FindEnforceAction::DbClean; diff --git a/projects/miopen/src/include/miopen/execution_context.hpp b/projects/miopen/src/include/miopen/execution_context.hpp index 23d78a67ac5..9885eaab6d0 100644 --- a/projects/miopen/src/include/miopen/execution_context.hpp +++ b/projects/miopen/src/include/miopen/execution_context.hpp @@ -97,7 +97,6 @@ struct MIOPEN_INTERNALS_EXPORT ExecutionContext bool disable_perfdb_access = false; bool use_dynamic_solutions_only = false; bool is_for_generic_search = false; - mutable bool search_cutoff = false; mutable float generic_search_worst_time = std::numeric_limits::max(); mutable float generic_search_best_time = std::numeric_limits::max(); diff --git a/projects/miopen/src/include/miopen/find_controls.hpp b/projects/miopen/src/include/miopen/find_controls.hpp index 2e5246fba5a..98ac73a8f0f 100644 --- a/projects/miopen/src/include/miopen/find_controls.hpp +++ b/projects/miopen/src/include/miopen/find_controls.hpp @@ -54,7 +54,6 @@ enum class FindEnforceAction DbUpdate, Search, SearchDbUpdate, - SearchCutoffDbUpdate, DbClean, Last_ = DbClean, Default_ = None, @@ -86,23 +85,15 @@ class MIOPEN_INTERNALS_EXPORT FindEnforce template bool IsSearch(const Context& context) const { - return IsEnabled(context) && (action == FindEnforceAction::Search || - action == FindEnforceAction::SearchDbUpdate || - action == FindEnforceAction::SearchCutoffDbUpdate); - } - - template - bool IsSearchCutoff(const Context& context) const - { - return IsEnabled(context) && (action == FindEnforceAction::SearchCutoffDbUpdate); + return IsEnabled(context) && + (action == FindEnforceAction::Search || action == FindEnforceAction::SearchDbUpdate); } template bool IsDbUpdate(const Context& context) const { return IsEnabled(context) && (action == FindEnforceAction::DbUpdate || - action == FindEnforceAction::SearchDbUpdate || - action == FindEnforceAction::SearchCutoffDbUpdate); + action == FindEnforceAction::SearchDbUpdate); } template diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index 11e2a70157e..f6429d1123f 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -140,9 +140,6 @@ auto FindSolutionImpl(rank<1>, if(context.do_search || enforce.IsSearch(context)) // TODO: Make it a customization point { - if(enforce.IsSearchCutoff(context)) - context.search_cutoff = true; - auto record = DbRecord(DbKinds::PerfDb, problem); if(env::enabled(MIOPEN_WARN_SEARCH)) MIOPEN_LOG_W("Search Start: " << record.GetKey() << " : " << s.SolverDbId() diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 212a457d9e1..b192e6725e5 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -482,6 +482,8 @@ auto GenericSearch(const Solver s, float worst_time = std::numeric_limits::max(); size_t n_failed = 0; size_t n_best = 0; + // enable early search termination + bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF); // terminate search when perf is less than cutoff float cutoff_time = context.generic_search_worst_time; if(cutoff_time < std::numeric_limits::max()) @@ -697,7 +699,7 @@ auto GenericSearch(const Solver s, }); // if using cutoff for search update timing - if(context.search_cutoff && best_time < context.generic_search_best_time) + if(using_search_cutoff == true && best_time < context.generic_search_best_time) { float new_worst = (perf_sols.end() - 1)->time; context_.generic_search_best_time = best_time; diff --git a/projects/miopen/src/include/miopen/generic_search_controls.hpp b/projects/miopen/src/include/miopen/generic_search_controls.hpp index a2ea1cded2b..a515f27fd3e 100644 --- a/projects/miopen/src/include/miopen/generic_search_controls.hpp +++ b/projects/miopen/src/include/miopen/generic_search_controls.hpp @@ -53,6 +53,8 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY) const size_t MIOPEN_DEFAULT_VERIFY_TOLERANCE_PCT = 15; MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_VERIFY_TOLERANCE_PCT, MIOPEN_DEFAULT_VERIFY_TOLERANCE_PCT) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) + MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_CUTOFF_MUL, 10) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_SKIP_PCT, 130) diff --git a/projects/miopen/test/gtest/smoke_tuning_policy.cpp b/projects/miopen/test/gtest/smoke_tuning_policy.cpp index d016c804bc4..401f56d3cf6 100644 --- a/projects/miopen/test/gtest/smoke_tuning_policy.cpp +++ b/projects/miopen/test/gtest/smoke_tuning_policy.cpp @@ -53,12 +53,6 @@ TEST_F(CPU_TuningPolicy_NONE, TestTuningPolicyGetterAndSetterValidValues) EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchDbUpdate); - EXPECT_EQ(miopenSetTuningPolicy(&handle, - miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate), - miopenStatusSuccess); - EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); - EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicySearchCutoffDbUpdate); - EXPECT_EQ(miopenSetTuningPolicy(&handle, miopenTuningPolicy_t::miopenTuningPolicyDbClean), miopenStatusSuccess); EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); @@ -74,7 +68,7 @@ TEST_F(CPU_TuningPolicy_NONE, TestTuningPolicyGetterAndSetterErrorHandling) { testSetInvalidValue(miopenTuningPolicy_t::miopenTuningPolicyDbUpdate, -1); testSetInvalidValue(miopenTuningPolicy_t::miopenTuningPolicySearch, 0); - testSetInvalidValue(miopenTuningPolicy_t::miopenTuningPolicySearchDbUpdate, 7); + testSetInvalidValue(miopenTuningPolicy_t::miopenTuningPolicySearchDbUpdate, 6); testSetInvalidValue(miopenTuningPolicy_t::miopenTuningPolicyDbClean, 1000); } From f6fceac4f504195b14fa9e1b6301a7530bca1523 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Sat, 8 Nov 2025 00:20:30 +0000 Subject: [PATCH 16/26] stop populating perf_sols when unused --- projects/miopen/include/miopen/miopen.h | 10 +++++----- projects/miopen/src/include/miopen/generic_search.hpp | 5 ++++- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/projects/miopen/include/miopen/miopen.h b/projects/miopen/include/miopen/miopen.h index cdd1ace9c2c..b7b9214905e 100644 --- a/projects/miopen/include/miopen/miopen.h +++ b/projects/miopen/include/miopen/miopen.h @@ -8528,11 +8528,11 @@ MIOPEN_EXPORT miopenStatus_t miopenMultiMarginLossForward(miopenHandle_t handle, */ typedef enum { - miopenTuningPolicyNone = 1, /*!< do not enforce anything */ - miopenTuningPolicyDbUpdate = 2, /*!< tune and update the db */ - miopenTuningPolicySearch = 3, /*!< search db first, if record not found tune but do not update the db */ - miopenTuningPolicySearchDbUpdate = 4, /*!< combination of Search and DbUpdate */ - miopenTuningPolicyDbClean = 5, /*!< remove existing entry, do not tune */ + miopenTuningPolicyNone = 1, /*!< do not enforce anything */ + miopenTuningPolicyDbUpdate = 2, /*!< tune and update the db */ + miopenTuningPolicySearch = 3, /*!< search db first, if record not found tune but do not update the db */ + miopenTuningPolicySearchDbUpdate = 4, /*!< combination of Search and DbUpdate */ + miopenTuningPolicyDbClean = 5, /*!< remove existing entry, do not tune */ } miopenTuningPolicy_t; /*! @ingroup handle diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index b192e6725e5..1e50fc5cd90 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -493,6 +493,8 @@ auto GenericSearch(const Solver s, if(skip_time < std::numeric_limits::max()) skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; + bool rec_results = perf_solsp || using_search_cutoff; + HeartBeat heartbeat; heartbeat.Start(); @@ -654,7 +656,8 @@ auto GenericSearch(const Solver s, } } } - perf_sols.push_back({current_config.ToString(), elapsed_time}); + if(rec_results) + perf_sols.push_back({current_config.ToString(), elapsed_time}); } // Banchmarked kernels will not be used anymore. From 829793c7f7e27ee936b5b6077060468e3866be5e Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 19 Nov 2025 23:20:05 +0000 Subject: [PATCH 17/26] add skip mechanism to find loop --- projects/miopen/src/conv/solver_finders.cpp | 32 ++++++++++++++++--- .../include/miopen/conv/solver_finders.hpp | 3 +- projects/miopen/src/ocl/convolutionocl.cpp | 13 ++++---- 3 files changed, 37 insertions(+), 11 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 2c008d4e8f8..c4e07929997 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -45,6 +45,9 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_SKIP_PCT, 130) + namespace miopen { namespace conv { @@ -205,8 +208,8 @@ const std::vector>& GetConvSolverFinders() tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); - tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); + tmp.emplace_back(std::make_unique()); return tmp; }(); @@ -221,13 +224,14 @@ std::vector EvaluateInvokers(const Handle& handle, const AlgorithmName& algorithm_name, const NetworkConfig& network_config, const AnyInvokeParams& invoke_ctx, - bool& is_result_optimal, + FindCoreResult& core_result, bool force_attach_binary) { const auto arch = env::value(MIOPEN_DEVICE_ARCH); if(!arch.empty()) return {}; + bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF); auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; auto best = std::numeric_limits::max(); auto best_invoker = Invoker{}; @@ -250,13 +254,27 @@ std::vector EvaluateInvokers(const Handle& handle, // That is why we do not write sub-optimal results into persistent find-db (on disk) // unless this is explicitly enabled via environment setting. if(!env::enabled(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE)) - is_result_optimal = false; + core_result.is_optimal = false; continue; } if(!sol.invoker_factory) MIOPEN_THROW("Invoker is not provided by solver " + sol.solver_id); + float skip_time = core_result.find_search_best_time; + if(skip_time < std::numeric_limits::max()) + { + skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; + // skip Naive if another solver has been timed. + if(sol.solver_id.find("Naive") != std::string::npos) + { + //s.SolverDbId() + MIOPEN_LOG_I("Skipping Naive solver: " << sol.solver_id); + continue; + } + } + MIOPEN_LOG_I("Evaluating solver: " << algorithm_name.ToString() << ":"<< sol.solver_id); + std::vector programs; const auto invoker = handle.PrepareInvoker(*sol.invoker_factory, sol.construction_params, @@ -272,6 +290,7 @@ std::vector EvaluateInvokers(const Handle& handle, auto first_elapsed = static_cast(0); int i = 0; samples.clear(); + while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX) { invoker(handle, invoke_ctx); @@ -280,6 +299,8 @@ std::vector EvaluateInvokers(const Handle& handle, if(i > 0) { samples.push_back(handle.GetKernelTime()); + if(i == 1 && using_search_cutoff && samples.front() > skip_time) + break; } else { @@ -308,6 +329,7 @@ std::vector EvaluateInvokers(const Handle& handle, best = elapsed; selected = sol; best_invoker = invoker; + core_result.find_search_best_time = best; } auto solution = Solution{solver::Id{sol.solver_id}, elapsed, sol.workspace_sz}; @@ -355,6 +377,7 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, for(auto it = solutions.begin(); it != solutions.end();) { + MIOPEN_LOG_I("testing alg: " << it->first.ToString()); if(it->second.empty()) { it = solutions.erase(it); @@ -392,12 +415,13 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, for(const auto& ss : solutions) { + MIOPEN_LOG_I("Eval alg: " << ss.first.ToString()); auto evaluated = EvaluateInvokers(handle, ss.second, ss.first, network_config, invoke_ctx, - ret.is_optimal, + ret, force_attach_binary); ret.solutions.insert(ret.solutions.end(), diff --git a/projects/miopen/src/include/miopen/conv/solver_finders.hpp b/projects/miopen/src/include/miopen/conv/solver_finders.hpp index 90fda1a238b..2e598afef4a 100644 --- a/projects/miopen/src/include/miopen/conv/solver_finders.hpp +++ b/projects/miopen/src/include/miopen/conv/solver_finders.hpp @@ -160,6 +160,7 @@ const std::vector>& GetConvSolverFinders(); struct FindCoreResult { std::vector solutions; + float find_search_best_time = std::numeric_limits::max(); bool is_optimal; }; @@ -168,7 +169,7 @@ std::vector EvaluateInvokers(const Handle& handle, const AlgorithmName& algorithm_name, const NetworkConfig& network_config, const AnyInvokeParams& invoke_ctx, - bool& is_result_optimal, + FindCoreResult& core_result, bool force_attach_binary); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index 467db5af1bc..32b1c8210a9 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -287,7 +287,8 @@ std::vector EvaluateConvSolutions(const ExecutionContext& ctx, // test timing of solver reported by system db const auto& handle = ctx.GetStream(); AutoEnableProfiling enableProfiling{handle}; - bool is_optimal = true; + FindCoreResult core_result; + core_result.is_optimal = true; // reverse solutions so that EvaluateInvokers registers the fastest solution last auto sol_itr = solutions.rbegin(); @@ -307,7 +308,7 @@ std::vector EvaluateConvSolutions(const ExecutionContext& ctx, AlgorithmName algo{ ConvolutionAlgoToDirectionalString(id.GetAlgo(), problem.GetDirection())}; std::vector eval_sol = EvaluateInvokers( - handle, conv_sols, algo, problem.MakeNetworkConfig(), invoke_ctx, is_optimal, false); + handle, conv_sols, algo, problem.MakeNetworkConfig(), invoke_ctx, core_result, false); if(!eval_sol.empty()) eval_sols.emplace_back(eval_sol.front()); @@ -377,15 +378,15 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, // add to user fdb so this check is skipped next time MIOPEN_LOG_I2("TrustVerify: Add system db entry to user db"); auto fallback = FallbackPath(); - auto ret = FindCoreResult(); - ret.is_optimal = true; + auto core_result = FindCoreResult(); + core_result.is_optimal = true; auto copy_sols = conv.GetSolutions(ctx, problem, 4, &fallback, &invoke_ctx); for(const auto& s : copy_sols) { auto solution = Solution{solver::Id{s.solution_id}, s.time, s.workspace_size}; - ret.solutions.emplace_back(std::move(solution)); + core_result.solutions.emplace_back(std::move(solution)); } - return ret; + return core_result; } else { From 4cae3cbd49127a0d1155b591e9be2ed5d764e4e2 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 20 Nov 2025 21:20:54 +0000 Subject: [PATCH 18/26] order solver finders by performance --- projects/miopen/src/conv/solver_finders.cpp | 39 ++++++++++----------- projects/miopen/src/ocl/convolutionocl.cpp | 6 ++-- 2 files changed, 21 insertions(+), 24 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index c4e07929997..acf0af5d6a9 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -232,10 +232,10 @@ std::vector EvaluateInvokers(const Handle& handle, return {}; bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF); - auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; - auto best = std::numeric_limits::max(); - auto best_invoker = Invoker{}; - auto ret = std::vector{}; + auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; + auto best = std::numeric_limits::max(); + auto best_invoker = Invoker{}; + auto ret = std::vector{}; std::vector samples; for(const auto& sol : solutions) @@ -266,14 +266,14 @@ std::vector EvaluateInvokers(const Handle& handle, { skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; // skip Naive if another solver has been timed. - if(sol.solver_id.find("Naive") != std::string::npos) + if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos) { - //s.SolverDbId() - MIOPEN_LOG_I("Skipping Naive solver: " << sol.solver_id); + MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); continue; } } - MIOPEN_LOG_I("Evaluating solver: " << algorithm_name.ToString() << ":"<< sol.solver_id); + MIOPEN_LOG_I("Evaluating Solver: " << algorithm_name.ToString() << ":" << sol.solver_id); std::vector programs; const auto invoker = handle.PrepareInvoker(*sol.invoker_factory, @@ -300,7 +300,11 @@ std::vector EvaluateInvokers(const Handle& handle, { samples.push_back(handle.GetKernelTime()); if(i == 1 && using_search_cutoff && samples.front() > skip_time) + { + MIOPEN_LOG_I("Skipping (Slow) Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); break; + } } else { @@ -326,9 +330,9 @@ std::vector EvaluateInvokers(const Handle& handle, MIOPEN_LOG_I(sol << ": " << elapsed << (elapsed < best ? " < " : " >= ") << best); if(elapsed < best) { - best = elapsed; - selected = sol; - best_invoker = invoker; + best = elapsed; + selected = sol; + best_invoker = invoker; core_result.find_search_best_time = best; } @@ -366,7 +370,7 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, auto& handle = ctx.GetStream(); // Find - auto solutions = std::map>{}; + auto solutions = std::vector>>{}; std::transform( finders.begin(), finders.end(), std::inserter(solutions, solutions.end()), [&](auto&& f) { return std::make_pair(f->GetAlgorithmName(problem), @@ -377,7 +381,6 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, for(auto it = solutions.begin(); it != solutions.end();) { - MIOPEN_LOG_I("testing alg: " << it->first.ToString()); if(it->second.empty()) { it = solutions.erase(it); @@ -415,14 +418,8 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, for(const auto& ss : solutions) { - MIOPEN_LOG_I("Eval alg: " << ss.first.ToString()); - auto evaluated = EvaluateInvokers(handle, - ss.second, - ss.first, - network_config, - invoke_ctx, - ret, - force_attach_binary); + auto evaluated = EvaluateInvokers( + handle, ss.second, ss.first, network_config, invoke_ctx, ret, force_attach_binary); ret.solutions.insert(ret.solutions.end(), std::make_move_iterator(evaluated.begin()), diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index 32b1c8210a9..2fcbe3b6c8f 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -377,10 +377,10 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, // system db result is good // add to user fdb so this check is skipped next time MIOPEN_LOG_I2("TrustVerify: Add system db entry to user db"); - auto fallback = FallbackPath(); - auto core_result = FindCoreResult(); + auto fallback = FallbackPath(); + auto core_result = FindCoreResult(); core_result.is_optimal = true; - auto copy_sols = conv.GetSolutions(ctx, problem, 4, &fallback, &invoke_ctx); + auto copy_sols = conv.GetSolutions(ctx, problem, 4, &fallback, &invoke_ctx); for(const auto& s : copy_sols) { auto solution = Solution{solver::Id{s.solution_id}, s.time, s.workspace_size}; From 330cc6c10d5f9bc6bedc40cae00a3c2405b2505b Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 20 Nov 2025 21:56:53 +0000 Subject: [PATCH 19/26] Search cutoff option for solver finders --- projects/miopen/src/conv/solver_finders.cpp | 59 +++++++++++++------ .../include/miopen/conv/solver_finders.hpp | 3 +- projects/miopen/src/ocl/convolutionocl.cpp | 17 +++--- 3 files changed, 51 insertions(+), 28 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 0dd79302d18..386418e2c0d 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -45,6 +45,9 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_SEARCH_SKIP_PCT, 130) + namespace miopen { namespace conv { @@ -202,11 +205,11 @@ const std::vector>& GetConvSolverFinders() { static const auto finders = []() { auto tmp = std::vector>{}; - tmp.emplace_back(std::make_unique()); - tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); + tmp.emplace_back(std::make_unique()); tmp.emplace_back(std::make_unique()); + tmp.emplace_back(std::make_unique()); return tmp; }(); @@ -221,17 +224,18 @@ std::vector EvaluateInvokers(const Handle& handle, const AlgorithmName& algorithm_name, const NetworkConfig& network_config, const AnyInvokeParams& invoke_ctx, - bool& is_result_optimal, + FindCoreResult& core_result, bool force_attach_binary) { const auto arch = env::value(MIOPEN_DEVICE_ARCH); if(!arch.empty()) return {}; - auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; - auto best = std::numeric_limits::max(); - auto best_invoker = Invoker{}; - auto ret = std::vector{}; + bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF); + auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; + auto best = std::numeric_limits::max(); + auto best_invoker = Invoker{}; + auto ret = std::vector{}; std::vector samples; for(const auto& sol : solutions) @@ -250,13 +254,27 @@ std::vector EvaluateInvokers(const Handle& handle, // That is why we do not write sub-optimal results into persistent find-db (on disk) // unless this is explicitly enabled via environment setting. if(!env::enabled(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE)) - is_result_optimal = false; + core_result.is_optimal = false; continue; } if(!sol.invoker_factory) MIOPEN_THROW("Invoker is not provided by solver " + sol.solver_id); + float skip_time = core_result.find_search_best_time; + if(skip_time < std::numeric_limits::max()) + { + skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; + // skip Naive if another solver has been timed. + if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos) + { + MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); + continue; + } + } + MIOPEN_LOG_I("Evaluating Solver: " << algorithm_name.ToString() << ":" << sol.solver_id); + std::vector programs; const auto invoker = handle.PrepareInvoker(*sol.invoker_factory, sol.construction_params, @@ -272,6 +290,7 @@ std::vector EvaluateInvokers(const Handle& handle, auto first_elapsed = static_cast(0); int i = 0; samples.clear(); + while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX) { invoker(handle, invoke_ctx); @@ -280,6 +299,12 @@ std::vector EvaluateInvokers(const Handle& handle, if(i > 0) { samples.push_back(handle.GetKernelTime()); + if(i == 1 && using_search_cutoff && samples.front() > skip_time) + { + MIOPEN_LOG_I("Skipping (Slow) Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); + break; + } } else { @@ -306,9 +331,10 @@ std::vector EvaluateInvokers(const Handle& handle, << (elapsed < best ? " < " : " >= ") << best); if(elapsed < best) { - best = elapsed; - selected = sol; - best_invoker = invoker; + best = elapsed; + selected = sol; + best_invoker = invoker; + core_result.find_search_best_time = best; } auto solution = Solution{solver::Id{sol.solver_id}, elapsed, sol.workspace_sz}; @@ -345,7 +371,7 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, auto& handle = ctx.GetStream(); // Find - auto solutions = std::map>{}; + auto solutions = std::vector>>{}; std::transform( finders.begin(), finders.end(), std::inserter(solutions, solutions.end()), [&](auto&& f) { return std::make_pair(f->GetAlgorithmName(problem), @@ -393,13 +419,8 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, for(const auto& ss : solutions) { - auto evaluated = EvaluateInvokers(handle, - ss.second, - ss.first, - network_config, - invoke_ctx, - ret.is_optimal, - force_attach_binary); + auto evaluated = EvaluateInvokers( + handle, ss.second, ss.first, network_config, invoke_ctx, ret, force_attach_binary); ret.solutions.insert(ret.solutions.end(), std::make_move_iterator(evaluated.begin()), diff --git a/projects/miopen/src/include/miopen/conv/solver_finders.hpp b/projects/miopen/src/include/miopen/conv/solver_finders.hpp index 90fda1a238b..2e598afef4a 100644 --- a/projects/miopen/src/include/miopen/conv/solver_finders.hpp +++ b/projects/miopen/src/include/miopen/conv/solver_finders.hpp @@ -160,6 +160,7 @@ const std::vector>& GetConvSolverFinders(); struct FindCoreResult { std::vector solutions; + float find_search_best_time = std::numeric_limits::max(); bool is_optimal; }; @@ -168,7 +169,7 @@ std::vector EvaluateInvokers(const Handle& handle, const AlgorithmName& algorithm_name, const NetworkConfig& network_config, const AnyInvokeParams& invoke_ctx, - bool& is_result_optimal, + FindCoreResult& core_result, bool force_attach_binary); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index aa208486a7c..1df3c411356 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -288,7 +288,8 @@ std::vector EvaluateConvSolutions(const ExecutionContext& ctx, // test timing of solver reported by system db const auto& handle = ctx.GetStream(); AutoEnableProfiling enableProfiling{handle}; - bool is_optimal = true; + FindCoreResult core_result; + core_result.is_optimal = true; // reverse solutions so that EvaluateInvokers registers the fastest solution last auto sol_itr = solutions.rbegin(); @@ -308,7 +309,7 @@ std::vector EvaluateConvSolutions(const ExecutionContext& ctx, AlgorithmName algo{ ConvolutionAlgoToDirectionalString(id.GetAlgo(), problem.GetDirection())}; std::vector eval_sol = EvaluateInvokers( - handle, conv_sols, algo, problem.MakeNetworkConfig(), invoke_ctx, is_optimal, false); + handle, conv_sols, algo, problem.MakeNetworkConfig(), invoke_ctx, core_result, false); if(!eval_sol.empty()) eval_sols.emplace_back(eval_sol.front()); @@ -377,16 +378,16 @@ std::vector VerifiedFDBSolution(const ExecutionContext& ctx, // system db result is good // add to user fdb so this check is skipped next time MIOPEN_LOG_I2("TrustVerify: Add system db entry to user db"); - auto fallback = FallbackPath(); - auto ret = FindCoreResult(); - ret.is_optimal = true; - auto copy_sols = conv.GetSolutions(ctx, problem, 4, &fallback, &invoke_ctx); + auto fallback = FallbackPath(); + auto core_result = FindCoreResult(); + core_result.is_optimal = true; + auto copy_sols = conv.GetSolutions(ctx, problem, 4, &fallback, &invoke_ctx); for(const auto& s : copy_sols) { auto solution = Solution{solver::Id{s.solution_id}, s.time, s.workspace_size}; - ret.solutions.emplace_back(std::move(solution)); + core_result.solutions.emplace_back(std::move(solution)); } - return ret; + return core_result; } else { From 01c280c8e1c218bfb8073be1f011bfb6fa0d5f92 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 24 Nov 2025 20:40:10 +0000 Subject: [PATCH 20/26] allow naives to be used when avg kernel time < 50ns --- projects/miopen/src/conv/solver_finders.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 386418e2c0d..c389162d8e1 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -264,14 +264,14 @@ std::vector EvaluateInvokers(const Handle& handle, float skip_time = core_result.find_search_best_time; if(skip_time < std::numeric_limits::max()) { - skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; - // skip Naive if another solver has been timed. - if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos) + // skip Naive if another solver has been timed and solution took more than 50ns. + if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos && skip_time > 0.05f) { MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" << sol.solver_id); continue; } + skip_time *= env::value(MIOPEN_SEARCH_SKIP_PCT) / 100.0f; } MIOPEN_LOG_I("Evaluating Solver: " << algorithm_name.ToString() << ":" << sol.solver_id); From 07d595a96f0ef08e82d7b972bfe7925af6c147f2 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 26 Nov 2025 10:51:19 -0600 Subject: [PATCH 21/26] add doc reference to MIOPEN_SEARCH_CUTOFF --- projects/miopen/docs/reference/env_variables.rst | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index a1ee2f250a8..1c8e51f1a37 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -101,6 +101,11 @@ and :doc:`Performance database <../conceptual/perfdb>`. | "SEARCH_DB_UPDATE" or 4: Combination of DB_UPDATE and SEARCH (unsafe with Fast/Hybrid/Trust modes) | "DB_CLEAN" or 5: Remove optimized values from User PerfDb (unsafe with Fast/Hybrid/Trust modes) + * - | ``MIOPEN_SEARCH_CUTOFF`` + | Allows speculative early termination of suboptimal searches. + - | 1: Enable + | 0 or unset: Disable + * - | ``MIOPEN_DEBUG_DISABLE_FIND_DB`` | Disables FindDb functionality. - | 1: Disable FindDb From 315e14aa345a91136513e0607e6ce57b9b5868c3 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 26 Nov 2025 20:23:31 +0000 Subject: [PATCH 22/26] format --- projects/miopen/src/conv/solver_finders.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index e81aed6a564..e1f66e2d407 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -265,7 +265,8 @@ std::vector EvaluateInvokers(const Handle& handle, if(skip_time < std::numeric_limits::max()) { // skip Naive if another solver has been timed and solution took more than 50ns. - if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos && skip_time > 0.05f) + if(using_search_cutoff && sol.solver_id.find("Naive") != std::string::npos && + skip_time > 0.05f) { MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" << sol.solver_id); From 30046e1589cc46dde9254e0b7631794d510ff354 Mon Sep 17 00:00:00 2001 From: Jonathan Lichtner <195780826+JonathanLichtnerAMD@users.noreply.github.com> Date: Fri, 28 Nov 2025 09:47:58 -0600 Subject: [PATCH 23/26] Change upcoming branch from 8.0.0 to 7.11.0 --- projects/miopen/CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index 111a2b28817..bcf70315396 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -2,7 +2,7 @@ # Change Log for MIOpen Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/projects/MIOpen/en/latest/) -## (Unreleased) MIOpen 3.5.1 for ROCm 8.0.0 +## (Unreleased) MIOpen 3.5.1 for ROCm 7.11.0 ## MIOpen 3.5.1 for ROCm 7.2.0 ### Changed From ee013692510770c189ece29954b8f8880e3c4156 Mon Sep 17 00:00:00 2001 From: Jonathan Lichtner <195780826+JonathanLichtnerAMD@users.noreply.github.com> Date: Fri, 28 Nov 2025 09:52:55 -0600 Subject: [PATCH 24/26] Update the changelog for this new tuning cutoff option --- projects/miopen/CHANGELOG.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index bcf70315396..c1f20273ac1 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -3,6 +3,8 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/projects/MIOpen/en/latest/) ## (Unreleased) MIOpen 3.5.1 for ROCm 7.11.0 +### Optimized +* Added `MIOPEN_SEARCH_CUTOFF` option which can reduce tuning times by skipping slow solvers and kernels ## MIOpen 3.5.1 for ROCm 7.2.0 ### Changed From 873632c09c1b41b720efb7b52d7a17c7f3fd39e6 Mon Sep 17 00:00:00 2001 From: NolanHannaAMD Date: Mon, 1 Dec 2025 16:47:48 +0000 Subject: [PATCH 25/26] Fixed an issue where an error would be displayed when generic search was either cutoff or skipped all kernels. Also adjusted to tense of some log messages to be more consistent. --- projects/miopen/src/include/miopen/find_solution.hpp | 4 ++-- projects/miopen/src/include/miopen/generic_search.hpp | 10 ++++++++++ 2 files changed, 12 insertions(+), 2 deletions(-) diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index f6429d1123f..551c06ab1b1 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -142,10 +142,10 @@ auto FindSolutionImpl(rank<1>, { auto record = DbRecord(DbKinds::PerfDb, problem); if(env::enabled(MIOPEN_WARN_SEARCH)) - MIOPEN_LOG_W("Search Start: " << record.GetKey() << " : " << s.SolverDbId() + MIOPEN_LOG_W("Search Started: " << record.GetKey() << " : " << s.SolverDbId() << ", enforce: " << enforce); else - MIOPEN_LOG_I("Search Start: " << record.GetKey() << " : " << s.SolverDbId() + MIOPEN_LOG_I("Search Started: " << record.GetKey() << " : " << s.SolverDbId() << ", enforce: " << enforce); try { diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 1e50fc5cd90..9077621382e 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -434,6 +434,7 @@ auto GenericSearch(const Solver s, using PerformanceConfig = decltype(s.GetDefaultPerformanceConfig(context, problem)); PerformanceConfig best_config; + PerformanceConfig last_config; // Used in cases where all kernels were intentionally skipped const auto default_solution = s.GetSolution(context, problem, s.GetDefaultPerformanceConfig(context, problem)); const auto invoke_ctx = [invoke_ctx_]() { @@ -540,6 +541,8 @@ auto GenericSearch(const Solver s, auto current_config = std::get<0>(kinder); auto current_solution = std::get<1>(kinder); + last_config = current_config; + if(std::get<2>(kinder)) { threads_remaining--; @@ -694,6 +697,13 @@ auto GenericSearch(const Solver s, MIOPEN_LOG_I("Done: " << n_runs_total << '/' << n_failed << '/' << n_runs_total << ", best #" << n_best << ' ' << best_time << ' ' << best_config); + // If no errors were encountered, but we either cutoff or skipped every kernel, don't throw. + if(!is_passed && n_failed == 0) + { + MIOPEN_LOG_I("Search cutoff or skipped for all kernels. Last config returned: " << last_config); + return last_config; + } + if(!is_passed) MIOPEN_THROW("Search failed"); From fccb9d601af293ff9e5a7108369cccbb91cfebed Mon Sep 17 00:00:00 2001 From: NolanHannaAMD Date: Mon, 1 Dec 2025 21:24:00 +0000 Subject: [PATCH 26/26] Fixing some clang-format errors. --- projects/miopen/src/include/miopen/find_solution.hpp | 4 ++-- projects/miopen/src/include/miopen/generic_search.hpp | 5 +++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index 551c06ab1b1..df794a03a8a 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -143,10 +143,10 @@ auto FindSolutionImpl(rank<1>, auto record = DbRecord(DbKinds::PerfDb, problem); if(env::enabled(MIOPEN_WARN_SEARCH)) MIOPEN_LOG_W("Search Started: " << record.GetKey() << " : " << s.SolverDbId() - << ", enforce: " << enforce); + << ", enforce: " << enforce); else MIOPEN_LOG_I("Search Started: " << record.GetKey() << " : " << s.SolverDbId() - << ", enforce: " << enforce); + << ", enforce: " << enforce); try { auto c = s.Search(context, problem, invoke_ctx); diff --git a/projects/miopen/src/include/miopen/generic_search.hpp b/projects/miopen/src/include/miopen/generic_search.hpp index 9077621382e..182472693de 100644 --- a/projects/miopen/src/include/miopen/generic_search.hpp +++ b/projects/miopen/src/include/miopen/generic_search.hpp @@ -700,10 +700,11 @@ auto GenericSearch(const Solver s, // If no errors were encountered, but we either cutoff or skipped every kernel, don't throw. if(!is_passed && n_failed == 0) { - MIOPEN_LOG_I("Search cutoff or skipped for all kernels. Last config returned: " << last_config); + MIOPEN_LOG_I( + "Search cutoff or skipped for all kernels. Last config returned: " << last_config); return last_config; } - + if(!is_passed) MIOPEN_THROW("Search failed");