Fix race condition in test_cudart_cudaMemcpy3DPeerAsync

[rwgk]

Description

test_cudart_cudaMemcpy3DPeerAsync was experiencing flaky failures, particularly on Windows when per-thread default stream (PTDS) mode is enabled via CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM=1.

Root Cause

The test performs a synchronous cudaMemcpy from pageable host memory to device memory, followed immediately by cudaMemcpy3DPeerAsync in an explicit stream. When using pageable host memory, cudaMemcpy may return before the DMA transfer to device memory has completed. The subsequent async copy operation in a different stream can then read from device memory before the host-to-device transfer finishes, resulting in incorrect (zero) data being copied.

This race condition is masked in legacy default stream mode due to implicit synchronization between blocking streams and the default stream, but is exposed in PTDS mode where streams do not implicitly synchronize.

Test Order Dependency

The flakiness was exposed through test order randomization. The cuda_bindings test suite uses pytest-randomly (see #1268), which automatically randomizes test execution order. The test suite shares a module-scoped CUDA context fixture across all tests in test_cudart.py.

When certain tests (notably test_cudaGraphConditionalHandleCreate_v2, test_cudart_cudaGetTextureObjectTextureDesc, test_cudart_cudaMemcpy2DToArray_DtoD, and test_cudart_cudaGraphAddMemcpyNode1D) execute before test_cudart_cudaMemcpy3DPeerAsync, they leave the CUDA execution context in a state that makes the race condition more likely to manifest. Removing any one of these tests from the execution sequence before test_cudart_cudaMemcpy3DPeerAsync causes the test to pass, demonstrating the order dependency.

Solution

Add cudaStreamSynchronize(0) after cudaMemcpy to explicitly ensure the DMA transfer to device memory completes before launching the async copy operation. This establishes proper ordering between the default stream and the explicit stream, making the test reliable under both legacy and PTDS stream semantics.

Comprehensive flakiness testing (win-64 WDDM)

• With fix: 1000 trials - ✓ All tests passed with no flakes, errors, or crashes
• Without fix: 1000 trials - 58 failures of test_cudart_cudaMemcpy3DPeerAsync

Each trial was running these commands:

cd cuda_bindings\
python -m pytest -ra -s -vv "tests\test_cudart.py"
set CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM=1 && python -m pytest -ra -s -vv "tests\test_cudart.py" & set CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM=

Internal-only xref with full logs: NVIDIA/cuda-python-private#239

[copy-pr-bot]

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

Contributors can view more details about this message here.

[rwgk]

/ok to test

[github-actions]

Successfully created backport PR for 12.9.x:

• [Backport 12.9.x] Fix race condition in test_cudart_cudaMemcpy3DPeerAsync #1554

[github-actions]

Doc Preview CI
Preview removed because the pull request was closed or merged.

[github-actions]

Backport failed for 12.9.x, because it was unable to cherry-pick the commit(s).

Please cherry-pick the changes locally and resolve any conflicts.

git fetch origin 12.9.x
git worktree add -d .worktree/backport-1550-to-12.9.x origin/12.9.x
cd .worktree/backport-1550-to-12.9.x
git switch --create backport-1550-to-12.9.x
git cherry-pick -x c94557c082c93bb7dfed7ce5f509cf8cf5a9deec

[github-actions]

Backport failed for 12.9.x, because it was unable to cherry-pick the commit(s).

Please cherry-pick the changes locally and resolve any conflicts.

git fetch origin 12.9.x
git worktree add -d .worktree/backport-1550-to-12.9.x origin/12.9.x
cd .worktree/backport-1550-to-12.9.x
git switch --create backport-1550-to-12.9.x
git cherry-pick -x c94557c082c93bb7dfed7ce5f509cf8cf5a9deec

[rwgk]

@leofang I'm mainly just curious / want to confirm: Can we ignore the "Backport failed" messages? Does editing the PR labels trigger them somehow?