From b00528eb7ef519f8026f3993d1bd5267cf9bd77f Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 15 Aug 2026 09:38:54 -0700 Subject: [PATCH] [blas][rocblas] Take the HIP stream from the interop handle 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 --- src/blas/backends/rocblas/rocblas_scope_handle.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/blas/backends/rocblas/rocblas_scope_handle.cpp b/src/blas/backends/rocblas/rocblas_scope_handle.cpp index 8b59e3f99..ac772b94e 100644 --- a/src/blas/backends/rocblas/rocblas_scope_handle.cpp +++ b/src/blas/backends/rocblas/rocblas_scope_handle.cpp @@ -155,7 +155,10 @@ rocblas_handle RocblasScopedContextHandler::get_handle(const sycl::queue& queue) } hipStream_t RocblasScopedContextHandler::get_stream(const sycl::queue& queue) { - return sycl::get_native(queue); + // The stream must come from the interop handle rather than from the queue: + // a queue owns a pool of streams and only the one the command was scheduled + // on is tracked by the SYCL event that gates this command's completion. + return interop_h.get_native_queue(); } sycl::context RocblasScopedContextHandler::get_context(const sycl::queue& queue) { return queue.get_context();