Skip to content

Commit

Permalink
Resolve error/warn message during rocblas_set_stream() call (#2664)
Browse files Browse the repository at this point in the history
  • Loading branch information
rkamd authored and JeniferC99 committed Aug 20, 2024
1 parent 01e00eb commit 30c8920
Show file tree
Hide file tree
Showing 2 changed files with 17 additions and 7 deletions.
17 changes: 11 additions & 6 deletions docs/how-to/Programmers_Guide.rst
Original file line number Diff line number Diff line change
Expand Up @@ -183,31 +183,36 @@ For the library to use a non-default device within a host thread, the device mus

The device in the host thread should not be changed between ``hipStreamCreate`` and ``hipStreamDestroy``. If the device in the host thread is changed between creating and destroying the stream, then the behavior is undefined.

If the user created a non-default stream, it is the user's responsibility to synchronize the non-default stream before destroying it:
If the user created a non-default stream, it is the user's responsibility to synchronize the old non-default stream, and update rocblas handle with default/new non-default stream before destroying the old non-default stream.

::

// Synchronize the non-default stream before destroying it
if(hipStreamSynchronize(stream) != hipSuccess) return EXIT_FAILURE;

// Reset the stream reference in the handle to either default or new non-default
if(rocblas_set_stream(handle, 0) != rocblas_status_success) return EXIT_FAILURE;

if(hipStreamDestroy(stream) != hipSuccess) return EXIT_FAILURE;

When a user changes the stream from one non-default stream to another non-default stream, it is the user's responsibility to synchronize the old stream before setting the new stream. Then, the user can optionally destroy the old stream:
.. note::
Resetting the rocblas handle's stream reference is essential to avoid the internally handled `hipErrorContextIsDestroyed` error. If this step is skipped, users may encounter this error in AMD_LOG_LEVEL logging or with hipPeekAtLastError( ).

When a user switches from one non-default stream to another, they must complete all rocblas operations previously submitted with this handle on the old stream using ``hipStreamSynchronize(old_stream)`` API before setting the new stream.
::

// Synchronize the old stream
if(hipStreamSynchronize(old_stream) != hipSuccess) return EXIT_FAILURE;

// Destroy the old stream (this step is optional but must come after synchronization)
if(hipStreamDestroy(old_stream) != hipSuccess) return EXIT_FAILURE;

// Create a new stream (this step can be done before the steps above)
if(hipStreamCreate(&new_stream) != hipSuccess) return EXIT_FAILURE;

// Set the handle to use the new stream (must come after synchronization)
// Set the handle to use the new stream (must come after synchronization & before deletion of old stream)
if(rocblas_set_stream(handle, new_stream) != rocblas_status_success) return EXIT_FAILURE;

// Destroy the old stream (this step is optional but must come after synchronization)
if(hipStreamDestroy(old_stream) != hipSuccess) return EXIT_FAILURE;

The above ``hipStreamSynchronize`` is necessary because the ``rocBLAS_handle`` contains allocated device
memory that must not be shared by multiple asynchronous streams at the same time.

Expand Down
7 changes: 6 additions & 1 deletion library/src/include/handle.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -369,7 +369,12 @@ struct _rocblas_handle
hipError_t status = hipStreamIsCapturing(stream, &capture_status);
if(status != hipSuccess)
{
PRINT_IF_HIP_ERROR(status);
//Avoid warning users about hipErrorContextIsDestroyed error on a destroyed stream,
//as it may not be in capture mode, and such notifications would be unnecessary.
if(status != hipErrorContextIsDestroyed)
{
PRINT_IF_HIP_ERROR(status);
}
return false;
}
}
Expand Down

0 comments on commit 30c8920

Please sign in to comment.