fix(miopen): keep naive-evaluation stream reclaim out of HIP graph captures by cderb · Pull Request #12265 · ROCm/rocm-libraries

cderb · 2026-09-17T19:34:23Z

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.

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

test_stream_tracker - ./bin/test_stream_tracker - Devices: gfx1100 - Status: Passed (26/26)
test_naive_timeout - ./bin/test_naive_timeout - Devices: gfx1100 - Status: Passed (6/6)
[Issue]: MIOpen lazily reclaims abandoned find streams on later kernel launches, breaking unrelated HIP graph captures #12121 regression against the pre-PR library - Devices: gfx1100 - Status: Fails pre-PR, passes here
clang-format 18.1.4 - clang-format -style=file - Status: Passed
Reporter's repro from [Issue]: MIOpen lazily reclaims abandoned find streams on later kernel launches, breaking unrelated HIP graph captures #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. feat(miopen): time-bound naive solver evaluation during find #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 feat(miopen): time-bound naive solver evaluation during find #9513 and is unchanged here.

…ptures #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 any later kernel launch might poll the abandoned stream with hipStreamQuery. HIP rejects that call while a capture is in flight and the rejection invalidates the capture, so a stream abandoned by one find could destroy an unrelated later convolution's graph capture. Closes #12121. Reclaiming now goes through SweepUnlessCapturing on the launch path, which skips while the handle's stream is recording. Keeping it on that path matters: it is the only event frequent enough to return the buffer promptly, since a find may never run again and frameworks that supply their own workspace never reach Handle::Create. A relaxed atomic counting slots that still pin a buffer gates the check, so a launch with nothing abandoned takes no lock and makes no HIP call. ScopedRelaxedCaptureMode covers what the stream check cannot see - a capture on a different stream than the handle's, or on another thread, which under the default global mode still makes our queries illegal and kills that capture when they are refused. Also in the same code: - acquire() waits for an abandoned stream rather than declining. Declining would skip a solver evaluation, and a skipped naive solver can lose a find it would have won, which is what MIOPEN_NAIVE_TIMEOUT exists to prevent. Find results reach the user find-db, so a transient backlog would persist as a worse recorded solver choice. - The backlog cap derives from GPU_MAX_HW_QUEUES, the same variable ROCclr uses to decide when streams start sharing hardware rings, so an evaluation is not left timing ring contention from an earlier runaway kernel and writing that number to the find-db. MIOPEN_NAIVE_MAX_DRAINING overrides it. - Reclaim tests the stop event TryNaiveWithTimeout already records after the evaluation's last kernel instead of polling the stream, and only ever releases a buffer once that reports the kernel retired. - StreamTracker takes a mutex. sweep() previously mutated three vectors from the kernel launch path with no lock, and MIOPEN_HANDLE_LOCK is a no-op unless MIOPEN_GPU_SYNC. - The destructor synchronizes every owned stream, not only the draining ones; release() does not require a slot to be quiescent. Tests cover the capture regression, the backlog wait, the event-based reclaim, and the launch-path gate. Each was checked to fail with its guard removed. Co-Authored-By: Claude <noreply@anthropic.com>

The sibling MIOPEN_NAIVE_TIMEOUT variables are listed here, so the cap added alongside them belongs in the same table. Records what it bounds, how the default is derived from GPU_MAX_HW_QUEUES, and why raising it needs evidence. Co-Authored-By: Claude <noreply@anthropic.com>

therock-pr-bot · 2026-09-17T19:39:46Z

✅ 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 · 2026-09-17T19:39:50Z

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

MIOpen regression, fix: ROCm/rocm-libraries#12265 Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

JonathanLichtnerAMD

Thanks for tracking this down. The approach of skipping the reclaim while capturing, plus a relaxed-capture guard around the queries, looks right to me. I found no blockers.

I built the PR head and a pre-PR baseline and ran a standalone repro. It does find on an empty find-db, so the naive solver times out and abandons a stream. It then captures miopenConvolutionForwardImmediate and calls hipStreamEndCapture. On the pre-PR build the capture fails in both thread-local and global mode, matching #12121 and ROCM-33231. On the PR head it succeeds in both modes.

Items I'd like addressed:

major acquire() can wait forever if a query returns any error other than hipErrorNotReady (#1).
major Nothing tests the Handle::Run reclaim guard, which is the actual #12121 fix (#3).
minor In a tune -> capture -> replay workload, scratch pinned at capture time is never released (#6). Gated tests hang when the backlog cap is 1 (#2, unlikely test failure). Negative env values wrap (#4). Changelog entry (#5).
question What MIOPEN_NAIVE_MAX_DRAINING adds beyond GPU_MAX_HW_QUEUES (#7).

5. Changelog (general, minor): This probably deserves an entry in projects/miopen/CHANGELOG.md.

Details on the others are inline.

These comments were generated by Claude (AI) at Jonathan Lichtner's direction, and were reviewed and approved by him before posting. Claude reproduced the following on gfx942 (MI300X, HIP runtime 71726416) by building and running them: the end-to-end capture failure and its fix, the cap=1 test hang, the negative-value wraparound, the Handle::Run coverage gap, and the replay-only scratch retention (#6). The acquire() hang was reproduced only with an injected query error, not with a real device fault. The proposed fixes for the acquire() hang and the host-callback reclaim in #6 were tested there too; the host-callback prototype passed test_naive_timeout and test_stream_tracker except for three existing tests that assert synchronous release. The remaining comments (changelog, question) were not empirically tested. I did not run the gfx1151/gfx950 repros or a PyTorch/XLA repro, and I did not test the host-callback approach with framework allocators.

- Retire an abandoned slot only on hipSuccess. hipErrorNotReady keeps waiting; any other answer is HIP's sticky device error, so acquire() throws instead of waiting on a backlog that can never drain. The slot and its scratch stay put: an error does not prove the kernel stopped. - Drop MIOPEN_NAIVE_MAX_DRAINING. The backlog cap is a per-tracker constructor argument defaulting to DefaultMaxDraining(), derived from GPU_MAX_HW_QUEUES, which is now parsed as signed and ignored when <= 0. - Make the gated tests independent of GPU_MAX_HW_QUEUES with a fixed test cap, and cover the Handle::Run capture guard and Handle::Create reclaim directly. - Document that pinned scratch is released at the next uncaptured launch, allocation, or find, never by graph replay. - Add a CHANGELOG entry for #12121.

…llback Scratch pinned by an abandoned naive evaluation was released only by a later Handle::Run or Handle::Create, so a workload that finds, captures, and then only replays the graph never released it, and with the default allocator the release's hipFree stalled the user's launch until every stream on the device drained. - abandon() queues a hipLaunchHostFunc callback behind the slot's work. The callback makes no HIP call; it hands a per-abandonment ticket to a worker thread owned by the tracker, which returns the slot to the pool and drops its scratch outside the tracker lock under the relaxed capture-mode guard. Tickets keep a late callback from retiring a later reuse of the same stream. - acquire() waits on a condition variable signalled by the worker, and re-tests the backlog every 10 ms for a callback that never arrives. SweepLocked returns on the first answer other than success or not-ready, HIP's sticky device error, which acquire() throws. - Remove the launch-path reclaim from Handle::Run and the sweep from Handle::Create, along with sweep(), SweepUnlessCapturing, HasPinnedScratch, and the pinned-scratch counter. - Test that scratch is released once its kernel retires with no further call, and that the release leaves an unrelated global-mode capture valid. Document that a miopenSetAllocator deallocator is now called from the worker thread.

Tracker: - SweepLocked hands reclaimed scratch to the worker instead of freeing it under the tracker lock inside find, where the default hipFree would wait for every other abandoned kernel. All deallocation now happens on the worker, outside the lock. - Document SweepLocked as the backstop for a late or lost callback; with the tickets, it guarantees a finished slot is reclaimed whatever happens to its callback. - release() returns a slot clean: scratch dropped outside the lock, done event and ticket cleared, so a stale done event cannot read as complete on the stream's next abandonment. - The destructor reads the backlog size under the lock, and the worker signals waiters only when it has moved a slot. - Move the backlog-cap parsing into MaxDrainingFor() so it can be tested; DefaultMaxDraining() reads GPU_MAX_HW_QUEUES once and calls it. Find and handle: - TryNaiveWithTimeout records its stop event in the catch path and hands it over as the done event, only if the record succeeded, so a slot abandoned on error is still reclaimable without its callback. - Guard Handle::GetScratchBuffer's cache with a mutex. Tests: - test_naive_timeout: CaptureAfterATimedOutNaiveEvaluationSucceeds, an end-to-end regression for #12121 in global and thread-local capture modes; fails against the pre-PR library, passes here. - test_stream_tracker: late callbacks and tickets, reclaim at the cap with held callbacks, shared scratch, clean release and reuse, teardown, concurrent finds, deallocation kept out of acquire() and off the lock, and the backlog cap table. - Gates are owned by the fixture and opened at teardown, and a scope guard opens them before a local buffer is freed, so a failed assertion fails instead of hanging. Timing thresholds leave wide margins.

…nt completion - abandon(): take the retire context out of its unique_ptr before queuing the callback, and re-own it if queuing fails, instead of discarding release()'s result. clang-tidy 23 flags the (void) discard as bugprone-unused-return-value. - AcquireLeavesTheDeallocatorToTheWorker and ALateCallbackDoesNotRetireALaterReuseOfItsStream assumed a done event recorded on an idle stream is complete by the next call. On Windows it can trail the record, so acquire()'s single up-front reclaim missed it and the tests failed. Both now wait for the event to complete first.

JonathanLichtnerAMD

LGTM with one comment on where the changelog is located (should be under 3.6.3/10.3.0)

…stream_reclaim

…pen 3.6.3

mattwill-amd · 2026-10-30T10:42:39Z

/docs-preview

assistant-librarian · 2026-10-30T10:43:20Z

Docs preview for miopen is building on Read the Docs: https://rocm.docs.amd.com/projects/MIOpen/en/users-cderb-naive_stream_reclaim/ (Note: If a pull request includes multiple projects, only the preview for the first project will be built.)

assistant-librarian · 2026-11-03T11:06:30Z

TheRock Submodule Bump Activity

Newest first

2026-11-03 — Included in TheRock via ROCm/TheRock#9698.

MIOpen regression, fix: ROCm/rocm-libraries#12265 (cherry picked from commit 5e01501) Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com>

* hlo_eval_tools: note the CLR fix for the command-buffer failures - The models that require CMD_BUFFER=off are fixed by CLR commit a9dad6342cd (ROCm/rocm-systems#12274), in TheRock ROCm nightlies since 10.3.0a20261103. - With MIOpen ROCm/rocm-libraries#12265 (since 10.3.0a20261104), every model in that table runs with command buffers on (gfx950). * hlo_eval_tools: say which nightlies the CLR fix note means - TheRock nightlies are development builds of TheRock main (rocm-systems and rocm-libraries develop), not a 10.3 release candidate (review question).

cderb requested a review from a team as a code owner September 17, 2026 19:34

github-actions Bot added the project: miopen label Sep 17, 2026

assistant-librarian Bot added the organization: ROCm label Sep 17, 2026

cderb requested a review from a team as a code owner September 17, 2026 19:40

github-actions Bot added the documentation label Sep 17, 2026

This was referenced Sep 24, 2026

[release/2.14] Skip test_ddp_optimizer_cudagraph on ROCm ROCm/pytorch#3708

Merged

[release/2.13] Skip test_ddp_optimizer_cudagraph on ROCm ROCm/pytorch#3709

Merged

[release/2.12] Skip test_ddp_optimizer_cudagraph on ROCm ROCm/pytorch#3710

Merged

pragupta pushed a commit to ROCm/pytorch that referenced this pull request Sep 24, 2026

[release/2.12] Skip test_ddp_optimizer_cudagraph on ROCm (#3710)

a55f030

MIOpen regression, fix: ROCm/rocm-libraries#12265 Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

pragupta pushed a commit to ROCm/pytorch that referenced this pull request Sep 24, 2026

[release/2.13] Skip test_ddp_optimizer_cudagraph on ROCm (#3709)

d808a54

MIOpen regression, fix: ROCm/rocm-libraries#12265 Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

pragupta pushed a commit to ROCm/pytorch that referenced this pull request Sep 24, 2026

[release/2.14] Skip test_ddp_optimizer_cudagraph on ROCm (#3708)

5e01501

MIOpen regression, fix: ROCm/rocm-libraries#12265 Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Merge branch 'develop' into users/cderb/naive_stream_reclaim

7a21b53

JonathanLichtnerAMD reviewed Oct 24, 2026

View reviewed changes

cderb added 4 commits October 26, 2026 12:57

JonathanLichtnerAMD approved these changes Oct 28, 2026

View reviewed changes

Comment thread projects/miopen/CHANGELOG.md Resolved

cderb added 2 commits October 29, 2026 10:42

Merge remote-tracking branch 'origin/develop' into users/cderb/naive_…

5f131af

…stream_reclaim

docs(miopen): move the naive stream reclaim changelog entry under MIO…

e191b73

…pen 3.6.3

mattwill-amd approved these changes Oct 30, 2026

View reviewed changes

magaonka-amd mentioned this pull request Nov 2, 2026

Create DNN handles before command buffers capture convolutions ROCm/xla#1387

Open

cderb merged commit 9953cb5 into develop Nov 2, 2026
103 checks passed

cderb deleted the users/cderb/naive_stream_reclaim branch November 2, 2026 16:42

albmalamd mentioned this pull request Nov 4, 2026

[ROCm] Relax HIP stream capture mode around MIOpen calls pytorch/pytorch#203370

Draft

magaonka-amd mentioned this pull request Nov 6, 2026

hlo_eval_tools: note the CLR fix for the command-buffer failures ROCm/xla#1407

Merged

Navigation Menu

Repository navigation

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

fix(miopen): keep naive-evaluation stream reclaim out of HIP graph captures#12265
cderb merged 9 commits into
developfrom
users/cderb/naive_stream_reclaim

cderb commented Sep 17, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 17, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 17, 2026

Uh oh!

JonathanLichtnerAMD left a comment •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

JonathanLichtnerAMD left a comment

Uh oh!

Uh oh!

mattwill-amd commented Oct 30, 2026

Uh oh!

assistant-librarian Bot commented Oct 30, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Nov 3, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

cderb commented Sep 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Risk Assessment

Related

Device / Architecture Coverage

Testing Summary

Testing Checklist

Flags / Guardrails

Performance Impact

Adjacent Tests Considered

Work Tracking

Technical Changes

Known limitations

Uh oh!

therock-pr-bot Bot commented Sep 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 17, 2026

Uh oh!

JonathanLichtnerAMD left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

JonathanLichtnerAMD left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

mattwill-amd commented Oct 30, 2026

Uh oh!

assistant-librarian Bot commented Oct 30, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Nov 3, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

cderb commented Sep 17, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 17, 2026 •

edited

Loading

JonathanLichtnerAMD left a comment •

edited

Loading