fix(miopen): keep naive-evaluation stream reclaim out of HIP graph captures

projects/miopen/CHANGELOG.md

@@ -9,6 +9,9 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise
   kernels to the `ConvHipConv` solver (hipconv v0.3.1).
 
+### Resolved Issues
+* [Conv] Fixed HIP graph captures of convolutions failing with `hipErrorStreamCaptureInvalidated` after a find had skipped a naive solver for exceeding `MIOPEN_NAIVE_TIMEOUT` (#12121). Scratch held by a skipped evaluation is now released by a background thread as soon as its kernel finishes, rather than from a later kernel launch, so it is also released in workloads that only replay a captured graph.
+
 ### Removed
 * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources.
 * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`.

projects/miopen/docs/reference/env_variables.rst

@@ -114,6 +114,21 @@ and :doc:`Performance database <../conceptual/perfdb>`.
         | solvers are still used as a fallback when no non-naive solver
         | succeeds.
         |
+        | At most ``GPU_MAX_HW_QUEUES`` minus 2 skipped streams (at least
+        | 1; 2 at the HIP default of 4 queues) may be outstanding. Beyond
+        | that, the next naive evaluation waits for one to finish rather
+        | than timing a kernel that shares a hardware queue with it.
+        |
+        | A skipped kernel's scratch buffer is released as soon as the
+        | kernel finishes, by a background thread MIOpen starts for the
+        | handle on the first skip; no further MIOpen call is needed, so
+        | a workload that only replays a captured HIP graph releases it
+        | too. A deallocator installed with ``miopenSetAllocator`` is
+        | therefore called from that thread, possibly while another
+        | thread is capturing. It must not make calls HIP refuses during
+        | any capture, such as a synchronous ``hipMemcpy`` or
+        | ``hipMemset``: those invalidate every capture in the process.
+        |
         | 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

projects/miopen/src/conv/solver_finders.cpp

@@ -320,6 +320,12 @@ static NaiveWarmup TryNaiveWithTimeout(const Handle& handle,
     }
     catch(...)
     {
+        // Kernels launched before the throw may still be running. Mark the end of
+        // whatever reached the stream, so the slot stays reclaimable without its
+        // callback. An event that was never recorded reads as complete, so hand
+        // it over only if this record succeeded.
+        if(hipEventRecord(ev_stop.get(), slot.stream) == hipSuccess)
+            slot.done = StreamTracker::EventPtr{std::move(ev_stop)};
         tracker.abandon(slot);
         throw;
     }
@@ -348,6 +354,10 @@ static NaiveWarmup TryNaiveWithTimeout(const Handle& handle,
         return {NaiveWarmup::Status::Completed, warmup_elapsed};
     }
 
+    // Hand the stop event over rather than destroying it: it is already recorded
+    // after this evaluation's last kernel, so it is exactly the point at which the
+    // stream becomes reusable, and testing it beats polling the stream.
+    slot.done = StreamTracker::EventPtr{std::move(ev_stop)};
     tracker.abandon(slot);
     return {NaiveWarmup::Status::TimedOut, 0.0f};
 }