Find Mode TrustVerify

docs/conceptual/tuningdb.rst

@@ -29,7 +29,6 @@ Enable this feature using these commands:
 
 .. code:: bash
 
-  export MIOPEN_FIND_MODE=3
   export MIOPEN_FIND_ENFORCE=3
   export MIOPEN_USER_DB_PATH="/user/specified/directory"
 
@@ -48,7 +47,6 @@ Enable this feature using these commands:
 
 .. code:: bash
 
-  export MIOPEN_FIND_MODE=3
   export MIOPEN_FIND_ENFORCE=3
   export MIOPEN_USER_DB_PATH="/user/specified/directory"
   export MIOPEN_SYSTEM_DB_PATH="$MIOPEN_USER_DB_PATH"
@@ -64,7 +62,6 @@ Then enable exhaustive tuning and run the unique commands.
 
 .. code:: bash
 
-  export MIOPEN_FIND_MODE=1
   export MIOPEN_FIND_ENFORCE=4
   export MIOPEN_USER_DB_PATH="/user/specified/directory"
 

docs/how-to/find-and-immediate.rst

@@ -237,6 +237,12 @@ modes by using the ``MIOPEN_FIND_MODE`` environment variable with one of these v
   entry. If there's a FindDb hit, it uses that entry. If there's a miss, it uses the existing find machinery,
   skipping non-dynamic kernels. It offers faster start-up times than hybrid find, but GPU performance
   might decrease.
+* ``TRUST_VERIFY``/``6`` (trust verify find): Checks :doc:`FindDb <../conceptual/finddb>` for an entry.
+  If there's a UserFindDb hit, it uses that entry.
+  If there's a FindDb hit the result is evaluated. If the ratio of evaluated to reported result time is
+  below the tolerance threshold, the result is used and added to the UserFindDb. Otherwise tuning will be triggered.
+  If there's a miss, tuning will be triggered, skipping non-dynamic kernels.
+  This mode may have slow start-up times, but typically selects the most performant solutions.
 
 The default find mode is ``DYNAMIC_HYBRID``. To run the full ``NORMAL`` find mode, use
 ``export MIOPEN_FIND_MODE=NORMAL`` or ``export MIOPEN_FIND_MODE=1``.

include/miopen/miopen.h

@@ -664,6 +664,7 @@ 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&);

test/gtest/conv_api_find_mode.cpp

@@ -64,6 +64,11 @@ TEST_F(CPU_ConvFindModeAPI_NONE, ConvFindModeAPI)
     MIOPEN_EXPECT_CHECK_RET(miopenGetConvolutionFindMode(conv_descr, &findMode));
     EXPECT_EQ(findMode, miopenConvolutionFindMode_t::miopenConvolutionFindModeDynamicHybrid);
 
+    MIOPEN_EXPECT_CHECK_RET(miopenSetConvolutionFindMode(
+        conv_descr, miopenConvolutionFindMode_t::miopenConvolutionFindModeTrustVerify));
+    MIOPEN_EXPECT_CHECK_RET(miopenGetConvolutionFindMode(conv_descr, &findMode));
+    EXPECT_EQ(findMode, miopenConvolutionFindMode_t::miopenConvolutionFindModeTrustVerify);
+
     MIOPEN_EXPECT_CHECK_RET(miopenSetConvolutionFindMode(
         conv_descr, miopenConvolutionFindMode_t::miopenConvolutionFindModeDefault));
     MIOPEN_EXPECT_CHECK_RET(miopenGetConvolutionFindMode(conv_descr, &findMode));