diff --git a/src/include/miopen/activ/solvers.hpp b/src/include/miopen/activ/solvers.hpp index 725d894330..8ae2642a32 100644 --- a/src/include/miopen/activ/solvers.hpp +++ b/src/include/miopen/activ/solvers.hpp @@ -48,6 +48,7 @@ using OldStyleSolver = SolverMixin; struct ActivFwdSolver0 : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -57,7 +58,7 @@ struct ActivFwdSolver0 : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -71,6 +72,7 @@ struct ActivFwdSolver0 : OldStyleSolver struct ActivFwdSolver1 : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -80,7 +82,7 @@ struct ActivFwdSolver1 : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -94,6 +96,7 @@ struct ActivFwdSolver1 : OldStyleSolver struct ActivBwdSolver0 : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -103,7 +106,7 @@ struct ActivBwdSolver0 : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -117,6 +120,7 @@ struct ActivBwdSolver0 : OldStyleSolver struct ActivBwdSolver1 : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -126,7 +130,7 @@ struct ActivBwdSolver1 : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } diff --git a/src/include/miopen/batchnorm/solvers.hpp b/src/include/miopen/batchnorm/solvers.hpp index 38aef49e05..d4463b6d1f 100644 --- a/src/include/miopen/batchnorm/solvers.hpp +++ b/src/include/miopen/batchnorm/solvers.hpp @@ -48,6 +48,7 @@ using OldStyleSolver = SolverMixin; struct BnFwdTrainingSpatialSingle : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -60,7 +61,7 @@ struct BnFwdTrainingSpatialSingle : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -74,6 +75,7 @@ struct BnFwdTrainingSpatialSingle : OldStyleSolver struct BnFwdTrainingSpatialMultiple : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -86,7 +88,7 @@ struct BnFwdTrainingSpatialMultiple : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -100,6 +102,7 @@ struct BnFwdTrainingSpatialMultiple : OldStyleSolver struct BnFwdTrainingPerActivation : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -112,7 +115,7 @@ struct BnFwdTrainingPerActivation : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -126,6 +129,7 @@ struct BnFwdTrainingPerActivation : OldStyleSolver struct BnBwdTrainingSpatialSingle : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -138,7 +142,7 @@ struct BnBwdTrainingSpatialSingle : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -152,6 +156,7 @@ struct BnBwdTrainingSpatialSingle : OldStyleSolver struct BnBwdTrainingSpatialMultiple : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -164,7 +169,7 @@ struct BnBwdTrainingSpatialMultiple : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -178,6 +183,7 @@ struct BnBwdTrainingSpatialMultiple : OldStyleSolver struct BnBwdTrainingPerActivation : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override @@ -190,7 +196,7 @@ struct BnBwdTrainingPerActivation : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } @@ -204,6 +210,7 @@ struct BnBwdTrainingPerActivation : OldStyleSolver struct BnFwdInference : OldStyleSolver { // To suppress -Woverloaded-virtual + using OldStyleSolver::GetSolution; using OldStyleSolver::IsApplicable; const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -213,7 +220,7 @@ struct BnFwdInference : OldStyleSolver return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - inline ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } diff --git a/src/include/miopen/pooling/solvers.hpp b/src/include/miopen/pooling/solvers.hpp index 56e109b2e5..ef2d819679 100644 --- a/src/include/miopen/pooling/solvers.hpp +++ b/src/include/miopen/pooling/solvers.hpp @@ -52,13 +52,14 @@ struct OldStyleSolver : SolverMixin // To suppress -Woverloaded-virtual using SolverMixin::GetWorkspaceSize; using SolverMixin::IsApplicable; + using SolverMixin::GetSolution; bool IsApplicable(const OldStyleProblemDescription& problem) const override { return IsApplicable(*std::get<0>(problem), *std::get<1>(problem)); } - ConvSolution GetSolution(const OldStyleProblemDescription& problem) const + ConvSolution GetSolution(const OldStyleProblemDescription& problem) const override { return GetSolution(*std::get<0>(problem), *std::get<1>(problem)); } diff --git a/src/include/miopen/solver.hpp b/src/include/miopen/solver.hpp index 1f45a1d71b..1706dcbb5f 100644 --- a/src/include/miopen/solver.hpp +++ b/src/include/miopen/solver.hpp @@ -128,7 +128,7 @@ struct SolverBase /// Takes problem config, optimization parameters and other info /// and computes information required to build and run the kernel(s). - /// ConvSolution GetSolution(const ConvolutionContext& params) const; + virtual ConvSolution GetSolution(const boost::any& ctx) const = 0; /// Searchable solvers provide a GetSolution that takes a Context and PerformanceConfig /// ConvSolution GetSolution(const ConvolutionContext& params, @@ -160,6 +160,10 @@ struct SolverMixin : SolverBase virtual bool IsApplicable(const Context& ctx) const = 0; virtual float GetWti(const Context&) const { return -2.0; }; virtual size_t GetWorkspaceSize(const Context&) const { return 0; }; + virtual ConvSolution GetSolution(const Context&) const + { + MIOPEN_THROW(miopenStatusInternalError, "Solver is being used incorrectly"); + } bool IsApplicable(const boost::any& ctx) const final { @@ -175,6 +179,11 @@ struct SolverMixin : SolverBase { return GetWorkspaceSize(boost::any_cast(ctx)); } + + ConvSolution GetSolution(const boost::any& ctx) const final + { + return GetSolution(boost::any_cast(ctx)); + } }; // Typedef for convolution solvers @@ -208,6 +217,9 @@ struct PerformanceConfigConvAsm3x3U : Serializable struct ConvAsm3x3U : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; @@ -277,6 +289,9 @@ struct PerformanceConfigConvAsm1x1U : Serializable struct ConvAsm1x1U : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } PerformanceConfigConvAsm1x1U GetPerformanceConfig(const ConvolutionContext&) const; @@ -378,6 +393,9 @@ struct PerformanceConfigConvAsm1x1UV2 : Serializable(); } PerformanceConfigConvAsm1x1UV2 GetPerformanceConfig(const ConvolutionContext&) const; @@ -396,7 +414,7 @@ struct ConvAsm5x10u2v2f1 : ConvSolver const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvAsm5x10u2v2b1 : ConvSolver @@ -404,7 +422,7 @@ struct ConvAsm5x10u2v2b1 : ConvSolver const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvAsm7x7c3h224w224k64u2v2p3q3f1 : ConvSolver @@ -415,7 +433,7 @@ struct ConvAsm7x7c3h224w224k64u2v2p3q3f1 : ConvSolver } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvOclDirectFwd11x11 : ConvSolver @@ -426,7 +444,7 @@ struct ConvOclDirectFwd11x11 : ConvSolver } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvOclDirectFwdGen : ConvSolver @@ -434,7 +452,7 @@ struct ConvOclDirectFwdGen : ConvSolver const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct PerformanceImplicitGemm : Serializable @@ -822,6 +840,9 @@ struct PerformanceImplicitGemmBwdDataV4R1Xdlops struct ConvHipImplicitGemmV4R1Fwd : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -842,6 +863,9 @@ struct ConvHipImplicitGemmV4R1Fwd : ConvSolver struct ConvHipImplicitGemmV4R4Fwd : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -903,6 +927,9 @@ struct PerformanceConvMlirIgemm : Serializable struct ConvMlirIgemmFwd : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& ctx) const override; @@ -965,13 +992,15 @@ struct PerformanceConvMlirIgemmXdlops : Serializable(); } bool IsApplicable(const ConvolutionContext& ctx) const override; - ConvSolution GetSolution(const ConvolutionContext& ctx) const; PerformanceConvMlirIgemmXdlops GetPerformanceConfig(const ConvolutionContext& ctx) const; bool IsValidPerformanceConfig(const ConvolutionContext& ctx, const PerformanceConvMlirIgemmXdlops& config) const; @@ -1029,6 +1058,9 @@ struct PerformanceImplicitGemmV4R4GenXdlopsFwdFp32 struct ConvHipImplicitGemmV4R4WrW : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1048,6 +1080,9 @@ struct ConvHipImplicitGemmV4R4WrW : ConvSolver struct ConvMlirIgemmWrW : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& ctx) const override; @@ -1063,6 +1098,9 @@ struct ConvMlirIgemmWrW : ConvSolver struct ConvMlirIgemmWrWXdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1348,6 +1386,9 @@ struct PerformanceImplicitGemmBwdV1R1Xdlops : Serializable(); @@ -1369,6 +1410,9 @@ struct ConvHipImplicitGemmForwardV4R4Xdlops : ConvSolver struct ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1392,6 +1436,9 @@ struct ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm : ConvSolver struct ConvHipImplicitGemmForwardV4R5Xdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1458,6 +1505,9 @@ struct PerformanceImplicitGemmV4R4GenXdlopsWrWFp32 struct ConvHipImplicitGemmV4R1WrW : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1477,6 +1527,9 @@ struct ConvHipImplicitGemmV4R1WrW : ConvSolver struct ConvHipImplicitGemmBwdDataV1R1 : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1498,6 +1551,9 @@ struct ConvHipImplicitGemmBwdDataV1R1 : ConvSolver struct ConvMlirIgemmBwd : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& ctx) const override; @@ -1513,6 +1569,9 @@ struct ConvMlirIgemmBwd : ConvSolver struct ConvMlirIgemmBwdXdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1531,6 +1590,9 @@ struct ConvMlirIgemmBwdXdlops : ConvSolver struct ConvHipImplicitGemmBwdDataV4R1 : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1551,6 +1613,9 @@ struct ConvHipImplicitGemmBwdDataV4R1 : ConvSolver struct ConvHipImplicitGemmBwdDataV4R1Xdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1573,6 +1638,9 @@ struct ConvHipImplicitGemmBwdDataV4R1Xdlops : ConvSolver struct ConvHipImplicitGemmBwdDataV1R1Xdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1601,7 +1669,7 @@ struct ConvAsmImplicitGemmV4R1DynamicFwd : ConvSolver bool IsApplicable(const ConvolutionContext& ctx) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvAsmImplicitGemmV4R1DynamicFwd_1x1 : ConvSolver @@ -1613,7 +1681,7 @@ struct ConvAsmImplicitGemmV4R1DynamicFwd_1x1 : ConvSolver bool IsApplicable(const ConvolutionContext& ctx) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvAsmImplicitGemmV4R1DynamicWrw : ConvSolver @@ -1627,7 +1695,7 @@ struct ConvAsmImplicitGemmV4R1DynamicWrw : ConvSolver bool IsDynamic() const override { return true; } size_t GetWorkspaceSize(const ConvolutionContext& ctx) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvAsmImplicitGemmGTCDynamicWrwXdlops : ConvSolver @@ -1641,7 +1709,7 @@ struct ConvAsmImplicitGemmGTCDynamicWrwXdlops : ConvSolver bool IsDynamic() const override { return true; } size_t GetWorkspaceSize(const ConvolutionContext& ctx) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvAsmImplicitGemmV4R1DynamicBwd : ConvSolver @@ -1653,7 +1721,7 @@ struct ConvAsmImplicitGemmV4R1DynamicBwd : ConvSolver bool IsApplicable(const ConvolutionContext&) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext&) const; + ConvSolution GetSolution(const ConvolutionContext&) const override; }; struct ConvAsmImplicitGemmGTCDynamicFwdXdlops : ConvSolver @@ -1665,7 +1733,7 @@ struct ConvAsmImplicitGemmGTCDynamicFwdXdlops : ConvSolver bool IsApplicable(const ConvolutionContext& ctx) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvAsmImplicitGemmGTCDynamicBwdXdlops : ConvSolver @@ -1677,7 +1745,7 @@ struct ConvAsmImplicitGemmGTCDynamicBwdXdlops : ConvSolver bool IsApplicable(const ConvolutionContext& ctx) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; /// Holds common member functions for the Solvers which share the same @@ -1695,6 +1763,9 @@ struct ConvOclDirectFwdLegacyExhaustiveSearch : ConvSolver struct ConvOclDirectFwd : ConvOclDirectFwdLegacyExhaustiveSearch { + // To suppress -Woverloaded-virtual + using ConvOclDirectFwdLegacyExhaustiveSearch::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; @@ -1720,6 +1791,9 @@ struct ConvOclDirectFwdFused : ConvOclDirectFwd struct ConvOclDirectFwd1x1 : ConvOclDirectFwdLegacyExhaustiveSearch { + // To suppress -Woverloaded-virtual + using ConvOclDirectFwdLegacyExhaustiveSearch::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; @@ -1737,7 +1811,7 @@ struct ConvBinWinograd3x3U : ConvSolver bool IsApplicable(const ConvolutionContext& params) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvBinWinogradRxS : ConvSolver @@ -1746,7 +1820,7 @@ struct ConvBinWinogradRxS : ConvSolver bool IsApplicable(const ConvolutionContext& params) const override; bool IsDynamic() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct PerformanceConfigConvBinWinogradRxSf3x2 @@ -1774,6 +1848,9 @@ struct PerformanceConfigConvBinWinogradRxSf3x2 struct ConvBinWinogradRxSf3x2 : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1822,6 +1899,9 @@ struct PerformanceConfigConvBinWinogradRxSf2x3 struct ConvBinWinogradRxSf2x3 : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -1856,7 +1936,7 @@ struct ConvBinWinogradRxSf2x3g1 : ConvSolver bool IsApplicable(const ConvolutionContext& params) const override; bool IsDynamic() const override { return true; } float GetWti(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvBinWinogradRxSf2x3g1Fused : ConvSolver @@ -1867,7 +1947,7 @@ struct ConvBinWinogradRxSf2x3g1Fused : ConvSolver } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvBinWinogradRxSFused : ConvSolver @@ -1878,7 +1958,7 @@ struct ConvBinWinogradRxSFused : ConvSolver } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; template @@ -1894,7 +1974,7 @@ struct ConvMPBidirectWinograd : ConvSolver bool IsDynamic() const override { return true; } size_t GetWorkspaceSize(const ConvolutionContext& params) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; // kernel_file_name for solver identification static std::string GetSolverFileNames(int id) @@ -1939,6 +2019,9 @@ extern template struct ConvMPBidirectWinograd<6, 3>; template struct ConvMPBidirectWinograd_xdlops : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId< @@ -2037,7 +2120,7 @@ struct ConvWinograd3x3MultipassWrW : ConvSolver bool IsDynamic() const override { return true; } size_t GetWorkspaceSize(const ConvolutionContext& params) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; // kernel_file_name for solver identification static std::string GetSolverFileNames(int id) @@ -2139,6 +2222,9 @@ struct PerformanceConfigAsmDirect3x3WrW : Serializable(); } PerformanceConfigAsmDirect3x3WrW GetPerformanceConfig(const ConvolutionContext&) const; @@ -2245,6 +2331,9 @@ struct PerformanceConfigConvAsmBwdWrW1x1 : Serializable(); } PerformanceConfigConvAsmBwdWrW1x1 GetPerformanceConfig(const ConvolutionContext&) const; @@ -2324,6 +2413,9 @@ struct PerformanceConfigConvOclBwdWrw2 template struct ConvOclBwdWrW2 : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId>(); @@ -2374,7 +2466,7 @@ struct ConvOclBwdWrW2NonTunable : ConvOclBwdWrW2<1> } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; private: // This function dervied from ConvOclBwdWrW2 is declared private @@ -2392,7 +2484,7 @@ struct ConvOclBwdWrW53 : ConvSolver bool IsApplicable(const ConvolutionContext& params) const override; size_t GetWorkspaceSize(const ConvolutionContext& params) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; }; struct ConvOclBwdWrW1x1 : ConvSolver @@ -2400,7 +2492,7 @@ struct ConvOclBwdWrW1x1 : ConvSolver const std::string& SolverDbId() const override { return GetSolverDbId(); } bool IsApplicable(const ConvolutionContext& params) const override; - ConvSolution GetSolution(const ConvolutionContext& params) const; + ConvSolution GetSolution(const ConvolutionContext& params) const override; size_t GetWorkspaceSize(const ConvolutionContext& params) const override; bool MayNeedWorkspace() const override { return true; } }; @@ -2412,7 +2504,7 @@ struct fft : ConvSolver bool IsApplicable(const ConvolutionContext& ctx) const override; size_t GetWorkspaceSize(const ConvolutionContext& ctx) const override; bool MayNeedWorkspace() const override { return true; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct PerformanceImplicitGemmWrwV4R4Xdlops : Serializable @@ -2471,6 +2563,9 @@ struct PerformanceImplicitGemmWrwV4R4Xdlops : Serializable(); @@ -2553,6 +2648,9 @@ struct PerformanceImplicitGemmWrwV4R4Xdlops_Padded_Gemm struct ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -2599,6 +2697,9 @@ struct PerformanceConvCkIgemmFwdV6r1DlopsNchw : Serializable(); @@ -2630,7 +2731,7 @@ struct ConvDirectNaiveConvFwd : ConvSolver /// Use very small fixed value enough to backup GEMM for cases when /// GEMM is disabled due to MIOpenGemm or OCL compiler issues. float GetWti(const ConvolutionContext&) const override { return 0.01; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvDirectNaiveConvBwd : ConvSolver @@ -2645,7 +2746,7 @@ struct ConvDirectNaiveConvBwd : ConvSolver /// Use very small fixed value enough to backup GEMM for cases when /// GEMM is disabled due to MIOpenGemm or OCL compiler issues. float GetWti(const ConvolutionContext&) const override { return 0.01; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct ConvDirectNaiveConvWrw : ConvSolver @@ -2660,7 +2761,7 @@ struct ConvDirectNaiveConvWrw : ConvSolver /// Use very small fixed value enough to backup GEMM for cases when /// GEMM is disabled due to MIOpenGemm or OCL compiler issues. float GetWti(const ConvolutionContext&) const override { return 0.01; } - ConvSolution GetSolution(const ConvolutionContext& ctx) const; + ConvSolution GetSolution(const ConvolutionContext& ctx) const override; }; struct GemmFwdBase : ConvSolver @@ -2681,6 +2782,7 @@ struct GemmFwdBase : ConvSolver struct GemmFwd1x1_0_2 : GemmFwdBase { // To suppress -Woverloaded-virtual + using GemmFwdBase::GetSolution; using GemmFwdBase::GetWorkspaceSize; using GemmFwdBase::IsApplicable; @@ -2697,7 +2799,7 @@ struct GemmFwd1x1_0_2 : GemmFwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2710,6 +2812,7 @@ struct GemmFwd1x1_0_2 : GemmFwdBase struct GemmFwd1x1_0_1_int8 : GemmFwdBase { // To suppress -Woverloaded-virtual + using GemmFwdBase::GetSolution; using GemmFwdBase::GetWorkspaceSize; using GemmFwdBase::IsApplicable; @@ -2726,7 +2829,7 @@ struct GemmFwd1x1_0_1_int8 : GemmFwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2739,6 +2842,7 @@ struct GemmFwd1x1_0_1_int8 : GemmFwdBase struct GemmFwd1x1_0_1 : GemmFwdBase { // To suppress -Woverloaded-virtual + using GemmFwdBase::GetSolution; using GemmFwdBase::GetWorkspaceSize; using GemmFwdBase::IsApplicable; @@ -2755,7 +2859,7 @@ struct GemmFwd1x1_0_1 : GemmFwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2768,6 +2872,7 @@ struct GemmFwd1x1_0_1 : GemmFwdBase struct GemmFwdRest : GemmFwdBase { // To suppress -Woverloaded-virtual + using GemmFwdBase::GetSolution; using GemmFwdBase::GetWorkspaceSize; using GemmFwdBase::IsApplicable; @@ -2784,7 +2889,7 @@ struct GemmFwdRest : GemmFwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2812,6 +2917,7 @@ struct GemmBwdBase : ConvSolver struct GemmBwd1x1_stride2 : GemmBwdBase { // To suppress -Woverloaded-virtual + using GemmBwdBase::GetSolution; using GemmBwdBase::GetWorkspaceSize; using GemmBwdBase::IsApplicable; @@ -2828,7 +2934,7 @@ struct GemmBwd1x1_stride2 : GemmBwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2841,6 +2947,7 @@ struct GemmBwd1x1_stride2 : GemmBwdBase struct GemmBwd1x1_stride1 : GemmBwdBase { // To suppress -Woverloaded-virtual + using GemmBwdBase::GetSolution; using GemmBwdBase::GetWorkspaceSize; using GemmBwdBase::IsApplicable; @@ -2857,7 +2964,7 @@ struct GemmBwd1x1_stride1 : GemmBwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2870,6 +2977,7 @@ struct GemmBwd1x1_stride1 : GemmBwdBase struct GemmBwdRest : GemmBwdBase { // To suppress -Woverloaded-virtual + using GemmBwdBase::GetSolution; using GemmBwdBase::GetWorkspaceSize; using GemmBwdBase::IsApplicable; @@ -2886,7 +2994,7 @@ struct GemmBwdRest : GemmBwdBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2914,6 +3022,7 @@ struct GemmWrwBase : ConvSolver struct GemmWrw1x1_stride1 : GemmWrwBase { // To suppress -Woverloaded-virtual + using GemmWrwBase::GetSolution; using GemmWrwBase::GetWorkspaceSize; using GemmWrwBase::IsApplicable; @@ -2930,7 +3039,7 @@ struct GemmWrw1x1_stride1 : GemmWrwBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -2943,6 +3052,7 @@ struct GemmWrw1x1_stride1 : GemmWrwBase struct GemmWrwUniversal : GemmWrwBase { // To suppress -Woverloaded-virtual + using GemmWrwBase::GetSolution; using GemmWrwBase::GetWorkspaceSize; using GemmWrwBase::IsApplicable; @@ -2959,7 +3069,7 @@ struct GemmWrwUniversal : GemmWrwBase return IsApplicable(ctx, ctx.conv_problem); } - ConvSolution GetSolution(const ConvolutionContext& ctx) const + ConvSolution GetSolution(const ConvolutionContext& ctx) const override { return GetSolution(ctx, ctx.conv_problem); } @@ -3343,6 +3453,9 @@ struct PerformanceConfigAsmImplicitGemmGTCFwdXdlopsNHWC : PerformanceConfigAsmIm struct ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -3533,6 +3646,9 @@ struct PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC : PerformanceConfigAsmIm struct ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); @@ -3728,6 +3844,9 @@ struct PerformanceConfigAsmImplicitGemmGTCWrwXdlopsNHWC : PerformanceConfigAsmIm struct ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC : ConvSolver { + // To suppress -Woverloaded-virtual + using ConvSolver::GetSolution; + const std::string& SolverDbId() const override { return GetSolverDbId(); diff --git a/test/solver.cpp b/test/solver.cpp index 1fb2aa8c0e..32596eb952 100644 --- a/test/solver.cpp +++ b/test/solver.cpp @@ -54,7 +54,7 @@ class TrivialTestSolver : public solver::ConvSolver return context.in_width == 1; } - solver::ConvSolution GetSolution(const ConvolutionContext&) const + solver::ConvSolution GetSolution(const ConvolutionContext&) const override { solver::ConvSolution ret; solver::KernelInfo kernel; @@ -80,6 +80,9 @@ struct TestConfig : solver::Serializable class SearchableTestSolver : public solver::ConvSolver { + // To suppress -Woverloaded-virtual + using solver::ConvSolver::GetSolution; + public: static int searches_done() { return _serches_done; } static const char* FileName() { return "SearchableTestSolver"; }