diff --git a/docs/how-to/Programmers_Guide.rst b/docs/how-to/Programmers_Guide.rst index 5355dac11..72fc59fef 100644 --- a/docs/how-to/Programmers_Guide.rst +++ b/docs/how-to/Programmers_Guide.rst @@ -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. diff --git a/library/src/include/handle.hpp b/library/src/include/handle.hpp index bf66fef25..763b4a983 100644 --- a/library/src/include/handle.hpp +++ b/library/src/include/handle.hpp @@ -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; } }