Skip to content

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

Merged
cderb merged 38 commits into
developfrom
users/cderb/naive-bailout
Aug 28, 2026
Merged

cderb merged 38 commits into
developfrom
users/cderb/naive-bailout

Conversation

@cderb

@cderb cderb commented Jul 16, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Fixes a DaViT throughput regression on MI355X (gfx950) between ROCm nightly builds 0615 and 0617. The regression was caused by MIOPEN_NAIVE_DISABLE_IF_ALT, which hard-skipped the naive solver whenever a non-naive alternative existed. For DaViT's convolution shapes, naive was actually the fastest solver — skipping it forced selection of slower optimized solvers, dropping throughput by 18.8%.

JIRA ID : ROCM-28401
JIRA ID : ROCM-26410

This PR replaces the hard-skip with a wall-clock timeout mechanism (MIOPEN_NAIVE_TIMEOUT). When a non-naive solver has already succeeded, naive's warmup run is launched on a separate HIP stream with a budget of 3× the best non-naive time. If the kernel exceeds that budget, the stream is abandoned and the solver is skipped. If it finishes in time — as it does for DaViT's shapes — naive proceeds through normal benchmarking and can win the find. This restores naive as a valid competitor on shapes where it is fast, while bounding its evaluation time on shapes where optimized solvers have orders-of-magnitude advantage.

Risk Assessment

Medium risk. This changes the default behavior of the conv find/evaluation path for all architectures by replacing a hard-skip with a timeout mechanism and stream pool. The change is behind an env var (MIOPEN_NAIVE_TIMEOUT=0 to disable) and validated end-to-end on gfx950 with DaViT throughput matching pre-regression baseline. Multi-arch CI coverage is needed to confirm no regressions on other architectures.

Testing Summary

  • DaViT throughput benchmark on gfx950: confirmed regression recovery (4374.4 sps vs 4380.9 pre-regression baseline, 3557.7 regressing).
  • Stress tests on gfx908: 16 cascading timeout configs with forced and natural timeouts, GPU health check post-cascade.
  • Timeout comparison on gfx908: 16 Inception-style configs, timeout-on (88.3s) vs timeout-off (98.6s) — no overhead from the mechanism.
  • StreamTracker unit tests: 5 GTest cases covering acquire/release round-trip, pool growth, abandon with async drain, draining-while-acquired, and cascade reclaim.

Testing Checklist

  • Pre-commit hooks - pre-commit run - Status: Passed
  • DaViT throughput benchmark - train.py --throughput - ASICs: gfx950 - Status: Passed (4374.4 sps)
  • Naive timeout stress test (factor=300 + factor=1) - ASICs: gfx908 - Status: Passed
  • Timeout comparison (16 cascade configs, isolated containers) - ASICs: gfx908 - Status: Passed
  • StreamTracker unit tests (5 cases) - test_stream_tracker - Status: Passed (build verified)
  • Multi-arch sweep - TheRock multi-arch CI - ASICs: gfx90a, gfx94X, gfx950, gfx120X - Status: Pending
  • PR CI - GitHub PR checks - Status: Pending

Technical Changes

  • TryNaiveWithTimeout() in solver_finders.cpp: launches naive's warmup on a pooled HIP stream, records start/stop events, polls hipEventQuery with a wall-clock deadline, and returns elapsed time or -1 on timeout. Replaces the hard-skip that unconditionally dropped naive when non-naive alternatives existed, which caused the DaViT regression by preventing naive from winning on shapes where it was fastest.
  • StreamTracker class (stream_tracker.hpp): generic pool slot manager tracking available and draining streams. acquire() reclaims finished draining streams via non-blocking hipStreamQuery only when the available list is empty, then falls back to growing the Handle's stream pool. release() returns a slot to available; abandon() moves it to draining.
  • StreamTracker is owned by the Handle (via HandleImpl), eliminating global state and making pool state naturally per-handle.
  • MIOPEN_NAIVE_TIMEOUT (bool, default: true) enables/disables the timeout mechanism. MIOPEN_NAIVE_TIMEOUT_FACTOR (uint64, default: 300) sets the budget as a percentage of the best non-naive time (300 = 3×).
  • Updated existing naive_disable_if_alt.cpp tests to cover the new timeout semantics (paths A, B, D through EvaluateInvokers).

🤖 Generated with Claude Code

cderb and others added 3 commits July 15, 2026 13:12
…am bailout

Replace the hard skip of naive convolution solvers (MIOPEN_NAIVE_DISABLE_IF_ALT)
with a timeout-based approach (MIOPEN_NAIVE_TIMEOUT). When a non-naive solver has
already succeeded, the naive solver's warmup is launched on a disposable HIP
stream with a wall-clock budget of 3x the best non-naive time. If the warmup
exceeds this budget, the stream is abandoned to a detached background thread for
async cleanup and the solver is skipped. If it finishes in time, normal timed
benchmarking proceeds on the root stream.

This preserves the find-time savings of deferring slow naive solvers while still
allowing naive to win when it is genuinely the fastest option.

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
Move the naive timeout stream pool tracker out of solver_finders.cpp
into a reusable StreamTracker class owned by the Handle. Eliminates a
leaked global map keyed by Handle pointer and makes pool state naturally
per-handle with no cross-thread concerns.

Also registers MIOPEN_NAIVE_TIMEOUT_FACTOR in the env var system and
adds five GTest cases exercising acquire/release, pool growth, abandon
with async drain, and cascade reclaim.

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
…bailout

# Conflicts:
#	projects/miopen/CHANGELOG.md
@therock-pr-bot

therock-pr-bot Bot commented Jul 16, 2026 •

Copy link
Copy Markdown

✅ 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

therock-pr-bot Bot commented Jul 16, 2026 •

Copy link
Copy Markdown

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

- Check hipFree return value (nodiscard)
- Cast hipMemsetAsync to void in fire-and-forget loop
- Remove unused variable from emplace_back
- Rename stream_tracker.cpp → stream_tracker_test.cpp to satisfy PR bot unit test naming convention

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
@therock-pr-bot

Copy link
Copy Markdown

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
@cderb
cderb marked this pull request as ready for review July 16, 2026 19:46
@cderb
cderb requested review from a team as code owners July 16, 2026 19:46
cderb and others added 3 commits July 16, 2026 15:58
Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
GPU+NONE is rejected by check_names.py (line 98). GPU tests require
a real datatype suffix.

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>

@jdcampbe jdcampbe left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think it looks good, but I would add the flagging interactions into the docs.

cderb and others added 3 commits July 27, 2026 10:01
…FACTOR env vars

Replace removed MIOPEN_NAIVE_DISABLE_IF_ALT with the two new env vars
and document their interaction as requested by reviewer.

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- Use per-test StreamTracker instances to avoid cross-test state leakage
- Replace racey large-memset with hipLaunchHostFunc barrier to
  deterministically hold the stream busy in AbandonStillDraining
- Defer releasing available slot so acquire is forced to scan draining

Verified: 5/5 tests pass on gfx942 (MI300X).

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>

@JonathanLichtnerAMD JonathanLichtnerAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Claude has found a few issues with this that need addressing.

Once that's fixed, I think we would want to run these tests with ASAN (both host and device), and we should add an additional test where naive runs for a long time and gets short-circuited.

Also, one last concern I have with this fix is that naive still runs, and there's cases where naive runs for 30+ seconds. This consumes gpu time and might affect multi-GPU runs. It might be simpler to implement a filter that predicts when naive will take over a certain amount of time, e.g., 1 second, or 0.5 second, or whatever seems reasonable. We could also test this filter using the find DB to see how accurate it is.

Comment thread projects/miopen/CHANGELOG.md Resolved
…eWithTimeout

When a naive kernel times out and its stream is abandoned, the kernel
continues writing to the shared output buffer while the next solver
reuses those same buffers on stream 0. Redirect naive output to a
per-Handle scratch buffer (capped at 1/8 device memory) so abandoned
writes are harmless. Scratch lifetime is caller-driven via shared_ptr
attached to StreamTracker slots — RAII frees when the last user finishes.

Also adds RAII exception safety for stream/profiling restore, renames
test files to match convention, and moves CHANGELOG entry to ROCm 10.1.

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
@cderb
cderb enabled auto-merge (squash) August 14, 2026 16:50
cderb and others added 13 commits August 18, 2026 10:39
…rchs

The StreamTracker and NaiveTimeout suites are plain TEST/TEST_F with no
INSTANTIATE_TEST_SUITE_P, so they carry no "Prefix/" and none of the
existing "*/GPU_*" patterns in test_categories.yaml could match them.
19 of the 20 new tests were silently never executed; only the CPU case
ran, via the blanket "CPU_*" pattern. Add unprefixed patterns so the
standard/comprehensive/full suites pick all 20 up.

Now that they actually run, guard the two assertions that are only
meaningful where ConvDirectNaive is a candidate. On gfx803/gfx900/
gfx906/gfx908, ConvBinWinograd3x3U is applicable to the test's
convolution, so IsWinograd3x3SupportedAndFast() sets use_winograd_only
and the Direct finder is disabled outright -- naive presence or absence
then says nothing about the timeout.

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
@cderb
cderb merged commit f28c5d2 into develop Aug 28, 2026
87 of 90 checks passed
@cderb
cderb deleted the users/cderb/naive-bailout branch August 28, 2026 23:44
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

@stoivonen-amd

Copy link
Copy Markdown

Some our scripts benchmarking models ran into an issue because of the lazy cleanup added in this change:
#12121

BradPepersAMD added a commit that referenced this pull request Sep 30, 2026
MIOpen had four partly-overlapping mechanisms for "this solver is applicable
but is a catastrophically bad choice here": MIOPEN_CONV_DIRECT_MAX_SIZE,
MIOPEN_NAIVE_TIMEOUT, the MIOPEN_SEARCH_CUTOFF slow branch, and GetWti. Each
acted too late, too bluntly, or without the cross-solver view the decision
needs. Replace all four with one declaration point and one consumption rule.

Declare: SolverInterface::GetSpeedClass(ctx, problem) returns Normal, Slow, or
ExceedsLaunchBudget. Consume: Find's mini-benchmark and immediate mode's
fallback both benchmark everything in the lowest class with an applicable
member and defer the rest.

Slow and ExceedsLaunchBudget are separate classes rather than one bool. With a
single flag, "all candidates are slow" means nothing is skipped and 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.

This is a pre-launch gate. Bounding how long Find *waits* cannot avert a TDR;
the dispatch has already happened. A sole-candidate ExceedsLaunchBudget solver
is still benchmarked, so no problem loses coverage.

The naive threshold is derived from device capability (CUs, wavefront width,
clock) through a measured MAC-per-lane-cycle rate, calibrated over 109 timed
shape/direction points on gfx1151 (RDNA3.5, wave32) and gfx950 (CDNA4,
wave64). Across the 39 pairs measured on both, gfx950 ran a median 5.04x
faster against a 5.30x lane-cycle ratio, so the scaling holds to within ~5%.

Removed: MIOPEN_CONV_DIRECT_MAX_SIZE; MIOPEN_NAIVE_TIMEOUT and its
separate-stream machinery (stream_tracker, naive_timeout); the SEARCH_CUTOFF
slow branch; a dead non_slow_succeeded out-parameter.

Escape hatch: MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0 restores prior behaviour.

Relates to #3509, #1860, #9513.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
@lostdisc

lostdisc commented Oct 6, 2026

Copy link
Copy Markdown

It looks like this change came out with ROCm 10.1. The removal of MIOPEN_NAIVE_DISABLE_IF_ALT hits gfx1150 pretty hard for first-time SDXL VAE decode of larger images. For a 1600x1280 image, my first-time decode went from 1 minute on ROCm 7.14/10, to about 5 minutes on 10.1. There's also a few minutes of benching prior to the start of sampling, making the overall 20-step gen take 10 minutes the first time (versus ~1.5 mins for subsequent runs).

Prior to ROCm 7.14 adding MIOPEN_NAIVE_DISABLE_IF_ALT, the same decode took a brutal ~37 minutes. ComfyUI still has MIOpen disabled by default thanks to those dark days. (Technically for about a year now, although an env var was added for enabling it.)

I see that @BradPepersAMD has further reworked and consolidated the slow-solver handling last week with 769afc1 and 669bd62. Hope to recover some speed again in 10.2!

@jdcampbe

jdcampbe commented Oct 6, 2026

Copy link
Copy Markdown
Contributor

It looks like this change came out with ROCm 10.1. The removal of MIOPEN_NAIVE_DISABLE_IF_ALT hits gfx1150 pretty hard for first-time SDXL VAE decode of larger images. For a 1600x1280 image, my first-time decode went from 1 minute on ROCm 7.14/10, to about 5 minutes on 10.1. There's also a few minutes of benching prior to the start of sampling, making the overall 20-step gen take 10 minutes the first time (versus ~1.5 mins for subsequent runs).

Prior to ROCm 7.14 adding MIOPEN_NAIVE_DISABLE_IF_ALT, the same decode took a brutal ~37 minutes. ComfyUI still has MIOpen disabled by default thanks to those dark days. (Technically for about a year now, although an env var was added for enabling it.)

I see that @BradPepersAMD has further reworked and consolidated the slow-solver handling last week with 769afc1 and 669bd62. Hope to recover some speed again in 10.2!

We had to go a different route with tuning naive because the MIOPEN_NAIVE_DISABLE_IF_ALT caused some common workloads to slow down. There is a merged change to naive (#7529) that should improve the naive solver's performance, and we are working on a new approach to how/when we select naive during tuning.

cderb added a commit that referenced this pull request Nov 2, 2026
…ptures (#12265)

## Summary

#9513 evaluates naive conv solvers on a disposable stream and, on
timeout, abandons it
still running. The scratch it pins was reclaimed lazily from
`Handle::Run`, so a later
kernel launch could poll the abandoned stream with `hipStreamQuery`. HIP
rejects that
call during a graph capture and the rejection invalidates the capture —
so a stream
abandoned by one find destroyed an unrelated later convolution's
capture. Any framework
capturing MIOpen convolutions into a HIP graph is affected.

This PR takes reclaim off the launch path entirely. An abandoned
evaluation's stream is
retired by the runtime: `abandon()` queues a host callback behind the
evaluation's work,
and a worker thread owned by the stream tracker returns the slot to the
pool and frees
its scratch as soon as the kernel retires. Nothing runs inside
`Handle::Run` any more,
and the scratch is released even if the caller never calls MIOpen again
— including a
workload that only replays a captured graph.

## Risk Assessment

Medium risk. Changes the conv find/evaluation path for all architectures
and adds a
worker thread per handle that has abandoned a naive evaluation. The
mechanism is
host-side stream, event, and scratch bookkeeping with no kernel or
solver-selection
change, so it is architecture-independent. `MIOPEN_NAIVE_TIMEOUT=0`
disables the whole
path without a rebuild. Not yet validated on gfx1151, the platform that
reported #12121.

A deallocator installed with `miopenSetAllocator` is now called from the
tracker's worker
thread rather than the caller's; see Known limitations.

## Related

- Fixes: #12121
- Caused by: #9513
- Jira ID: ROCM-33231

## Device / Architecture Coverage

Architecture-independent: host-side stream/scratch bookkeeping, no
kernel code, no
solver-selection or support-surface change. Standard multi-arch PR CI is
the right bar.
The exception is gfx1151 — the reporting platform in #12121 — where the
reporter's repro
should be confirmed before merge.

The new end-to-end test captures whichever solution find ranks best. On
gfx1100 that is
an assembly Winograd kernel; multi-arch CI is the first run on other
architectures.

## Testing Summary

- `test_naive_timeout` (6 cases, 1 new):
`CaptureAfterATimedOutNaiveEvaluationSucceeds`
runs a find that abandons a naive evaluation, then captures the
convolution on a
capturable stream in global and thread-local mode. Against the pre-PR
library (develop
  at the merge base, same test file) it fails in both modes with
  `hipErrorStreamCaptureInvalidated` (901); on this branch it passes.
- `test_stream_tracker` (26 cases): release with no further MIOpen call,
release during
an unrelated global capture, late callbacks and tickets, reclaim at the
backlog cap
with callbacks held, shared scratch, clean release and reuse, teardown,
concurrent
finds, deallocation kept out of `acquire()` and off the tracker lock,
and the backlog
  cap derived from `GPU_MAX_HW_QUEUES`.
- Each guard was checked by reintroducing the regression it catches and
confirming the
  corresponding test fails.
- Determinism: both suites pass repeated (×20) in shuffled order, with
and without all
CPU cores saturated, and with `GPU_MAX_HW_QUEUES=3`. Gated tests open
their gates on
  any exit, so a failed assertion reports instead of hanging.
- Offline, not part of CI: HIP's sticky device error, a callback that
cannot be queued,
and a failed naive launch inside find were exercised with an
injected-failure shim.
A ThreadSanitizer build reports only the hand-off through
`hipLaunchHostFunc`, whose
ordering lives in the uninstrumented HIP runtime; nothing on the
tracker's
  lock-protected state.

## Testing Checklist

- [x] `test_stream_tracker` - `./bin/test_stream_tracker` - Devices:
gfx1100 - Status: Passed (26/26)
- [x] `test_naive_timeout` - `./bin/test_naive_timeout` - Devices:
gfx1100 - Status: Passed (6/6)
- [x] #12121 regression against the pre-PR library - Devices: gfx1100 -
Status: Fails pre-PR, passes here
- [x] clang-format 18.1.4 - `clang-format -style=file` - Status: Passed
- [ ] Reporter's repro from #12121 - Devices: gfx1151 - Status: Not run
(no access)
- [ ] PR CI - GitHub PR checks - Status: Pending

## Flags / Guardrails

`MIOPEN_NAIVE_TIMEOUT=0` is a complete runtime kill switch for
everything this PR
touches. It makes `defer_naive` — and so `cutoff_naive` — false, so
`TryNaiveWithTimeout`
never runs, no stream is ever abandoned, no callback is queued, and the
worker thread is
never started.

The backlog of abandoned streams is capped at `GPU_MAX_HW_QUEUES` minus
2, at least 1
(2 at the HIP default of 4). An unset, non-positive, or unparseable
value counts as the
HIP default. `MIOPEN_NAIVE_MAX_DRAINING`, added earlier in this PR, has
been removed.

## Performance Impact

- `Handle::Run`: no tracker work at all. #9513 called `sweep()` on every
launch; this PR
  removes it.
- Per abandoned evaluation, measured on gfx1100: queuing the callback
costs ~9 µs on the
find thread; the callback runs ~13 µs after the stream reaches it, and
the worker wakes
  ~4 µs later.
- With the default allocator, freeing scratch (`hipFree`) waits for
every stream on the
device. That wait now happens on the worker thread, not in a caller's
kernel launch.
- `acquire()` at the backlog cap waits on a condition variable signalled
by the worker
  instead of polling.

## Adjacent Tests Considered

`test_naive_timeout` covers the find path changed in
`solver_finders.cpp` and was run.
Conv correctness suites were not: no solver, kernel, or selection logic
changed, and the
scratch involved is a write-only warmup destination that is never read
back.

## Work Tracking

Tracked by public issue #12121 and Jira ROCM-33231.

## Technical Changes

- `abandon()` queues a `hipLaunchHostFunc` callback behind the slot's
work. The callback
makes no HIP call — a deallocator's `hipFree` there would wait on the
stream running
it — and only hands a per-abandonment ticket to a worker thread owned by
the tracker.
The worker returns the slot to the pool and frees its scratch outside
the tracker lock,
under a relaxed capture-mode guard so it cannot invalidate another
thread's capture.
- Tickets keep a callback that arrives late from retiring a later reuse
of the same
  stream.
- `acquire()` still reclaims on the slot's done event (`SweepLocked`)
before waiting,
and again on each wait timeout. That is the backstop for a callback that
is late — the
runtime runs host callbacks one at a time — or never arrives. Reclaimed
scratch goes
  to the worker; `acquire()` never runs a deallocator.
- Only `hipSuccess` retires a slot. `hipErrorNotReady` keeps waiting;
any other answer is
HIP's sticky device error, and `acquire()` throws rather than waiting on
a backlog that
  can never drain.
- `TryNaiveWithTimeout` hands its stop event to the slot as the done
event on timeout,
and on the error path if recording it succeeds, so every abandonment in
find carries
  one.
- `release()` returns a slot clean (scratch, done event, ticket), and
the destructor
synchronizes every owned stream, lets every callback run, then joins the
worker.
- `Handle::GetScratchBuffer` is guarded by a mutex.
- `MIOPEN_NAIVE_MAX_DRAINING` removed; the cap derives from
`GPU_MAX_HW_QUEUES` via
  `StreamTracker::MaxDrainingFor`.

### Known limitations

- A deallocator installed with `miopenSetAllocator` is called from the
tracker's worker
thread, possibly while another thread is capturing. It must not make
calls HIP refuses
during any capture regardless of mode, such as a synchronous `hipMemcpy`
or
`hipMemset`: those invalidate every capture in the process. Documented
in
  `docs/reference/env_variables.rst`.
- Host callbacks run one at a time process-wide, so a blocking host
function elsewhere in
the process delays release. A delayed or lost callback's slot is still
reclaimed at
the next `acquire()` through its done event, but its scratch stays
pinned until then.
- Find itself is not safe to run while another thread captures in global
mode; its
solver timing synchronizes events. That predates #9513 and is unchanged
here.

---------

Co-authored-by: Claude <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants