[BUG] host_buffer_sink isn't stream-safe
- Dominant language
- C++
- Stars
- 9.8k
- Forks
- 1.1k
- Avg merge
- 3d 6m
- Merged PRs (30d)
- 278
Description
**Describe the bug**
While working on #20800, I found an issue with stream ordering in `host_buffer_sink::device_write_async`. Currently `cudaMemcpyAsync` is doing a synchronization (it's not actually async in this case) and we are reliant on that synchronization for stream-ordering correctness. I tried replacing it with `cudf::detail::memcpy_async` which uses `cudaMemcpyBatchAsync` internally, but that caused tests to fail (segfault? can't recall).
The implementation looks like:
```cpp
std::future device_write_async(void const* gpu_data,
size_t size,
rmm::cuda_stream_view stream) override
{
auto const current_size = buffer_->size();
buffer_->resize(current_size + size);
// TODO: Replace with memcpy_batch_async after fixing stream ordering
// The issue: buffer_->resize() can reallocate, invalidating pointers from previous async copies
// that are still in-flight when using cudaMemcpySrcAccessOrderStream. Need to ensure stream
// ordering or pre-reserve buffer to avoid reallocation.
CUDF_CUDA_TRY(cudaMemcpyAsync(
buffer_->data() + current_size, gpu_data, size, cudaMemcpyDeviceToHost, stream.value()));
return std::async(std::launch::deferred, [stream]() -> void { stream.synchronize(); });
}
```
See https://github.com/rapidsai/cudf/pull/20800/changes#r2673497421 for more discussion.
We may need to track what streams were most recently used and insert some syncs along the way, but I haven't fully diagnosed it.
Contributor guide
Assessment
This issue has not been assessed yet.