[BradPepersAMD]

Summary

MIOpen has accumulated four independent, partly-overlapping mechanisms for the same
underlying problem: a solver can be applicable to a problem while being a catastrophically
bad choice for it. This PR replaces all four with a single declaration point and a single
consumption rule, and in doing so fixes a GPU driver watchdog reset (TDR) that could kill
the calling process during miopenFindConvolution*.

Net: +1081 / −1068 across 25 files.

The problem

Some solvers are deliberately applicable far outside the region where they are a sensible
choice, and this is by design:

• Naive convolution is MIOpen's universal fallback and must stay applicable at any
  problem size, or some convolutions become unsupported. Its kernel tiles within a grid,
  but a problem is still covered by a single launch, so on a large shape (VAE-decode, 3D
  convolution) that one launch runs for many seconds.
• Explicit GEMM stays applicable into shapes where im2col expansion and per-batch BLAS
  calls make it wildly uncompetitive.

Find measures every applicable solver by actually executing it. For naive this means the
OS GPU watchdog fires and resets the driver — even when a fast solver also applied and would
have won the benchmark. The solver does not have to be selected to hang the machine; being
benchmarked is enough. This is most visible on Windows, where the default TDR delay is 2 s
and the reset takes down the calling process, but the same mechanism exists on Linux.

Why the existing mechanisms are insufficient

Mechanism	Limitation
MIOPEN_CONV_DIRECT_MAX_SIZE	An output-element-count threshold that disables the entire Direct algorithm. A poor proxy for launch cost — it ignores C, K and filter extent, so it both over- and under-fires.
MIOPEN_NAIVE_TIMEOUT / _FACTOR	Bounds how long Find waits, on a separate stream. Structurally cannot prevent a TDR: the dispatch has already happened, and the kernel keeps occupying the GPU after the wait is abandoned — which is precisely what the watchdog measures.
MIOPEN_SEARCH_CUTOFF slow branch	Explicit GEMM already declares "I am off the pace here", but the consumer drops such a solver unconditionally, so it can discard the only applicable candidate and fail the convolution. Opt-in, therefore off in production.
GetWti	A per-solver hook that cannot see what else is available, so it cannot distinguish "slow, but the only option" from "slow, and something better applies".

Each either acts too late (timeout), too bluntly (MAX_SIZE, SEARCH_CUTOFF), or without
the cross-solver view the decision actually requires (GetWti).

The design

One declaration, one consumption rule.

Declare — SolverInterface::GetSpeedClass(ctx, problem) returns:

Class	Meaning
Normal	Compete on merit.
Slow	Wasteful to benchmark, but safe to launch — bounded per-batch chunks.
ExceedsLaunchBudget	A single launch is estimated to outlast the OS watchdog.

This keeps the criterion with the solver that understands it: device-derived MAC work for
naive, arch-aware shape heuristics for explicit GEMM.

Consume — one comparison, applied identically by Find's mini-benchmark and by immediate
mode's fallback: benchmark everything in the best (lowest) class that has an applicable
member; defer everything worse.

Consequences, which are the whole policy:

• Any Normal candidate exists → every deferred solver is skipped and behaviour is
  byte-identical to today.
• No Normal, several Slow → all of them are benchmarked and the least slow wins on
  measurement. Deferral ranks candidates; it does not discard a tier in favour of an
  arbitrary survivor.
• Only ExceedsLaunchBudget remains → it is benchmarked. The universal-fallback
  guarantee holds: no problem loses coverage. A huge sole-naive shape may still TDR — an
  honest "extend coverage here" signal we deliberately do not mask.
• Slow and ExceedsLaunchBudget are separate classes, not one bool. This is the
  load-bearing detail. With a single flag, "all candidates are slow" ⇒ nothing is skipped ⇒
  naive is launched anyway — and all-slow is most likely exactly when the problem is huge,
  i.e. precisely the regime the gate exists for. Separating them means a merely-wasteful
  solver outranks a watchdog risk.
• Depthwise keeps competing with no special case. Real depthwise convs
  (group == C == K ⇒ c_per_group == 1 ⇒ ~C× less work) never approach the limit, so
  naive still wins where it is genuinely fastest.

This is deliberately a pre-launch gate. Nothing that bounds waiting can avert a TDR.

The naive threshold

Rather than a magic constant, the limit is derived from the device: compute units,
wavefront width, and clock, converted through a measured MAC-per-lane-cycle rate into a work
budget for a target launch time (default 1000 ms, clamped at 60 s). It therefore scales
across the CDNA/RDNA range instead of being tuned for one part. The re-measurement recipe is
documented in-tree above the constant in conv_direct_naive_conv.cpp.

On gfx1151 this yields a 22.27 GMAC limit at the default budget; on gfx950, 117.96 GMAC.

The constant was calibrated by forcing naive and timing it over 96 shape/direction points
on gfx1151 (RDNA3.5, wave32) and gfx950 (CDNA4, wave64). The result that matters is not the
constant but the scaling: across the 39 pairs measured on both parts, gfx950 ran a median
5.04× faster on shapes ≥ 5 GMAC against a 5.30× lane-cycle ratio, so deriving the limit from
device capability holds to within ~5% over a 5× spread of part size and across both
wavefront widths.

The known gap is that MAC count is not occupancy-aware: naive's grid is
group * n * k_per_group blocks wide, so a shape whose (n, k_per_group) product is below
the part's CU count under-fills it and runs under the model. Every measured point that fell
short is of that form — 3D convs at n=1, k=64 on a 256-CU part. These are under-gated:
safe for coverage, unsafe for TDR. Recorded in-tree rather than patched, because the fix is
a different metric, not a different constant.

The constant's role is a floor, and it was checked as one: forcing naive on the six
gfx1151 shapes the calibration sweep had refused to run — up to 154.6 GMAC, 6.9× the gate
limit — put the slowest of 18 points at an implied 2.97 against the modelled 3 (see
Validation). So the limit is not optimistic anywhere measured. It is pessimistic nearly
everywhere: the other 17 points run 1.5–3.9× faster than
predicted, and the binding term is direction, which the model does not carry — forward
sustains 7.9–11.6, backward-data 6.8–8.7, weights 3.0–7.1. The constant is in effect fitted
to weights-on-3D, the one case that touches it.

Escape hatch

MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0 disables the mechanism entirely — every applicable
solver is benchmarked and offered. Intended for working around a
mis-classification (a solver deferred on a shape where it was in fact right) without a
rebuild. It re-exposes the TDR.

Verified live: with the gate on, a 128×512×512 / k=128 / 3×3 forward conv logs
Skipping last-resort (slow) Solver: miopenConvolutionFwdAlgoDirect:ConvDirectNaiveConvFwd;
with the switch off, that solver is benchmarked.

Removed

• MIOPEN_CONV_DIRECT_MAX_SIZE
• MIOPEN_NAIVE_TIMEOUT, MIOPEN_NAIVE_TIMEOUT_FACTOR, and the separate-stream evaluation
  machinery behind them (stream_tracker.hpp, handlehip.cpp stream plumbing,
  naive_timeout.cpp, stream_tracker.cpp)
• The slow-solver branch of MIOPEN_SEARCH_CUTOFF (its measured tuning/sampling bounds are
  retained unchanged)
• A dead non_slow_succeeded out-parameter that was written but never read

Known limitations

• find-db hits bypass the gate. A recorded solver is returned without consulting its
  speed class. This is accepted: a previously-recorded entry implies the benchmark already
  completed on that machine, and every alternative considered (invalidating entries,
  re-checking on load) was judged worse than the window it leaves open.
• The two retained measured-time cutoffs track wall-clock, not GPU time, so a heavily
  loaded CPU can drop a solver whose GPU time was fine. Pre-existing, deliberately not
  addressed here — it is a separate concern and a separate fix.
• The Gemm* speed-class heuristics mis-fire, and this PR is what makes that reachable.
  Their thresholds are carried over unchanged, but on develop they are consumed only
  behind env::enabled(MIOPEN_SEARCH_CUTOFF) (find_solution.hpp:332,391); here they are
  consumed unconditionally. A gfx1151 sweep found two shapes where the deferred Gemm*
  solver was in fact the faster one (see Validation). The severity is sharply asymmetric:
  when another performance solver survives the deferral the cost is ≤1.28×, but when the
  only surviving Normal candidate is naive it is 7.07×, because naive is a coverage
  fallback rather than a competitor. The heuristics were evidently fitted on 2D data —
  s = GetOutHeight() * GetOutWidth() carries no depth term, so a 3D conv's spatial work is
  undercounted by a factor of D — and GemmFwdRest's own comment quotes FPR = 11–25%. A
  false-positive rate in that range is a different proposition as a default than as an
  opt-in. Remedy is deliberately left open pending the CDNA sweep; the two candidates are
  scoping these heuristics to GetSpatialDimension() == 2, or refusing to defer when the
  surviving Normal set is naive-only.

Validation

Completed (gfx1151, Linux/WSL2)

Suite	Result
test_unit_conv_solver_naive_work_gate	16 / 16
ConvDirectNaiveConvFwd	404
ConvDirectNaiveConvBwd	293 (1 disabled)
ConvDirectNaiveConvWrw	305
GemmFwdRest / GemmBwdRest / GemmWrwUniversal	15 / 13 / 12

Plus: MIOPEN_DEBUG_DEFER_SLOW_SOLVERS on/off behaviour confirmed against driver logs;
clang-format clean (18.1.4, matching the pinned revision).

Gate on/off selection sweep, 13 shapes × 3 directions = 39 points, comparing the solver
Find selects and its benchmarked time against MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0:

	Result
Gate inert (nothing deferred)	25 points, selection identical 25/25
Gate fired	14 points, selection identical 12/14
Run-to-run spread on the inert points	median 2.7%, max 29.9% — the yardstick for the rows below

Of the 14 points where the gate fired, 7 deferred naive (0.94–1.06×, the intended
behaviour) and 7 deferred a Gemm* solver. Six of those left a Winograd solver
competing and cost 0.85–1.28×. The seventh left naive as the only candidate:

Shape	Dir	Deferred	Gate on	Gate off
conv3d_small	bwd	GemmBwdRest	6.74 ms (naive)	0.95 ms (GemmBwdRest)	7.07× slower
iso_batched	wrw	GemmWrwUniversal	23.5 ms (Winograd)	18.3 ms	1.28× slower

The 7.07× is 24× the worst inert-row excursion. The 1.28× sits inside that excursion, so it
was re-measured at 9 reps and reproduced (23.5 vs 18.3 ms) with a genuine solver change.
Both are analysed under Known limitations.

Above-the-limit check. The calibration sweep skips shapes it predicts will exceed the
budget, which leaves the threshold validated only from below. Forcing naive on all six
skipped gfx1151 shapes, smallest-first:

Shape	GMAC	×limit	predicted	measured fwd / bwd / wrw	implied rate
ramp_448	29.6	1.3	1328.8 ms	344 / 521 / 559 ms	11.6 / 7.6 / 7.1
dilate2	38.7	1.7	1735.6 ms	502 / 676 / 944 ms	10.4 / 7.7 / 5.5
conv3d_large	58.0	2.6	2603.4 ms	989 / 903 / 2630 ms	7.9 / 8.6 / 2.97
ramp_768	87.0	3.9	3905.0 ms	1086 / 1437 / 2475 ms	10.8 / 8.2 / 4.7
sdxl_vae_decode	87.0	3.9	3905.0 ms	1237 / 1713 / 2487 ms	9.5 / 6.8 / 4.7
dw_dense_ref	154.6	6.9	6942.3 ms	1922 / 2631 / 4526 ms	10.8 / 7.9 / 4.6

Three results.

The gate is firing on real hazards, not hypothetical ones. Five of these eighteen
launches meet or exceed the stock Windows 2 s TDR delay — dw_dense_ref wrw at 4526 ms is
2.26× over it, and a sixth point sits at 96% of it. These runs therefore are the failure
this PR exists to prevent, reproduced deliberately. They did not reset this host because
Linux/WSL2 has no 2 s watchdog; that is a property of the test box, not evidence of safety,
and it is why item 3 below is still required.

No point undershoots the constant over a 5× range of work above the limit, so the floor
holds. The single point that reaches it is the 3D weights conv (2629.7 ms measured against
2603.4 predicted, implied 2.97) — the documented low-occupancy case, still on the boundary
on a build that includes #7529's intra-grid spatial tiling.

The false-positive region is real: 8 of 18 points finish inside the 1000 ms budget
despite being gated, ramp_448 by 2.4–3.9×. Scored against the budget the gate is right on
10/18; scored against the thing that actually breaks, 5/18. The difference is the deliberate
~2× headroom between a 1000 ms budget and a 2 s watchdog, which the directional spread above
justifies keeping. The cost of the over-fire is a benchmark naive would lose anyway at these
sizes, so this is not an argument for raising the limit.

ramp_768 and sdxl_vae_decode are the same shape, so their six points double as an
unplanned replicate: 14% / 19% / 0.5% spread fwd / bwd / wrw.

Required before merge

1. Naive threshold on a third part. Calibrated on gfx1151 (RDNA) and gfx950 (CDNA); a
   gfx942 sweep would confirm the scaling holds on a part neither datapoint anchors.
   Separately, the gfx950 sweep ran on a driver predating Optimize naive conv solver: file split, float acc, spatial tiling, bfloat16 #7529's intra-grid spatial tiling,
   so its low-occupancy 3D outliers still need re-timing on a current build. The
   above-the-limit check weakens the hope that they are a stale-build artefact: on gfx1151
   with Optimize naive conv solver: file split, float acc, spatial tiling, bfloat16 #7529, the 3D weights case lands at an implied 2.97 — on the constant, not above
   it. Re-time it on gfx950 before treating the margin there as recovered.
2. Explicit-GEMM selection regression — confirmed, remedy outstanding. This is the
   highest-risk item and is not a TDR concern. It is no longer hypothetical: the gfx1151
   sweep above demonstrates it, so what remains is scope and fix, not existence.
   • Decide the remedy (scope the Gemm* heuristics to 2D, or refuse to defer when the
     surviving Normal set is naive-only, or re-fit the thresholds).
   • Repeat the same sweep on gfx942 and gfx950, whose thresholds differ from the
     gfx11/gfx12 branch and are therefore not covered by the gfx1151 result. Compare
     selected solver and achieved TFLOP/s, gate on vs MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0.
3. Windows TDR reproduction. Confirm the original failure no longer reproduces on a
   Windows host with the stock 2 s TDR delay.
4. Full conv CI across the supported architecture matrix.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.