Skip to content
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

Jump to bottom

[tuning] Interface for Fin #3330

Draft
wants to merge 34 commits into
base: develop
Choose a base branch
from
+1,567 −121
Draft

[tuning] Interface for Fin #3330

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
34 commits
Select commit Hold shift + click to select a range
bf1cc6d
Add interface for Fin
averinevg Oct 23, 2024
ec2ef18
Fix formatting
averinevg Oct 23, 2024
2b89653
Merge branch 'develop' into ea_interface_for_fin
averinevg Oct 23, 2024
be07add
Fix the description
averinevg Oct 25, 2024
655bb36
Merge branch 'develop' into ea_interface_for_fin
averinevg Oct 25, 2024
3652303
Fix formatting
averinevg Oct 25, 2024
d89c0ce
Fix tidy
averinevg Oct 25, 2024
ee46f04
Fix windows build
averinevg Oct 25, 2024
a56b391
Refactor
averinevg Oct 25, 2024
3a89c66
Fix formatting
averinevg Oct 25, 2024
d0d4463
Merge branch 'develop' into ea_interface_for_fin
averinevg Oct 25, 2024
ee2196f
Fix tidy
averinevg Oct 25, 2024
2a9b9a7
Fix windows build
averinevg Oct 25, 2024
c9abd64
Merge branch 'develop' into ea_interface_for_fin
averinevg Oct 28, 2024
7b47b41
Fix tidy
averinevg Oct 28, 2024
80261ed
Fix windows build
averinevg Oct 28, 2024
a8a432e
Refactor, add GetConvSolvers & GetBatchNormSolvers
averinevg Oct 28, 2024
f57bbac
Fix formatting
averinevg Oct 28, 2024
81ffac5
Fix tidy
averinevg Oct 29, 2024
230acd5
Add test for GetSolvers()
averinevg Oct 29, 2024
0f9bdf2
Refactor
averinevg Oct 29, 2024
61bfaea
Fix linker error
averinevg Oct 29, 2024
8fa4325
Fix formatting
averinevg Oct 29, 2024
da58038
Merge branch 'develop' into ea_interface_for_fin
averinevg Oct 30, 2024
c7525a9
Add test for IsApplicable() & GetWorkspaceSize()
averinevg Oct 30, 2024
6da9a98
Fix formatting
averinevg Oct 30, 2024
9a3591e
Fix windows build
averinevg Oct 31, 2024
41afab4
Merge branch 'develop' into ea_interface_for_fin
averinevg Nov 1, 2024
279d459
Implement methods
averinevg Nov 13, 2024
0ff5960
Fix formatting
averinevg Nov 13, 2024
8fffbc7
Merge branch 'develop' into ea_interface_for_fin
averinevg Nov 13, 2024
9109326
Merge branch 'develop' into ea_interface_for_fin
averinevg Nov 14, 2024
f67b5b5
Add tests for GetAllSolutions(), GetPerfCfgParams() & TestPerfCfgPara…
averinevg Nov 14, 2024
b315df9
Registry: remove smart pointer
averinevg Nov 14, 2024
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
1 change: 1 addition & 0 deletions src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,6 +120,7 @@ set( MIOpen_Source
env.cpp
execution_context.cpp
expanduser.cpp
fin/fin_interface.cpp
find_controls.cpp
find_db.cpp
fused_api.cpp
Expand Down
584 changes: 584 additions & 0 deletions src/fin/fin_interface.cpp
View file
Open in desktop

Large diffs are not rendered by default.

70 changes: 34 additions & 36 deletions src/include/miopen/batchnorm/solvers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,6 +133,7 @@ struct BnFwdInference final : BatchnormSolver

bool IsApplicable(const ExecutionContext& context,
const miopen::batchnorm::ProblemDescription& problem) const override;
bool IsDynamic() const override { return true; }
ConvSolution GetSolution(const ExecutionContext& context,
const miopen::batchnorm::ProblemDescription& problem) const override;
};
Expand All @@ -149,13 +150,11 @@ struct PerformanceConfigBnCKFwdInference : PerfConfigBaseCK<PerformanceConfigBnC
PerformanceConfigBnCKFwdInference() : PerformanceConfigBnCKFwdInference(0, "") {}
PerformanceConfigBnCKFwdInference(bool) : PerformanceConfigBnCKFwdInference(0, "") {}
MIOPEN_INTERNALS_EXPORT void
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem_desc);
MIOPEN_INTERNALS_EXPORT bool
SetNextValue(const miopen::batchnorm::ProblemDescription& problem_desc);
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool SetNextValue(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool IsValidValue() const;
MIOPEN_INTERNALS_EXPORT bool
IsValid(const ExecutionContext&,
const miopen::batchnorm::ProblemDescription& problem_desc) const;
IsValid(const ExecutionContext&, const miopen::batchnorm::ProblemDescription& problem) const;

template <typename Self, typename F>
static void Visit(Self&& s, F f)
Expand Down Expand Up @@ -185,23 +184,24 @@ struct BnCKFwdInference final : BatchNormTunableSolver<PerformanceConfigBnCKFwdI
{
const std::string& SolverDbId() const override { return GetSolverDbId<BnCKFwdInference>(); }

MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem) const override;
bool IsDynamic() const override { return true; }
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKFwdInference GetDefaultPerformanceConfig(
const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
const miopen::batchnorm::ProblemDescription& problem) const override;
MIOPEN_INTERNALS_EXPORT bool
IsValidPerformanceConfig(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKFwdInference& config) const override;
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKFwdInference
Search(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx) const override;
MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
MIOPEN_INTERNALS_EXPORT ConvSolution
GetSolution(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKFwdInference& config) const override;
};

Expand All @@ -217,13 +217,11 @@ struct PerformanceConfigBnCKBwdBackward : PerfConfigBaseCK<PerformanceConfigBnCK
PerformanceConfigBnCKBwdBackward() : PerformanceConfigBnCKBwdBackward(0, "") {}
PerformanceConfigBnCKBwdBackward(bool) : PerformanceConfigBnCKBwdBackward(0, "") {}
MIOPEN_INTERNALS_EXPORT void
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem_desc);
MIOPEN_INTERNALS_EXPORT bool
SetNextValue(const miopen::batchnorm::ProblemDescription& problem_desc);
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool SetNextValue(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool IsValidValue() const;
MIOPEN_INTERNALS_EXPORT bool
IsValid(const ExecutionContext&,
const miopen::batchnorm::ProblemDescription& problem_desc) const;
IsValid(const ExecutionContext&, const miopen::batchnorm::ProblemDescription& problem) const;

template <typename Self, typename F>
static void Visit(Self&& s, F f)
Expand Down Expand Up @@ -255,23 +253,24 @@ struct BnCKBwdBackward final : BatchNormTunableSolver<PerformanceConfigBnCKBwdBa
{
const std::string& SolverDbId() const override { return GetSolverDbId<BnCKBwdBackward>(); }

MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem) const override;
bool IsDynamic() const override { return true; }
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKBwdBackward GetDefaultPerformanceConfig(
const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
const miopen::batchnorm::ProblemDescription& problem) const override;
MIOPEN_INTERNALS_EXPORT bool
IsValidPerformanceConfig(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKBwdBackward& config) const override;
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKBwdBackward
Search(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx) const override;
MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
MIOPEN_INTERNALS_EXPORT ConvSolution
GetSolution(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKBwdBackward& config) const override;
};

Expand All @@ -287,13 +286,11 @@ struct PerformanceConfigBnCKFwdTraining : PerfConfigBaseCK<PerformanceConfigBnCK
PerformanceConfigBnCKFwdTraining() : PerformanceConfigBnCKFwdTraining(0, "") {}
PerformanceConfigBnCKFwdTraining(bool) : PerformanceConfigBnCKFwdTraining(0, "") {}
MIOPEN_INTERNALS_EXPORT void
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem_desc);
MIOPEN_INTERNALS_EXPORT bool
SetNextValue(const miopen::batchnorm::ProblemDescription& problem_desc);
HeuristicInit(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool SetNextValue(const miopen::batchnorm::ProblemDescription& problem);
MIOPEN_INTERNALS_EXPORT bool IsValidValue() const;
MIOPEN_INTERNALS_EXPORT bool
IsValid(const ExecutionContext&,
const miopen::batchnorm::ProblemDescription& problem_desc) const;
IsValid(const ExecutionContext&, const miopen::batchnorm::ProblemDescription& problem) const;

template <typename Self, typename F>
static void Visit(Self&& s, F f)
Expand Down Expand Up @@ -323,23 +320,24 @@ struct BnCKFwdTraining final : BatchNormTunableSolver<PerformanceConfigBnCKFwdTr
{
const std::string& SolverDbId() const override { return GetSolverDbId<BnCKFwdTraining>(); }

MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem) const override;
bool IsDynamic() const override { return true; }
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKFwdTraining GetDefaultPerformanceConfig(
const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
const miopen::batchnorm::ProblemDescription& problem) const override;
MIOPEN_INTERNALS_EXPORT bool
IsValidPerformanceConfig(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKFwdTraining& config) const override;
MIOPEN_INTERNALS_EXPORT PerformanceConfigBnCKFwdTraining
Search(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const AnyInvokeParams& invoke_ctx) const override;
MIOPEN_INTERNALS_EXPORT bool
IsApplicable(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc) const override;
MIOPEN_INTERNALS_EXPORT ConvSolution
GetSolution(const ExecutionContext& ctx,
const miopen::batchnorm::ProblemDescription& problem_desc,
const miopen::batchnorm::ProblemDescription& problem,
const PerformanceConfigBnCKFwdTraining& config) const override;
};

Expand Down
1 change: 0 additions & 1 deletion src/include/miopen/conv/solvers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2807,7 +2807,6 @@ struct ConvCkIgemmFwdV6r1DlopsNchw final : ConvTunableSolver<PerformanceConvCkIg
MIOPEN_INTERNALS_EXPORT size_t GetWorkspaceSize(
const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
bool MayNeedWorkspace() const override { return true; }
bool IsDynamic() const override { return false; }
MIOPEN_INTERNALS_EXPORT PerformanceConvCkIgemmFwdV6r1DlopsNchw GetDefaultPerformanceConfig(
const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
MIOPEN_INTERNALS_EXPORT bool
Expand Down
Loading