refactor(miopen): unify slow-solver handling behind a single speed class, fixing conv Find watchdog resets (TDR)

projects/miopen/CHANGELOG.md

@@ -8,8 +8,18 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 ### Added
 * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise
   kernels to the `ConvHipConv` solver (hipconv v0.3.1).
+* [Conv] Added the `MIOPEN_DEBUG_DEFER_SLOW_SOLVERS` environment variable (default on). Setting it to 0 disables speed-class based solver deferral entirely, so find and immediate mode consider every applicable solver as they did before that mechanism existed. It is an escape hatch for a solver being deferred on a shape where it was in fact the right choice, and re-exposes the naive-kernel watchdog reset (TDR) the mechanism prevents.
+
+### Changed
+* [Conv] Solvers that are applicable but known to be far off the pace for a given problem now declare this through `SolverInterface::GetSpeedClass`, which returns `Normal`, `Slow` (wasteful to benchmark, but safe to launch) or `ExceedsLaunchBudget` (a single launch is estimated to outlast the OS GPU watchdog). Both selection paths honor it: find's mini-benchmark and immediate mode's fallback now evaluate only the best class that has an applicable member and defer the rest. Ranking is preserved within a class, so when several equally-safe candidates remain they are all benchmarked and the least slow wins; separating `Slow` from `ExceedsLaunchBudget` is what keeps a watchdog risk from being launched merely because the alternative was also off the pace. A solver is never dropped when it is the only candidate, so no problem loses coverage. The naive convolution solvers report `ExceedsLaunchBudget` via the new work gate below; the explicit-GEMM solvers report `Slow` via their existing heuristics, which previously took effect only under the opt-in `MIOPEN_SEARCH_CUTOFF`.
+
+### Resolved Issues
+* [Conv] Fixed a GPU driver watchdog reset (TDR) that could occur while find benchmarked the un-tiled naive convolution solver on very large problems (e.g. VAE-decode and 3D convolutions). The three naive solvers now report `SolverSpeedClass::ExceedsLaunchBudget` when the problem's total MAC work exceeds what the device can retire inside a launch-time budget, so they are skipped during find and immediate mode whenever any better-classed solver applies. The limit is derived from device compute units, wavefront width, and clock rather than a fixed constant, so it scales across the CDNA/RDNA range; override the budget with `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS` (default 1000 ms) or pin the MAC limit outright with `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK`. Naive still runs when it is the only applicable solver, and small problems where it is fastest continue to compete.
 
 ### Removed
+* [Conv] Removed the `MIOPEN_CONV_DIRECT_MAX_SIZE` environment variable (an output-element-count threshold that disabled the entire Direct algorithm). It was a poor proxy for kernel launch cost and is superseded by the naive-conv work gate described above.
+* [Conv] Removed the `MIOPEN_NAIVE_TIMEOUT` and `MIOPEN_NAIVE_TIMEOUT_FACTOR` environment variables along with the separate-stream evaluation machinery behind them. That deferral measured the naive warmup only after dispatching it, so it could not prevent the launch that causes a TDR; the work gate above declines the launch up front and makes the timeout path unreachable.
+* [Conv] Removed the slow-solver handling from `MIOPEN_SEARCH_CUTOFF`, which dropped a slow solver unconditionally and could therefore discard the only applicable candidate. That declaration is now `GetSpeedClass` and is consumed by the last-resort policy described above, which is on by default and never drops a sole candidate. `MIOPEN_SEARCH_CUTOFF` continues to bound measured tuning and sampling time.
 * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources.
 * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`.
 * [Conv] Removed the gfx803 find-database files.

projects/miopen/docs/reference/env_variables.rst

@@ -104,39 +104,39 @@ and :doc:`Performance database <../conceptual/perfdb>`.
       - | 1: Enable
         | 0 or unset: Disable
 
-    * - | ``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_CONV_DIRECT_NAIVE_MAX_TIME_MS``
+        | Wall-clock budget, in milliseconds, for a single launch of the
+        | un-tiled naive convolution kernel. The naive solvers convert
+        | this budget into a MAC limit using the device's compute units,
+        | wavefront width, and clock, then report the speed class
+        | ``ExceedsLaunchBudget`` for any problem exceeding it, which
+        | keeps find and immediate mode from selecting them while any
+        | better-classed solver applies. Raise it to let
+        | naive compete on larger problems, at the risk of a GPU driver
+        | watchdog reset (TDR); lower it on a device with a short
+        | watchdog. Values above 60000 are clamped.
+      - | Milliseconds (default: 1000)
+
+    * - | ``MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK``
+        | Pins the naive convolution MAC limit to a fixed value instead
+        | of deriving it from the device and the time budget above. Takes
+        | precedence over ``MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS``.
+        | Intended for testing and for bisecting the derived limit.
+      - | MAC count (unset: derive from device)
+
+    * - | ``MIOPEN_DEBUG_DEFER_SLOW_SOLVERS``
+        | Master switch for speed-class based solver deferral. When on,
+        | each solver reports a speed class for the problem at hand and
+        | find and immediate mode consider only the best class that has
+        | an applicable member, so a solver that is merely applicable is
+        | not benchmarked or offered while a better-classed one exists.
+        | Set to 0 to consider every applicable solver regardless of its
+        | speed class, as releases before this mechanism did. That is a
+        | workaround for a solver being deferred on a shape where it was
+        | in fact the right choice; it also re-exposes the naive-kernel
+        | GPU driver watchdog reset (TDR) the mechanism prevents.
+      - | 1 or unset: Defer slow solvers
+        | 0: Consider all applicable solvers
 
     * - | ``MIOPEN_DEBUG_DISABLE_FIND_DB``
         | Disables FindDb functionality.