Make internal host_allocator stream ordered

[davidwendt]

Description

Changes the cudf::detail::rmm_host_allocator implementation to make it stream-ordered. Uses the feature of cudaMemPoolCreate which allows creating a pinned host memory pool that can be used appropriately with cudaMemcpyAsync to keep the host memory alive in the scope of the stream.

This will allow usages of the various cudf::detail::host_vector to not require a stream synchronize when copying to the device. Several synchronize calls added recently can likely be removed with this change.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[davidwendt]

/ok to test 8ebdb6c

[davidwendt]

Summary

This PR attempted to replace the cudaMallocHost-backed pinned pool with a CUDA native host memory pool
(cudaMemPoolCreate with cudaMemAllocationTypePinned + cudaMemLocationTypeHost, allocating via cudaMallocFromPoolAsync
/ freeing via cudaFreeAsync). The goal was twofold:

1. Remove stream.synchronize() from rmm_host_allocator::allocate() — the CUDA pool guarantees it only recycles blocks
   with no outstanding GPU accesses, so no external sync is needed before CPU writes.
2. Stream-ordered deallocation — cudaFreeAsync defers the reclaim until the stream reaches the free point, meaning
   callers can destroy a host_vector immediately after enqueuing a device copy without needing to sync first.

The implementation compiled and all local tests passed on CUDA 12.8 / sm_86 (Ampere). However, CI exposed three
failure modes, all sharing the same root cause:

Root cause

cudaMallocFromPoolAsync with cudaMemLocationTypeHost does not behave as a drop-in replacement for cudaMallocHost.
cudaMallocHost registers pinned pages through CUDA's unified addressing system, giving them proper pointer attributes
(cudaMemoryTypeHost + a valid devicePointer) that all CUDA memory transfer paths understand. The native host pool
allocates page-locked memory and can be made accessible to the GPU via cudaMemPoolSetAccess, but that grants peer
access, not the same unified-addressing registration. As a result:

• cudaMemcpyAsync (H→D / D→H DMA) fails on some GPU/driver combinations because the DMA engine doesn't find the
  expected pinned registration.
• cudaMemcpyDeviceToDevice fails unconditionally when one pointer is pool host memory, because D→D DMA expects both
  endpoints in device DRAM — the peer access mapping is invisible to that path.

Specific failures

Failure: cudaMemPoolCreate
Environment: CUDA 12.2.2 amd64
Error: cudaErrorInvalidValue — host memory pools require a minimum driver version not present on that runner
────────────────────────────────────────
Failure: cudaMemcpyAsync in IO paths (HybridScanFiltersTest, ParquetTest.ChunkedOperations)
Environment: CUDA 12.9+ amd64
Error: cudaErrorIllegalAddress — DMA path rejects pool host memory on those GPUs
────────────────────────────────────────
Failure: TDigestMergeTest (groupby + reductions)
Environment: Local CUDA 12.8
Error: cudaErrorInvalidValue in thrust::copy D→D — tdigest merge uses cudf::get_pinned_memory_resource() to allocate
device_uvectors in pinned memory and copies into them with cudaMemcpyDeviceToDevice, which rejects pool host pointers

Future path

There is a device attribute — cudaDevAttrHostMemoryPoolsSupported (144, added in CUDA 13.x headers) — that correctly
gates whether this feature is usable. The right implementation would query this attribute at pool construction time
and fall back to the cudaMallocHost-backed pool when it returns 0. Even then, the cudaMemcpyDeviceToDevice
incompatibility would need to be addressed separately, either by using cudaMemcpyDefault everywhere, or by ensuring
pool host memory goes through a different allocation path than the one device_uvector uses for its unified-addressing
trick. A safe minimum would likely be CUDA 13.x + sm_80+ (Ampere) + cudaDevAttrHostMemoryPoolsSupported == 1.

[bdice]

We do need to replace cudaMemcpyDeviceToDevice with cudaMemcpyDefault. Where is that source code? Thrust or cuDF?

[bdice]

For CUDA 12.x we need to use a different set of APIs for pinned memory, I think. There is a bit of nuance here. See rapidsai/rmm#2054.

[bdice]

I opened #23605 to use cudaMemcpyDefault everywhere.

[davidwendt]

/ok to test