Skip to content

[blas][rocblas] Take the HIP stream from the interop handle - #765

Open
zjin-lcf wants to merge 1 commit into
uxlfoundation:developfrom
zjin-lcf:fix/rocblas-interop-stream
Open

[blas][rocblas] Take the HIP stream from the interop handle#765
zjin-lcf wants to merge 1 commit into
uxlfoundation:developfrom
zjin-lcf:fix/rocblas-interop-stream

Conversation

@zjin-lcf

Copy link
Copy Markdown
Contributor

Summary

RocblasScopedContextHandler::get_stream currently asks the queue for a native stream:

return sycl::get_native<sycl::backend::ext_oneapi_hip>(queue);

A SYCL queue owns a pool of HIP streams, so the stream this returns is not necessarily the one the enqueued native command was scheduled on — and only the scheduled stream is covered by the SYCL event that gates the command's completion. Work submitted to a different stream would not be ordered by that event.

The interop handle knows the stream the command was actually scheduled on, so this takes it from there. That matches what the cuBLAS backend already does:

// src/blas/backends/cublas/cublas_scope_handle.cpp
return ih.get_native_queue<sycl::backend::ext_oneapi_cuda>();

and what the AdaptiveCpp build of this same backend does:

// src/blas/backends/rocblas/rocblas_scope_handle_adaptivecpp.cpp
return interop_h.get_native_queue<sycl::backend::hip>();

so this makes the DPC++ path consistent with both.

Caveat

This is a consistency fix found while investigating #486, not a fix for an observed failure. I could not measure any behavioural difference from it on the hardware available (MI210 and MI300A) — the mismatch appears to be latent in practice, presumably because the pool commonly hands back the same stream. I am raising it separately from the #486 fix for that reason.

Test plan

  • Builds against develop with the rocBLAS backend (gfx90a, ROCm 7.1.1)
  • BLAS functional tests pass on an AMD Instinct MI210

Made with Cursor

get_stream asks the queue for a native stream. A SYCL queue owns a pool of HIP
streams, so the returned stream is not necessarily the one the enqueued native
command was scheduled on, and only that stream is covered by the SYCL event
gating the command.

Take the stream from the interop handle instead, which is what the cuBLAS
backend and the AdaptiveCpp build of this backend already do.

Co-authored-by: Cursor <cursoragent@cursor.com>
@zjin-lcf
zjin-lcf requested a review from a team as a code owner August 15, 2026 16:39
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.

1 participant