fix(cuda.core): preserve ordering for asynchronous memory operations in tests

[Andy-Jost]

Summary

Prevent races caused by allocating memory asynchronously on one CUDA stream and consuming it on another stream or outside CUDA stream ordering. This expands the original IPC fix to cover affected memory tests and examples.

Several tests incorrectly assumed that memory allocations against a default stream behaved synchronously. The legacy default stream has special synchronization semantics with respect to blocking streams, but nothing orders it against non-blocking streams. This PR fixes several latent races related to that.

Changes

• Require PatternGen callers to provide a stream and submit fills, copies, and verification on that stream.
• Allocate test buffers on the same stream as their subsequent copies and kernel operations.
• Synchronize before IPC handoff or host access where the consumer cannot inherit CUDA stream dependencies.
• Use CuPy's current stream for the CUDA Core allocation in the strided memory-view example, and wait before NumPy accesses asynchronously allocated managed and pinned memory.

Related Work

Follow-up to #1308.

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[github-actions]

Doc Preview CI
🚀 View preview at https://nvidia.github.io/cuda-python/pr-preview/pr-2657/
https://nvidia.github.io/cuda-python/pr-preview/pr-2657/cuda-core/
https://nvidia.github.io/cuda-python/pr-preview/pr-2657/cuda-bindings/
https://nvidia.github.io/cuda-python/pr-preview/pr-2657/cuda-pathfinder/
Preview will be ready when the GitHub Pages deployment is complete.

[rwgk]

codex gpt-5.6-sol ultra found:

This allocates asynchronously on dev1.default_stream, while PatternGen(dev1, ...) creates a separate nonblocking stream and initializes the buffer there. The later dev1.sync() waits for both streams, but does not establish allocation-before-copy ordering. CUDA defines that cross-stream access as undefined unless explicitly ordered. CUDA documentation (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html).

I’d use one stream throughout:

        stream = dev1.default_stream
        buffer = mr.allocate(NBYTES, stream=stream)
        pgen = PatternGen(dev1, NBYTES, stream=stream)
        pgen.fill_buffer(buffer, seed=False)
        stream.sync()

That also avoids a device-wide barrier.

[Andy-Jost]

Thanks for pointing this out. I did a sweep and found several instances of this happening. The updated PR removes the stream=None option from PatternGen, forcing tests to provide a stream. I also updated several tests to use a single stream rather than multiple.

[rwgk]

LGTM — but please see the one inline suggestion (it was classified as "Low").

codex gpt-5.6-sol ultra:

The original issue is now fixed properly:

• Allocation and PatternGen initialization use the same stream.
• That stream is synchronized before the child process starts.
• All 33 PatternGen call sites now provide an explicit stream, making recurrence less likely.
• Host accesses and CuPy interoperability are likewise ordered correctly.
• A repo-wide sweep found no analogous missed allocation/use race.

The rationale is technically sound. CUDA documents that cross-stream access to a stream-ordered allocation must be explicitly ordered, that nonblocking streams do not implicitly synchronize with the legacy default stream, and that importing an IPC allocation does not wait for the exporter’s allocation to become ready. A host-side stream synchronization before process handoff is a simple, valid test-level barrier. CUDA allocator documentation (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html), default-stream behavior (https://docs.nvidia.com/cuda/cuda-driver-api/stream-sync-behavior.html).

[rwgk]

dev0 is still current here (set at line 134), while Device.sync() calls cudaDeviceSynchronize() without switching to self. As a result, dev1.sync() on the next line synchronizes dev0 rather than the dev1 default stream on which buffer.close() queues the imported buffer's asynchronous free. This is low severity because process teardown/join still protects the exporter lifetime, but the cleanup synchronization itself is ineffective. Could we restore dev1 before closing and synchronizing?

        dev1.set_current()
        buffer.close()
        # TODO(seberg): 2026-06: mr close may be unsafe with incomplete `buf.close()`
        dev1.sync()