NVIDIA / NVIDIA/cudf

[BUG] host_buffer_sink isn't stream-safe

Open
#21,680 1 comment 1 reaction 1 assignee Claimed by @vuule View on GitHub
bug cuIO libcudf
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.