Find Mode TrustVerify

include/miopen/miopen.h

@@ -664,6 +664,8 @@ typedef enum
               miss, uses the existing Find machinery with skipping non-dynamic kernels, thus saving
               compilation time. Faster start-up times than Hybrid Find, but GPU performance might be
               a bit worse. >*/
+    miopenConvolutionFindModeTrustVerify =
+        6,
     miopenConvolutionFindModeDefault =
         miopenConvolutionFindModeDynamicHybrid /*!< Default FindMode >*/
 } miopenConvolutionFindMode_t;

src/conv/solver_finders.cpp

@@ -216,13 +216,13 @@ const std::vector<std::unique_ptr<ISolversFinder>>& GetConvSolverFinders()
 } // namespace conv
 
 /// Register invoker only for the best solution within algorithm.
-static std::vector<Solution> EvaluateInvokers(const Handle& handle,
-                                              const std::vector<solver::ConvSolution>& solutions,
-                                              const AlgorithmName& algorithm_name,
-                                              const NetworkConfig& network_config,
-                                              const AnyInvokeParams& invoke_ctx,
-                                              bool& is_result_optimal,
-                                              bool force_attach_binary)
+std::vector<Solution> EvaluateInvokers(const Handle& handle,
+                                       const std::vector<solver::ConvSolution>& solutions,
+                                       const AlgorithmName& algorithm_name,
+                                       const NetworkConfig& network_config,
+                                       const AnyInvokeParams& invoke_ctx,
+                                       bool& is_result_optimal,
+                                       bool force_attach_binary)
 {
     const auto arch = env::value(MIOPEN_DEVICE_ARCH);
     if(!arch.empty())

src/convolution.cpp

@@ -408,7 +408,7 @@ std::size_t ConvolutionDescriptor::GetWorkSpaceSize(ExecutionContext ctx,
     ctx.do_search             = false;
     ctx.disable_perfdb_access = true;
 
-    while(findMode.IsFast(ctx) || findMode.IsHybrid(ctx))
+    while(findMode.IsFast(ctx) || (findMode.IsHybrid(ctx) && !findMode.IsTrustVerify(ctx)))
     {
         /// \section ffind_gwss_why_not_0
         /// Basically we can return 0 here because

src/find_controls.cpp

@@ -188,6 +188,7 @@ const char* ToCString(const FindMode::Values mode)
     case FindMode::Values::Hybrid: return "HYBRID";
     case FindMode::Values::DeprecatedFastHybrid: break;
     case FindMode::Values::DynamicHybrid: return "DYNAMIC_HYBRID";
+    case FindMode::Values::TrustVerify: return "TRUST_VERIFY";
     case FindMode::Values::End_: break;
     }
     return "<Unknown>";
@@ -222,6 +223,10 @@ std::optional<FindMode::Values> GetFindModeValueImpl2(Variable variable)
     {
         return FindMode::Values::DynamicHybrid;
     }
+    else if(str == "TRUST_VERIFY")
+    {
+        return FindMode::Values::TrustVerify;
+    }
     else
     { // Nop. Fall down & try numerics.
     }
@@ -270,5 +275,8 @@ static_assert(miopenConvolutionFindModeHybrid ==
 static_assert(miopenConvolutionFindModeDynamicHybrid ==
                   static_cast<miopenConvolutionFindMode_t>(FindMode::Values::DynamicHybrid),
               "API is not in sync with the implementation.");
+static_assert(miopenConvolutionFindModeTrustVerify ==
+                  static_cast<miopenConvolutionFindMode_t>(FindMode::Values::TrustVerify),
+              "API is not in sync with the implementation.");
 
 } // namespace miopen

src/include/miopen/conv/solver_finders.hpp

@@ -164,6 +164,14 @@ struct FindCoreResult
     bool is_optimal;
 };
 
+std::vector<Solution> EvaluateInvokers(const Handle& handle,
+                                       const std::vector<solver::ConvSolution>& solutions,
+                                       const AlgorithmName& algorithm_name,
+                                       const NetworkConfig& network_config,
+                                       const AnyInvokeParams& invoke_ctx,
+                                       bool& is_result_optimal,
+                                       bool force_attach_binary);
+
 FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx,
                         const ExecutionContext& ctx,
                         const ProblemDescriptionBase& problem,

src/include/miopen/find_controls.hpp

@@ -119,6 +119,7 @@ class MIOPEN_INTERNALS_EXPORT FindMode
         Hybrid               = miopenConvolutionFindModeHybrid,
         DeprecatedFastHybrid = 4,
         DynamicHybrid        = miopenConvolutionFindModeDynamicHybrid,
+        TrustVerify          = miopenConvolutionFindModeTrustVerify,
         End_,
         Default_ = miopenConvolutionFindModeDefault,
     };
@@ -152,13 +153,22 @@ class MIOPEN_INTERNALS_EXPORT FindMode
     template <class Context>
     bool IsHybrid(const Context& context) const
     {
-        return (value == Values::Hybrid || value == Values::DynamicHybrid) && IsEnabled(context);
+        return (value == Values::Hybrid || value == Values::DynamicHybrid ||
+                value == Values::TrustVerify) &&
+               IsEnabled(context);
     }
 
     template <class Context>
     bool IsDynamicHybrid(const Context& context) const
     {
-        return value == Values::DynamicHybrid && IsEnabled(context);
+        return (value == Values::DynamicHybrid || value == Values::TrustVerify) &&
+               IsEnabled(context);
+    }
+
+    template <class Context>
+    bool IsTrustVerify(const Context& context) const
+    {
+        return value == Values::TrustVerify && IsEnabled(context);
     }
 
     MIOPEN_INTERNALS_EXPORT friend std::ostream& operator<<(std::ostream&, const FindMode&);