cderb · Aug 28, 2026 · Jul 15, 2026 · Jul 16, 2026 · Jul 16, 2026 · Jul 16, 2026
@@ -12,7 +12,7 @@ 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] Naive convolution solvers are now time-bounded during find when a non-naive solver has already succeeded. The warmup run is launched on a disposable 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.
 * [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.
 

@@ -4,9 +4,15 @@
 #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>
@@ -26,7 +32,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 +209,59 @@ const std::vector<std::unique_ptr<ISolversFinder>>& GetConvSolverFinders()
 
 } // namespace conv
 
+static float TryNaiveWithTimeout(const Handle& handle,
+                                 const Invoker& invoker,
+                                 const AnyInvokeParams& invoke_ctx,
+                                 float best_time)
+{
+    auto& tracker = handle.GetStreamTracker();
+    auto slot     = tracker.acquire(handle);
+
+    handle.SetStreamFromPool(slot.pool_id);
+    handle.EnableProfiling(false);
+
+    HipEventPtr ev_start = make_hip_event();
+    HipEventPtr ev_stop  = make_hip_event();
+
+    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, invoke_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");
+
+    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();
+    }
+
+    handle.SetStreamFromPool(0);
+    handle.EnableProfiling(true);
+
+    if(finished)
+    {
+        tracker.release(slot);
+        float warmup_elapsed = 0.0f;
+        (void)hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get());
+        return warmup_elapsed;
+    }
+
+    tracker.abandon(slot);
+    return -1.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 +282,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 +309,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 +379,23 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
             int i                           = 0;
             samples.clear();
 
+            if(cutoff_naive)
+            {
+                const float warmup = TryNaiveWithTimeout(
+                    handle, invoker, invoke_ctx, core_result.find_search_best_time);
+                if(warmup < 0)
+                {
+                    MIOPEN_LOG_I("Timeout: Naive Solver "
+                                 << algorithm_name.ToString() << ":" << sol.solver_id
+                                 << " exceeded " << env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)
+                                 << "% best budget (" << core_result.find_search_best_time
+                                 << " ms)");
+                    continue;
+                }
+                first_elapsed = warmup;
+                i             = 1;
+            }
+
             while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX)
             {
                 invoker(handle, invoke_ctx);

@@ -71,7 +71,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 {
@@ -208,7 +209,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},

@@ -283,6 +283,7 @@ struct HandleImpl
     Allocator allocator{};
     KernelCache cache;
     TargetProperties target_properties;
+    mutable StreamTracker stream_tracker_;
 };
 
 Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique<HandleImpl>())
@@ -400,6 +401,37 @@ miopenAcceleratorQueue_t Handle::GetStream() const
     return this->impl->ms_resourse_ptr->stream_pool.at(meopenHandle_current_stream_id - 1).get();
 }
 
+StreamTracker& Handle::GetStreamTracker() const { return impl->stream_tracker_; }
+
+StreamTracker::Slot StreamTracker::acquire(const Handle& handle)
+{
+    if(available_.empty())
+    {
+        for(auto it = draining_.begin(); it != draining_.end();)
+        {
+            if(hipStreamQuery(it->stream) == hipSuccess)
+            {
+                available_.push_back(*it);
+                it = draining_.erase(it);
+            }
+            else
+                ++it;
+        }
+    }
+
+    if(!available_.empty())
+    {
+        auto slot = available_.back();
+        available_.pop_back();
+        return slot;
+    }
+
+    int id = next_id_++;
+    handle.ReserveExtraStreamsInPool(id);
+    handle.SetStreamFromPool(id);
+    return {id, handle.GetStream()};
+}
+
 void Handle::SetAllocator(miopenAllocatorFunction allocator,
                           miopenDeallocatorFunction deallocator,
                           void* allocatorContext) const

@@ -38,6 +38,7 @@
 #include <miopen/simple_hash.hpp>
 #include <miopen/solver_id.hpp>
 #include <miopen/stringutils.hpp>
+#include <miopen/stream_tracker.hpp>
 #include <miopen/target_properties.hpp>
 
 #include <cstdio>
@@ -94,6 +95,8 @@ struct MIOPEN_EXPORT Handle : miopenHandle
     void SetStreamFromPool(int streamID) const;
     void ReserveExtraStreamsInPool(int cnt) const;
 
+    StreamTracker& GetStreamTracker() const;
+
     void SetAllocator(miopenAllocatorFunction allocator,
                       miopenDeallocatorFunction deallocator,
                       void* allocatorContext) const;

@@ -0,0 +1,36 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#ifndef GUARD_MIOPEN_STREAM_TRACKER_HPP_
+#define GUARD_MIOPEN_STREAM_TRACKER_HPP_
+
+#include <miopen/config.h>
+
+#include <vector>
+#include <hip/hip_runtime_api.h>
+
+namespace miopen {
+
+struct Handle;
+
+struct MIOPEN_EXPORT StreamTracker
+{
+    struct Slot
+    {
+        int pool_id;
+        hipStream_t stream;
+    };
+
+    Slot acquire(const Handle& handle);
+    void release(Slot slot) { available_.push_back(slot); }
+    void abandon(Slot slot) { draining_.push_back(slot); }
+
+private:
+    std::vector<Slot> available_;
+    std::vector<Slot> draining_;
+    int next_id_ = 1;
+};
+
+} // namespace miopen
+
+#endif // GUARD_MIOPEN_STREAM_TRACKER_HPP_