Skip to content

[release/2.12] reuse one stream for warmup and captures in make_graphed_callables#3354

Merged
pragupta merged 1 commit into
ROCm:release/2.12from
dnikolaev-amd:dnikolaev/use_one_stream_make_graphed_callables_rel2.12
Jun 22, 2026
Merged

[release/2.12] reuse one stream for warmup and captures in make_graphed_callables#3354
pragupta merged 1 commit into
ROCm:release/2.12from
dnikolaev-amd:dnikolaev/use_one_stream_make_graphed_callables_rel2.12

Conversation

@dnikolaev-amd

Copy link
Copy Markdown

Summary

This PR updates torch.cuda.make_graphed_callables so warmup and graph capture use the same stream, avoiding ROCm graph capture failures tied to hipBLASLt.

Why

On ROCm, hipBLASLt handle caching is per (device, stream). If stream= is omitted, torch.cuda.graph uses its own default capture stream, which is not the stream used for warmup. The first real use on that capture stream will trigger a lazy hipblasLtCreate there. Handle creation runs capture‑unsafe internal allocation/setup, which fail with stream-capture errors (and sometimes hang) during capture.

Running warmup on the same stream passed into torch.cuda.graph(..., stream=...) creates and caches the hipBLASLt handle before capture, so hipblasLtCreate (and its associated allocations) are not hit mid‑capture.

Fixes pytorch#179961

Pull Request resolved: pytorch#187745
Approved by: https://github.com/jeffdaily

(cherry picked from commit d9d756c)

…bles (pytorch#187745)

### Summary

This PR updates `torch.cuda.make_graphed_callables` so warmup and graph capture use the same stream, avoiding ROCm graph capture failures tied to hipBLASLt.

### Why

On ROCm, hipBLASLt handle caching is per (device, stream). If `stream=` is omitted, `torch.cuda.graph` uses its own default capture stream, which is not the stream used for warmup. The first real use on that capture stream will trigger a lazy hipblasLtCreate there. Handle creation runs capture‑unsafe internal allocation/setup, which fail with stream-capture errors (and sometimes hang) during capture.

Running warmup on the same stream passed into `torch.cuda.graph(..., stream=...)` creates and caches the hipBLASLt handle before capture, so hipblasLtCreate (and its associated allocations) are not hit mid‑capture.

Fixes pytorch#179961

Pull Request resolved: pytorch#187745
Approved by: https://github.com/jeffdaily

(cherry picked from commit d9d756c)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants