Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

PR #16520: [ROCM] ResetStream function for GemmAlgorithmPicker (BlasSupport interface) #17712

Open
wants to merge 1 commit into
base: main
Choose a base branch
from

Conversation

copybara-service[bot]
Copy link

@copybara-service copybara-service bot commented Sep 27, 2024

PR #16520: [ROCM] ResetStream function for GemmAlgorithmPicker (BlasSupport interface)

Imported from GitHub PR #16520

Here I added ResetStream function which sets the underlying stream for cublas/rocblas libraries to default stream 0.

This is useful for GemmAlgorithmPicker which uses a temporary stream object for autotuning. In rocblas, rocblas_set_stream function is persistent, meaning that once the stream value is set, it will be used in all subsequent computations until new stream value is set.

In case of GemmAlgorithmPicker, we leave a destroyed stream object set into the math library. This does not produce any error behaviour but merely just a warning on ROCM side: "Stream Capture Check Failed".

With this new ResetStream function, one can reset the stream value in GemmAlgorithmPicker destructor. Potentially, it can also be useful in other places where temporary stream value is used.

Besides, I have also made some small code restructure for GemmAlgorithmPicker

@xla-rotation: could you have a look please?

Copybara import of the project:

--
2bd0cf2 by Pavel Emeliyanenko [email protected]:

set stream to null at the end of rocm_blas gemm function call

--
436d073 by Pavel Emeliyanenko [email protected]:

fixing buildbreaks

--
9347c71 by Pavel Emeliyanenko [email protected]:

added test for reset_stream

--
bb009b0 by Pavel Emeliyanenko [email protected]:

changed IsMainStreamSet interface

Merging this change closes #16520

FUTURE_COPYBARA_INTEGRATE_REVIEW=#16520 from ROCm:ci_blas_reset_stream bb009b0

@copybara-service copybara-service bot force-pushed the test_679049775 branch 13 times, most recently from 94a18af to dbbad03 Compare September 30, 2024 19:34
…upport interface)

Imported from GitHub PR #16520

Here I added **ResetStream** function which sets the underlying stream for cublas/rocblas libraries to default stream 0.

This is useful for GemmAlgorithmPicker which uses a temporary stream object for autotuning. In rocblas, **rocblas_set_stream** function is **persistent**, meaning that once the stream value is set, it will be used in all subsequent computations until new stream value is set.

In case of GemmAlgorithmPicker, we leave a **destroyed** stream object set into the math library. This does not produce any error behaviour but merely just a warning on ROCM side: "Stream Capture Check Failed".

With this new ResetStream function, one can reset the stream value in GemmAlgorithmPicker destructor. Potentially, it can also be useful in other places where temporary stream value is used.

Besides, I have also made some small code restructure for GemmAlgorithmPicker

@xla-rotation: could you have a look please?

Copybara import of the project:

--
2bd0cf2 by Pavel Emeliyanenko <[email protected]>:

set stream to null at the end of rocm_blas gemm function call

--
436d073 by Pavel Emeliyanenko <[email protected]>:

fixing buildbreaks

--
9347c71 by Pavel Emeliyanenko <[email protected]>:

added test for reset_stream

--
bb009b0 by Pavel Emeliyanenko <[email protected]>:

changed IsMainStreamSet interface

Merging this change closes #16520

FUTURE_COPYBARA_INTEGRATE_REVIEW=#16520 from ROCm:ci_blas_reset_stream bb009b0
PiperOrigin-RevId: 679049775
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