feat(miopen): time-bound naive solver evaluation during find

projects/miopen/CHANGELOG.md

@@ -17,6 +17,7 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 * [Conv] Restored gfx12x support in the Winograd Rage solver, recovering performance that regressed when earlier gfx12 support was reverted.
 * [Conv] Refreshed the gfx1100, gfx1102, and gfx1201 (Navi) SystemDBs with updated tuned find/perf-database entries.
 * [Conv] Refreshed the gfx950 SystemDB with additional tuned entries to cover more models.
+* [Conv] Naive convolution solvers are now time-bounded during find when a non-naive solver has already succeeded. The warmup run is launched on a separate stream with a budget of 3x the best non-naive time; if it exceeds this budget the solver is skipped. Set `MIOPEN_NAIVE_TIMEOUT=0` to disable this behavior.
 
 ### Resolved Issues
 * [BatchNorm] Fixed an off-by-stride indexing bug in the backward CalcStats mean/variance remainder loop that caused a ~1% systematic bias in NCHW batch normalization backward results.
@@ -35,7 +36,6 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 
 ### Changed
 * [Conv] Load the Composable Kernel (CK) dynamic library during handle creation (`miopenCreate`) instead of on first CK use, so the one-time library-load cost is paid at a predictable setup point rather than as an unpredictable mid-run stall.
-* [Conv] Naive convolution solvers are now skipped by default during find when any non-naive solver succeeds across any algorithm. Set `MIOPEN_NAIVE_DISABLE_IF_ALT=0` to restore the previous behavior.
 * [Conv] Refreshed the gfx942 SystemDB with updated tuned entries from recent perf-eval runs.
 * [Conv] Refined the 2D AI heuristics (Tunanet solver selection and kernel-tuning-network candidate ranking) with feature-engineered inputs for gfx942 and gfx950, improving selection for grouped implicit-GEMM convolutions.
 

projects/miopen/docs/reference/env_variables.rst

@@ -104,13 +104,40 @@ and :doc:`Performance database <../conceptual/perfdb>`.
       - | 1: Enable
         | 0 or unset: Disable
 
-    * - | ``MIOPEN_NAIVE_DISABLE_IF_ALT``
-        | Skips naive convolution solvers during find when at least one
-        | non-naive solver succeeds across any algorithm. Naive solvers
-        | are still used as a fallback when no non-naive solver succeeds.
+    * - | ``MIOPEN_NAIVE_TIMEOUT``
+        | When enabled and a non-naive solver has already succeeded,
+        | naive solver evaluation is launched on a separate HIP stream
+        | with a wall-clock budget of ``MIOPEN_NAIVE_TIMEOUT_FACTOR``
+        | percent of the best non-naive time. If the naive kernel does
+        | not complete within the budget, it is skipped and the stream
+        | is returned to an async pool for later reclamation. Naive
+        | solvers are still used as a fallback when no non-naive solver
+        | succeeds.
+        |
+        | Limitation: a skipped kernel keeps reading the caller's input
+        | tensors until it finishes. Its results are discarded, so stale
+        | or reused data is harmless, but unmapping those pages while it
+        | reads them faults the process. Freeing with ``hipFree`` is safe
+        | because it synchronizes the device first. ``hipFreeAsync`` is
+        | ordered only against the stream it is given and does not wait
+        | for a skipped kernel; the pool normally keeps the address range
+        | reserved, but it releases pages back to the operating system
+        | once ``hipMemPoolAttrReleaseThreshold`` is exceeded, which its
+        | default value of 0 always is. Disable this option if the
+        | application frees convolution inputs through the
+        | stream-ordered allocator.
       - | 1 or unset: Enable (default)
         | 0: Disable
 
+    * - | ``MIOPEN_NAIVE_TIMEOUT_FACTOR``
+        | Sets the wall-clock budget for naive solver evaluation as a
+        | percentage of the best non-naive solver time. Only takes
+        | effect when ``MIOPEN_NAIVE_TIMEOUT`` is enabled and a
+        | non-naive solver has already succeeded. For example, the
+        | default value of 300 gives naive a budget of 3× the best
+        | non-naive time.
+      - | Integer percentage (default: 300)
+
     * - | ``MIOPEN_DEBUG_DISABLE_FIND_DB``
         | Disables FindDb functionality.
       - | 1: Disable FindDb

projects/miopen/src/conv/solver_finders.cpp

@@ -4,15 +4,23 @@
 #include <miopen/conv/solver_finders.hpp>
 
 #include <algorithm>
+#include <chrono>
 #include <numeric>
+#include <thread>
 
 #include <miopen/conv_algo_name.hpp>
+#include <miopen/handle.hpp>
+#include <miopen/hipoc_kernel.hpp>
+
+#include <hip/hip_runtime.h>
 #include <miopen/config.h>
 #include <miopen/env.hpp>
 #include <miopen/kernel_tuning_mode.hpp>
 #include <miopen/mlo_internal.hpp>
 #include <miopen/perf_field.hpp>
+#include <miopen/conv/data_invoke_params.hpp>
 #include <miopen/conv/problem_description.hpp>
+#include <miopen/conv/wrw_invoke_params.hpp>
 #include <miopen/solution.hpp>
 #include <miopen/utility/modified_z.hpp>
 
@@ -26,7 +34,8 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY)
 
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE)
 
-MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_DISABLE_IF_ALT, true)
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true)
+MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_NAIVE_TIMEOUT_FACTOR, 300)
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_FIND_SKIP_PCT, 130)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CONV_DIRECT_MAX_SIZE, 0)
@@ -202,6 +211,147 @@ const std::vector<std::unique_ptr<ISolversFinder>>& GetConvSolverFinders()
 
 } // namespace conv
 
+namespace {
+
+struct NaiveWarmup
+{
+    enum class Status
+    {
+        Completed,
+        TimedOut,
+        ScratchUnavailable,
+        UnsupportedInvokeParams,
+    };
+
+    Status status;
+    float elapsed;
+};
+
+/// Redirects the handle onto a tracker-owned stream with profiling off, restoring
+/// both on scope exit so no early return can strand the handle off the root stream.
+struct AutoExclusiveStream
+{
+    AutoExclusiveStream(const Handle& h, hipStream_t stream)
+        : handle(h), prev_profiling(h.IsProfilingEnabled())
+    {
+        handle.SetExclusiveStream(stream);
+        handle.EnableProfiling(false);
+    }
+
+    ~AutoExclusiveStream()
+    {
+        handle.SetExclusiveStream(nullptr);
+        handle.EnableProfiling(prev_profiling);
+    }
+
+    AutoExclusiveStream(const AutoExclusiveStream&)            = delete;
+    AutoExclusiveStream& operator=(const AutoExclusiveStream&) = delete;
+
+private:
+    const Handle& handle;
+    bool prev_profiling;
+};
+
+std::string NaiveSkipReason(NaiveWarmup::Status status, float best_time)
+{
+    switch(status)
+    {
+    case NaiveWarmup::Status::TimedOut:
+        return "exceeded " + std::to_string(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) +
+               "% of best non-naive time (" + std::to_string(best_time) + " ms)";
+    case NaiveWarmup::Status::ScratchUnavailable:
+        return "output tensor exceeds the scratch buffer cap";
+    case NaiveWarmup::Status::UnsupportedInvokeParams:
+        return "invoke params type does not support scratch redirection";
+    case NaiveWarmup::Status::Completed: break;
+    }
+    return "completed";
+}
+
+} // namespace
+
+static NaiveWarmup TryNaiveWithTimeout(const Handle& handle,
+                                       const Invoker& invoker,
+                                       const AnyInvokeParams& invoke_ctx,
+                                       float best_time)
+{
+    std::shared_ptr<ScratchAllocation> scratch;
+    AnyInvokeParams scratch_ctx;
+
+    if(invoke_ctx.IsOfType<conv::DataInvokeParams>())
+    {
+        auto params = invoke_ctx.CastTo<conv::DataInvokeParams>();
+        scratch     = handle.GetScratchBuffer(params.tensors.outDesc.GetNumBytes());
+        if(!scratch)
+            return {NaiveWarmup::Status::ScratchUnavailable, 0.0f};
+        params.tensors.out = scratch->buffer.get();
+        scratch_ctx        = AnyInvokeParams{params};
+    }
+    else if(invoke_ctx.IsOfType<conv::WrWInvokeParams>())
+    {
+        auto params = invoke_ctx.CastTo<conv::WrWInvokeParams>();
+        scratch     = handle.GetScratchBuffer(params.tensors.dwDesc.GetNumBytes());
+        if(!scratch)
+            return {NaiveWarmup::Status::ScratchUnavailable, 0.0f};
+        params.tensors.dw = scratch->buffer.get();
+        scratch_ctx       = AnyInvokeParams{params};
+    }
+    else
+        return {NaiveWarmup::Status::UnsupportedInvokeParams, 0.0f};
+
+    auto& tracker = handle.GetStreamTracker();
+    auto slot     = tracker.acquire(handle);
+    slot.scratch  = scratch;
+
+    AutoExclusiveStream stream_guard{handle, slot.stream};
+
+    HipEventPtr ev_start = make_hip_event();
+    HipEventPtr ev_stop  = make_hip_event();
+
+    try
+    {
+        auto ev_status = hipEventRecord(ev_start.get(), slot.stream);
+        if(ev_status != hipSuccess)
+            MIOPEN_THROW_HIP_STATUS(ev_status, "Failed to record naive start event");
+        invoker(handle, scratch_ctx);
+        ev_status = hipEventRecord(ev_stop.get(), slot.stream);
+        if(ev_status != hipSuccess)
+            MIOPEN_THROW_HIP_STATUS(ev_status, "Failed to record naive stop event");
+    }
+    catch(...)
+    {
+        tracker.abandon(slot);
+        throw;
+    }
+
+    const float timeout_factor =
+        static_cast<float>(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) / 100.0f;
+    const float naive_budget = best_time * timeout_factor;
+    const auto deadline      = std::chrono::steady_clock::now() +
+                          std::chrono::microseconds(static_cast<long long>(naive_budget * 1000));
+    bool finished = false;
+    while(std::chrono::steady_clock::now() < deadline)
+    {
+        if(hipEventQuery(ev_stop.get()) == hipSuccess)
+        {
+            finished = true;
+            break;
+        }
+        std::this_thread::yield();
+    }
+
+    if(finished)
+    {
+        tracker.release(slot);
+        float warmup_elapsed = 0.0f;
+        (void)hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get());
+        return {NaiveWarmup::Status::Completed, warmup_elapsed};
+    }
+
+    tracker.abandon(slot);
+    return {NaiveWarmup::Status::TimedOut, 0.0f};
+}
+
 /// Register invoker only for the best solution within algorithm.
 std::vector<Solution> EvaluateInvokers(const Handle& handle,
                                        const std::vector<solver::ConvSolution>& solutions,
@@ -222,11 +372,11 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
         return s.solver_id.find("Naive") != std::string::npos;
     };
 
-    bool naive_disable       = env::value(MIOPEN_NAIVE_DISABLE_IF_ALT);
+    bool naive_timeout       = env::value(MIOPEN_NAIVE_TIMEOUT);
     bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF);
     // Defer Naive only when a non-Naive alternative exists across all algorithms or this one.
     const bool defer_naive =
-        naive_disable &&
+        naive_timeout &&
         (non_naive_succeeded || std::any_of(solutions.begin(), solutions.end(), [&](const auto& s) {
              return !is_naive_solver(s);
          }));
@@ -249,18 +399,8 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
     {
         const auto& sol = solutions[idx];
 
-        const bool is_naive = is_naive_solver(sol);
-        if(naive_disable && is_naive)
-        {
-            if(defer_naive && non_naive_succeeded)
-            {
-                MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":"
-                                                       << sol.solver_id);
-                continue;
-            }
-            MIOPEN_LOG_I("Unable to Skip Naive Solver: " << algorithm_name.ToString() << ":"
-                                                         << sol.solver_id);
-        }
+        const bool is_naive     = is_naive_solver(sol);
+        const bool cutoff_naive = defer_naive && is_naive && non_naive_succeeded;
 
         if(!conv::IsEnoughWorkspace(
                "EvaluateInvokers", solver::Id{sol.solver_id}, sol.workspace_sz, &invoke_ctx))
@@ -329,6 +469,22 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
             int i                           = 0;
             samples.clear();
 
+            if(cutoff_naive)
+            {
+                const auto warmup = TryNaiveWithTimeout(
+                    handle, invoker, invoke_ctx, core_result.find_search_best_time);
+                if(warmup.status != NaiveWarmup::Status::Completed)
+                {
+                    MIOPEN_LOG_I(
+                        "Skipped naive solver "
+                        << algorithm_name.ToString() << ":" << sol.solver_id << ": "
+                        << NaiveSkipReason(warmup.status, core_result.find_search_best_time));
+                    continue;
+                }
+                first_elapsed = warmup.elapsed;
+                i             = 1;
+            }
+
             while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX)
             {
                 invoker(handle, invoke_ctx);

projects/miopen/src/env_debug.cpp

@@ -75,7 +75,8 @@ MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_CUSTOM_CACHE_DIR)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_LOG_LEVEL)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_LOG_BUFFER_SIZE, 0);
 MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS)
-MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_DISABLE_IF_ALT, true)
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true)
+MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_NAIVE_TIMEOUT_FACTOR, 300)
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SKIP_ASAN_DISABLED_TESTS, true)
 
 namespace miopen {
@@ -215,7 +216,8 @@ const LibEnvVar& FindEnvVariable(std::string_view name)
         {MIOPEN_FIND_ENFORCE.GetName(), MIOPEN_FIND_ENFORCE},
         {MIOPEN_FIND_MODE.GetName(), MIOPEN_FIND_MODE},
         {MIOPEN_FIND_MODE_FUSION.GetName(), MIOPEN_FIND_MODE_FUSION},
-        {MIOPEN_NAIVE_DISABLE_IF_ALT.GetName(), MIOPEN_NAIVE_DISABLE_IF_ALT},
+        {MIOPEN_NAIVE_TIMEOUT.GetName(), MIOPEN_NAIVE_TIMEOUT},
+        {MIOPEN_NAIVE_TIMEOUT_FACTOR.GetName(), MIOPEN_NAIVE_TIMEOUT_FACTOR},
         {MIOPEN_LOG_LEVEL.GetName(), MIOPEN_LOG_LEVEL},
         {MIOPEN_LOG_BUFFER_SIZE.GetName(), MIOPEN_LOG_BUFFER_SIZE},
         {MIOPEN_USER_DB_PATH.GetName(), MIOPEN_USER_DB_PATH},