Fix temporary host source lifetime with a separate copy API by PointKernel · Pull Request #23561 · NVIDIA/cudf

PointKernel · 2026-08-05T22:08:43Z

Summary

add separate *_async_consume_source APIs for host-to-device copies whose CPU source may disappear when the function returns
keep the existing async APIs and their existing rule: callers must keep the source alive until the stream reaches the copy
use cudaMemcpyBatchAsync with cudaMemcpySrcAccessOrderDuringApiCall on CUDA 13+; use a stream synchronization on older CUDA/default-stream paths to provide the same source-lifetime promise
move temporary-source callers to the new API, including Arrow import, groupby, merge, ORC, Parquet, row operators, strings, and the hybrid dictionary/page-index paths
teach the stream checker to validate cudaMemcpyBatchAsync and its per-thread-default-stream entry point on CUDA 13+
keep the GH200/GB300 CI setup from Fix GB300 and GH200 #23517

Root cause

This is a cuDF caller lifetime bug, not a bug in the normal async-copy API.

Some functions gave CUDA a pointer to a local CPU vector, queued an asynchronous copy, and then destroyed the vector when the function returned. CUDA is allowed to read that source later. GB300's timing exposed the race: the GPU sometimes read memory after the CPU vector was gone.

The first real hybrid-scan failure was HybridScanFiltersTest.FilterRowGroupsWithDictionary. Its temporary list of dictionary-result pointers had this lifetime problem. The later hybrid-scan failures happened after the first illegal access had already damaged the CUDA context.

The new API gives callers an explicit choice:

normal async API: the caller keeps the source alive
consume-source async API: CUDA finishes reading the host source before the function returns; destination work may still finish later

The later STREAM_MERGE_TEST wrong-stream exception was emitted by the test memory-resource guard, not the CUDA API callback checker. The same GB300 job passed all 120 C++ test binaries when rerun without changing the commit, and the exact merge case did not fail in 4,000 local stress runs. That exception is therefore intermittent and is not evidence that the new copy API selected the wrong stream. During the audit, however, merge was found to have two real host-source lifetime gaps; both now use the consume-source API.

Validation

Built and installed latest main plus this patch in the CUDA 13.3 devcontainer for one local GH200-compatible architecture (sm90), with sccache enabled.

complete C++ CTest set with six-way scheduling: 120/120 test binaries passed
full STREAM_MERGE_TEST: 49/49 passed; exact MergeTest.KeysWithNulls passed 100 additional post-fix repetitions
pre-fix investigation: exact merge case passed 1,000 sequential and 3,000 six-process contention repetitions locally
full HYBRID_SCAN_TEST: 85/85 passed
exact original hybrid CI failure under CUDA memcheck: passed, 0 errors
source-lifetime regression test: passed
full pylibcudf suite: 11,390 passed, 1,148 skipped
all pre-commit hooks passed on all 21 PR files after merging latest main
unchanged GB300 rerun of the prior head: all 120 C++ test binaries passed

Related: #23498 and #23517.

copy-pr-bot · 2026-08-05T22:08:46Z

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.

PointKernel · 2026-08-05T22:09:42Z

/ok to test

PointKernel · 2026-08-05T22:18:39Z

/ok to test 210be0d

wence- · 2026-08-05T22:22:24Z

+                                  .flags          = cudaMemcpyFlagPreferOverlapWithCompute};
+    CUDF_CUDA_TRY(
+      cudaMemcpyBatchAsync(dsts, srcs, sizes, 1, &attrs, &attrs_idxs, 1, stream.value()));
+    return;


needs a sync before return in this case?

ah no because of the memcpy attributes?

I would prefer a solution with explicit syncs.

I also don't like that we're doing an explicit call inside an async API. Let me rework this.

…memcpy-during-api-call

PointKernel · 2026-08-06T00:12:41Z

+  std::size_t const* sizes,
+  std::size_t count,
+  rmm::cuda_stream_view stream,
+  host_source_access_order source_access_order = host_source_access_order::STREAM);


testing only, not the final solution

PointKernel · 2026-08-06T00:14:04Z

/ok to test eeadee3

PointKernel · 2026-08-06T02:07:01Z

/ok to test

…memcpy-during-api-call

PointKernel · 2026-08-06T06:59:03Z

/ok to test df79dac

KyleFromNVIDIA · 2026-08-06T13:13:57Z

/ok to test fb5ea3f

PointKernel · 2026-08-06T18:09:41Z

/ok to test

Adds the *_async_consume_source copy APIs (cudaMemcpySrcAccessOrderDuringApiCall on CUDA 13+) and converts the call sites that pass short-lived host buffers. Deferred pageable reads on coherent-memory systems otherwise become use-after-free once the caller frees the source (observed on GB10 as intermittent sort-order corruption under runtime PTDS).

…urces Covers the call sites NVIDIA#23561 does not convert (row operators, ORC/Parquet writers, contiguous_split, row_bit_count and tests) with explicit synchronization before the host buffers go out of scope. Where the two backports overlap, the NVIDIA#23561 form is kept.

PointKernel · 2026-08-19T16:33:53Z

Closing this as #23517 is the right solution to go.

## Description Fix asynchronous host-to-device copies whose host source could be destroyed or mutated before the copy completed. The CUDA 13 `cudaMemcpyBatchAsync` changes in rapidsai/rmm#2511 exposed these invalid lifetime assumptions as nondeterministic failures in pylibcudf, cudf-polars, and hybrid scan tests. Synchronize affected copies at the ownership boundary, and preserve backing storage for Python buffer slices until queued copies can consume them. Also recognize `cudaMemcpyBatchAsync` in the stream-usage checker. This does not introduce a new API or change source ownership semantics. This borrows some lifetime fixes from #23517 and #23561 that we observed were necessary on GB300 but haven't been merged upstream yet. ## Checklist - [x] I am familiar with the [Contributing Guidelines](https://github.com/NVIDIA/cudf/blob/HEAD/CONTRIBUTING.md). - [x] New or existing tests cover these changes. - [x] The documentation is up to date with these changes. --------- Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>

The upstream-selected resolution of the deferred pageable host-to-device copy hazards (supersedes the in-review NVIDIA#23561/NVIDIA#23517 forms shipped in build 4): synchronize affected copies at the ownership boundary instead of adding a separate copy API. Adapted to 26.08 types (rmm::cuda_stream_view / synchronize()); the Python-side buffer preservation is out of scope for the libcudf package.

The upstream-selected resolution of the deferred pageable host-to-device copy hazards (supersedes the in-review NVIDIA#23561/NVIDIA#23517 forms shipped in build 4): synchronize affected copies at the ownership boundary instead of adding a separate copy API. Adapted to 26.08 types (rmm::cuda_stream_view::synchronize()); the Python-side buffer preservation is out of scope for the libcudf package.

The upstream-selected resolution of the deferred pageable host-to-device copy hazards (supersedes the in-review NVIDIA#23561/NVIDIA#23517 forms shipped in build 4): synchronize affected copies at the ownership boundary instead of adding a separate copy API. Adapted to 26.08 types (rmm::cuda_stream_view::synchronize()); the Python-side buffer preservation is out of scope for the libcudf package. (cherry picked from commit 3d94601)

Fix pageable source lifetime in cuda_memcpy_async

c4fc8f1

github-actions Bot assigned PointKernel Aug 5, 2026

github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Aug 5, 2026

PointKernel added bug Something isn't working non-breaking Non-breaking change labels Aug 5, 2026

Update copyright notice

210be0d

wence- reviewed Aug 5, 2026

View reviewed changes

PointKernel added 2 commits August 6, 2026 00:09

Fix host source lifetime in async copies

43c6858

Merge remote-tracking branch 'upstream/main' into agent/fix-pageable-…

eeadee3

…memcpy-during-api-call

PointKernel commented Aug 6, 2026

View reviewed changes

Add consume-source async copy APIs

31ef341

PointKernel changed the title ~~Fix pageable source lifetime in cuda_memcpy_async~~ Fix temporary host source lifetime with a separate copy API Aug 6, 2026

PointKernel added 2 commits August 6, 2026 06:25

Fix merge host source lifetime

599d2b8

Merge remote-tracking branch 'upstream/main' into agent/fix-pageable-…

df79dac

…memcpy-during-api-call

KyleFromNVIDIA added 2 commits August 6, 2026 09:13

Fix test matrices for GB300

c855da0

Merge branch 'main' into agent/fix-pageable-memcpy-during-api-call

fb5ea3f

github-actions Bot assigned KyleFromNVIDIA Aug 6, 2026

Make host-to-device async copies source-safe

7330580

a-hirota mentioned this pull request Aug 13, 2026

Select few large segments in segmented_top_k with cub::DeviceTopK ins… #23602

Open

3 tasks

bdice mentioned this pull request Aug 19, 2026

Fix host-source lifetime races #23718

Merged

3 tasks

PointKernel closed this Aug 19, 2026

PointKernel deleted the agent/fix-pageable-memcpy-during-api-call branch August 19, 2026 16:34

Uh oh!

Conversation

PointKernel commented Aug 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Root cause

Validation

Uh oh!

copy-pr-bot Bot commented Aug 5, 2026

Uh oh!

PointKernel commented Aug 5, 2026

Uh oh!

PointKernel commented Aug 5, 2026

Uh oh!

wence- Aug 5, 2026

Choose a reason for hiding this comment

Uh oh!

wence- Aug 5, 2026

Choose a reason for hiding this comment

Uh oh!

bdice Aug 5, 2026

Choose a reason for hiding this comment

Uh oh!

PointKernel Aug 5, 2026

Choose a reason for hiding this comment

Uh oh!

PointKernel Aug 6, 2026

Choose a reason for hiding this comment

Uh oh!

PointKernel commented Aug 6, 2026

Uh oh!

PointKernel commented Aug 6, 2026

Uh oh!

PointKernel commented Aug 6, 2026

Uh oh!

KyleFromNVIDIA commented Aug 6, 2026

Uh oh!

PointKernel commented Aug 6, 2026

Uh oh!

PointKernel commented Aug 19, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

PointKernel commented Aug 5, 2026 •

edited

Loading