Skip to content

[MIOpen] Cutoff perf for GenericSearch #1536

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
NolanHannaAMD merged 48 commits into from
Dec 8, 2025
Merged

[MIOpen] Cutoff perf for GenericSearch #1536

Show file tree
Hide file tree
Changes from 10 commits
Commits
Show all changes
48 commits
Select commit Hold shift + click to select a range
6115f17
add cutoff time for GenericSearch relative to worst case timing from …
cderb Sep 4, 2025
02a92b9
remove const from execution context to allow message passing between …
cderb Sep 5, 2025
ca3384b
removing const modifiers
cderb Sep 5, 2025
a303d23
use context copy for ExecutePrimitive
cderb Sep 8, 2025
4ab9fd1
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Sep 8, 2025
f8e3c30
whitespace
cderb Sep 9, 2025
bbba8f6
adjust function signatures, corrections for cutoff updates
cderb Sep 9, 2025
233cdcc
clang format
cderb Sep 9, 2025
770f56f
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Sep 9, 2025
1352364
split cutoff time to best_time and worst_time
cderb Sep 10, 2025
2d449a4
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Sep 11, 2025
a953883
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Sep 11, 2025
2362fd6
restore const to ExecutionContext, add mutable to search time fields
cderb Sep 11, 2025
a6ea744
add envs to modify GenericSearch skips
cderb Sep 11, 2025
5c03cc8
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Sep 16, 2025
7f78659
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Oct 6, 2025
a72cef3
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Oct 14, 2025
2f548b2
format
cderb Oct 14, 2025
9ef9664
format
cderb Oct 14, 2025
0507b84
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Oct 14, 2025
179b0d6
adjust tuning policy smoke test
cderb Oct 15, 2025
dc3b189
format
cderb Oct 17, 2025
7e23c0a
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Oct 27, 2025
e981caf
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Nov 3, 2025
950c6aa
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Nov 3, 2025
cc29aa9
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Nov 4, 2025
5121522
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Nov 6, 2025
8795a3a
switch to using env MIOPEN_SEARCH_CUTOFF
cderb Nov 7, 2025
b5ae9ab
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Nov 7, 2025
f6fceac
stop populating perf_sols when unused
cderb Nov 8, 2025
8477e12
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Nov 8, 2025
3891835
Merge branch 'develop' into users/cderb/tuning_cutoff
cderb Nov 10, 2025
6d06795
Merge branch 'develop' into users/cderb/tuning_cutoff
NolanHannaAMD Nov 19, 2025
829793c
add skip mechanism to find loop
cderb Nov 19, 2025
4cae3cb
order solver finders by performance
cderb Nov 20, 2025
c953669
Merge remote-tracking branch 'origin/develop' into users/cderb/find_c…
cderb Nov 20, 2025
330cc6c
Search cutoff option for solver finders
cderb Nov 20, 2025
01c280c
allow naives to be used when avg kernel time < 50ns
cderb Nov 24, 2025
ecffb43
Merge remote-tracking branch 'origin/users/cderb/find_cutoff2' into u…
cderb Nov 25, 2025
f881198
Merge branch 'users/cderb/find_cutoff' into users/cderb/tuning_cutoff
cderb Nov 25, 2025
3842a09
Merge remote-tracking branch 'origin/develop' into users/cderb/tuning…
cderb Nov 25, 2025
07d595a
add doc reference to MIOPEN_SEARCH_CUTOFF
cderb Nov 26, 2025
315e14a
format
cderb Nov 26, 2025
8fd6c42
Merge branch 'develop' into users/cderb/tuning_cutoff
JonathanLichtnerAMD Nov 28, 2025
30046e1
Change upcoming branch from 8.0.0 to 7.11.0
JonathanLichtnerAMD Nov 28, 2025
ee01369
Update the changelog for this new tuning cutoff option
JonathanLichtnerAMD Nov 28, 2025
873632c
Fixed an issue where an error would be displayed when generic search …
NolanHannaAMD Dec 1, 2025
fccb9d6
Fixing some clang-format errors.
NolanHannaAMD Dec 1, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 5 additions & 3 deletions projects/miopen/include/miopen/miopen.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
*/
Expand All @@ -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
Expand Down
16 changes: 8 additions & 8 deletions projects/miopen/src/conv/solver_finders.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,7 +66,7 @@ class DirectSolverFinder : public SolversFinderMixin<ProblemDescription, ConvFin
return !parameters.use_winograd_only && !env::disabled(MIOPEN_DEBUG_CONV_DIRECT);
}

std::vector<solver::ConvSolution> FindImpl(const ExecutionContext& ctx,
std::vector<solver::ConvSolution> FindImpl(ExecutionContext& ctx,
const ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx,
const ConvFindParameters&,
Expand Down Expand Up @@ -95,7 +95,7 @@ class ImplicitGemmSolverFinder : public SolversFinderMixin<ProblemDescription, C
return !parameters.use_winograd_only && !env::disabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM);
}

std::vector<solver::ConvSolution> FindImpl(const ExecutionContext& ctx,
std::vector<solver::ConvSolution> FindImpl(ExecutionContext& ctx,
const ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx,
const ConvFindParameters&,
Expand Down Expand Up @@ -126,7 +126,7 @@ class FftSolverFinder : public SolversFinderMixin<ProblemDescription, ConvFindPa
!env::disabled(MIOPEN_DEBUG_CONV_FFT);
}

std::vector<solver::ConvSolution> FindImpl(const ExecutionContext& ctx,
std::vector<solver::ConvSolution> FindImpl(ExecutionContext& ctx,
const ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx,
const ConvFindParameters&,
Expand All @@ -153,7 +153,7 @@ class GemmSolverFinder : public SolversFinderMixin<ProblemDescription, ConvFindP
return !parameters.use_winograd_only && !env::disabled(MIOPEN_DEBUG_CONV_GEMM);
}

std::vector<solver::ConvSolution> FindImpl(const ExecutionContext& ctx,
std::vector<solver::ConvSolution> FindImpl(ExecutionContext& ctx,
const ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx,
const ConvFindParameters&,
Expand All @@ -180,7 +180,7 @@ class WinogradSolverFinder : public SolversFinderMixin<ProblemDescription, ConvF
return !env::disabled(MIOPEN_DEBUG_CONV_WINOGRAD);
}

std::vector<solver::ConvSolution> FindImpl(const ExecutionContext& ctx,
std::vector<solver::ConvSolution> FindImpl(ExecutionContext& ctx,
const ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx,
const ConvFindParameters& parameters,
Expand All @@ -202,10 +202,10 @@ const std::vector<std::unique_ptr<ISolversFinder>>& GetConvSolverFinders()
{
static const auto finders = []() {
auto tmp = std::vector<std::unique_ptr<ISolversFinder>>{};
tmp.emplace_back(std::make_unique<WinogradSolverFinder>());
tmp.emplace_back(std::make_unique<DirectSolverFinder>());
tmp.emplace_back(std::make_unique<ImplicitGemmSolverFinder>());
tmp.emplace_back(std::make_unique<GemmSolverFinder>());
tmp.emplace_back(std::make_unique<WinogradSolverFinder>());
tmp.emplace_back(std::make_unique<DirectSolverFinder>());
tmp.emplace_back(std::make_unique<FftSolverFinder>());
Comment thread
JonathanLichtnerAMD marked this conversation as resolved.
return tmp;
}();
Expand Down Expand Up @@ -334,7 +334,7 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
}

FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx,
const ExecutionContext& ctx,
ExecutionContext& ctx,
const ProblemDescriptionBase& problem,
const PrimitiveFindParameters& parameters,
const std::vector<std::unique_ptr<ISolversFinder>>& finders,
Expand Down
7 changes: 7 additions & 0 deletions projects/miopen/src/find_controls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,8 @@ static_assert(FindEnforceAction::Search ==
static_cast<FindEnforceAction>(miopenTuningPolicySearch));
static_assert(FindEnforceAction::SearchDbUpdate ==
static_cast<FindEnforceAction>(miopenTuningPolicySearchDbUpdate));
static_assert(FindEnforceAction::SearchCutoffDbUpdate ==
static_cast<FindEnforceAction>(miopenTuningPolicySearchCutoffDbUpdate));
static_assert(FindEnforceAction::DbClean ==
static_cast<FindEnforceAction>(miopenTuningPolicyDbClean));

Expand All @@ -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 "<Unknown>";
Expand Down Expand Up @@ -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;
Expand Down
10 changes: 5 additions & 5 deletions projects/miopen/src/fusion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<std::string>& perf_cfg_override,
const FusionDescription& problem,
Expand Down Expand Up @@ -829,13 +829,13 @@ class FusionSolverFinder : public SolversFinderMixin<FusionDescription, FusionFi
}

std::vector<solver::ConvSolution>
FindImpl(const ExecutionContext& ctx,
FindImpl(ExecutionContext& ctx,
const FusionDescription& problem,
const AnyInvokeParams& invoke_ctx,
const FusionFindParameters&,
const std::optional<FindOptions>& options) const override
{
const auto fusion_ctx = FusionContext(ctx);
auto fusion_ctx = FusionContext(ctx);
return solvers.SearchForAllSolutions(fusion_ctx,
problem,
MakeConvDbGetter(ctx),
Expand Down Expand Up @@ -867,7 +867,7 @@ static const std::vector<std::unique_ptr<ISolversFinder>>& GetFusionSolverFinder
}

static std::vector<Solution>
FindFusion(const ExecutionContext& ctx,
FindFusion(ExecutionContext& ctx,
const FusionDescription& fusion_problem,
const std::function<fusion::FusionInvokeParams()>& invoke_params,
const std::optional<FindOptions>& options = std::nullopt)
Expand Down Expand Up @@ -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
Expand Down
47 changes: 24 additions & 23 deletions projects/miopen/src/include/miopen/any_solver.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -109,7 +109,7 @@ struct AnySolver
}
#endif

ConvSolution FindSolution(const ExecutionContext& ctx,
ConvSolution FindSolution(ExecutionContext& ctx,
const miopen::conv::ProblemDescription& problem,
std::function<PerformanceDb&()>& db_getter,
const miopen::AnyInvokeParams& invoke_ctx,
Expand All @@ -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,
Expand All @@ -129,13 +129,13 @@ 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<miopen::solver::SolutionPerf>* perf_sols = nullptr) const
std::vector<miopen::solver::SolutionPerf>* 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,
Expand Down Expand Up @@ -194,26 +194,26 @@ 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<PerformanceDb&()>& 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,
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<miopen::solver::SolutionPerf>* perf_sols) const = 0;
std::vector<miopen::solver::SolutionPerf>* perf_solsp) const = 0;
virtual InvokerFactory GetInvokeFactory(const ExecutionContext& ctx,
const miopen::conv::ProblemDescription& problem,
const std::string& perf_cfg) const = 0;
Expand Down Expand Up @@ -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,
Expand All @@ -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<PerformanceDb&()>& db_getter,
const miopen::AnyInvokeParams& invoke_ctx,
Expand All @@ -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,
Expand All @@ -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<miopen::solver::SolutionPerf>*,
Expand All @@ -454,19 +454,20 @@ 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<miopen::solver::SolutionPerf>* perf_sols,
std::vector<miopen::solver::SolutionPerf>* 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();
}

// non-tunable solver has no search
std::string GenericSearch(const ExecutionContext&,
std::string GenericSearch(ExecutionContext&,
const miopen::conv::ProblemDescription&,
const miopen::AnyInvokeParams&,
std::vector<miopen::solver::SolutionPerf>*,
Expand All @@ -476,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<miopen::solver::SolutionPerf>*,
Expand All @@ -487,15 +488,15 @@ struct AnySolver
}

std::string
GenericSearch(const ExecutionContext& ctx,
GenericSearch(ExecutionContext& ctx,
const miopen::conv::ProblemDescription& problem,
const miopen::AnyInvokeParams& invoke_ctx,
std::vector<miopen::solver::SolutionPerf>* perf_sols) const override
std::vector<miopen::solver::SolutionPerf>* perf_solsp) const override
{
return GenericSearch(ctx,
problem,
invoke_ctx,
perf_sols,
perf_solsp,
std::integral_constant<bool, TunableSolver::Is>(),
std::integral_constant<bool, LegacySolver::Is>());
}
Expand Down
4 changes: 2 additions & 2 deletions projects/miopen/src/include/miopen/batchnorm/solvers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -130,7 +130,7 @@ struct BnBwdTrainingSpatial final : BatchNormTunableSolver<PerformanceConfigBnBw
const PerformanceConfigBnBwdBackward& config) const override;

MIOPEN_INTERNALS_EXPORT PerformanceConfigBnBwdBackward
Search(const ExecutionContext& ctx,
Search(ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx) const override;

Expand All @@ -157,7 +157,7 @@ struct BnFwdTrainingSpatial final : BatchNormTunableSolver<PerformanceConfigBnFw
const PerformanceConfigBnFwdTraining& config) const override;

MIOPEN_INTERNALS_EXPORT PerformanceConfigBnFwdTraining
Search(const ExecutionContext& ctx,
Search(ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx) const override;

Expand Down
Loading
Loading