Skip to content

Synchronize CUDA async deallocate_sync #2401

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.

Sign up for GitHub

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

Merged
rapids-bot merged 2 commits into from
May 22, 2026
+9 −3
Merged

Synchronize CUDA async deallocate_sync #2401

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ec97f1a
Synchronize cuda async view deallocate_sync
bdice May 17, 2026
9337403
Merge branch 'main' into pr-sync-cuda-async-deallocate
bdice May 21, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion cpp/include/rmm/mr/cuda_async_view_memory_resource.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,7 +133,9 @@ class cuda_async_view_memory_resource final {
std::size_t bytes,
std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept
{
deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
deallocate(stream, ptr, bytes, alignment);
RMM_ASSERT_CUDA_SUCCESS_SAFE_SHUTDOWN(cudaStreamSynchronize(stream.get()));
Comment on lines -136 to +138
Copy link
Copy Markdown
Contributor

@wence- wence- May 22, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

question: what is the point of this synchronisation? It can't really be stream safety because after this call the ptr is not alive and cannot be reused. Is it that you want the freed memory to be immediately available on all streams?

Copy link
Copy Markdown
Collaborator Author

@bdice bdice May 22, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, the name says "sync" so we should guarantee there is no asynchronous work ordered after the end of this function's scope. This should align with the allocation behavior.

}

/**
Expand Down
4 changes: 3 additions & 1 deletion cpp/src/mr/detail/cuda_async_managed_memory_resource_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,9 @@ void cuda_async_managed_memory_resource_impl::deallocate_sync(void* ptr,
std::size_t bytes,
std::size_t alignment) noexcept
{
deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
deallocate(stream, ptr, bytes, alignment);
RMM_ASSERT_CUDA_SUCCESS_SAFE_SHUTDOWN(cudaStreamSynchronize(stream.get()));
}

} // namespace detail
Expand Down
4 changes: 3 additions & 1 deletion cpp/src/mr/detail/cuda_async_memory_resource_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,7 +101,9 @@ void cuda_async_memory_resource_impl::deallocate_sync(void* ptr,
std::size_t bytes,
std::size_t alignment) noexcept
{
deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
deallocate(stream, ptr, bytes, alignment);
RMM_ASSERT_CUDA_SUCCESS_SAFE_SHUTDOWN(cudaStreamSynchronize(stream.get()));
}

} // namespace detail
Expand Down
Loading