Migrate stream APIs from rmm::cuda_stream_view to cuda::stream_ref - #23929
Conversation
|
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. |
84de667 to
bb334ee
Compare
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughThe pull request updates CUDA stream usage across native code, Python streaming bindings, and streaming join tests. It also updates the SPDX attribution in ChangesCUDA stream migration
Priority: ⬇️ Low Estimated code review effort: 1 (Trivial) | ~5 minutes Change: Refactor Merge Risk: ⚪ Minimal · up to The stream-reference migration preserves the inspected native and Python API contracts, with no concrete build or runtime regression established. 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
Full details: Docstring CoverageExplanation Docstring coverage is 25.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff. Analyzed 4 functions across 15 files. (1 skipped: 1 unsupported.) ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
a61eee2 to
56b7d90
Compare
|
/merge |
Signed-off-by: Bradley Dice <bdice@bradleydice.com>
Signed-off-by: Bradley Dice <bdice@bradleydice.com>
Notes on CUDA contextsWhile switching everything from I already have a Python approval from earlier from @vyasr and I am going to merge this PR as-is because it is blocking downstream repos until the migration is complete. However, there are some follow-up discussions we need to have. The current implementation of
|
| if isinstance(stream, runtime.cudaStream_t): | ||
| return Stream._from_cudaStream_t( | ||
| <cudaStream_t><uintptr_t>int(stream), owner=stream | ||
| ) |
There was a problem hiding this comment.
This makes _get_stream accept runtime.cudaStream_t objects directly, because older versions of cuda-bindings that we support do not have __cuda_stream__ support in cuda.bindings.runtime.cudaStream_t.
Previously we used Cython wrappers around rmm::cuda_stream_view which has support for __cuda_stream__ and can be implicitly converted to cudaStream_t.
|
/merge |
9d3e855
into
NVIDIA:release/26.10
Follow-up to the [CUDA context discussion in #23929](#23929 (comment)). Use `cuCtxGetCurrent()` from `cuda.bindings.cydriver` in `_ensure_cuda_context()` to skip the CUDA Runtime call when a context is already current. This also avoids releasing and reacquiring the GIL on the warm path. The existing `cudaFree(NULL)` initialization, with the GIL released, remains for fresh threads or an uninitialized CUDA driver. Other driver errors are propagated. This requires us to add `cuda-profiler-api` to pylibcudf's conda `host` requirements. CUDA 12 `cydriver` declarations include `cudaProfiler.h` even when importing only context APIs. This is only a build-time dependency; no profiling functionality is used. A local microbenchmark of `_get_stream(None)` decreased from **166–220 ns** to **43–50 ns** per call. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) URL: #24151
Summary
Track the coordinated migration of stream APIs and call sites from
rmm::cuda_stream_viewto CCCL'scuda::stream_ref. This propagatescuda::stream_refthrough RMM containers and memory resources, RAFT resource and handle APIs, downstream C++ interfaces, Python/Cython bindings, benchmarks, tests, and documentation.This updates cuDF C++ and Cython stream accessors, building on the default-stream work merged in #23770. It continues work from #23636.
Depends on rapidsai/rmm#2372.
Tracked in rapidsai/build-planning#318.
Migrations
cuda::stream_refthrough stream pools, resource accessors, conditionals, and downstream APIs without converting tormm::cuda_stream_viewcuda::stream_refconstructions for default/legacy/per-thread streamsrmm::cuda_stream_default➡️cuda::stream_ref{cudaStream_t{cudaStreamDefault}}rmm::cuda_stream_legacy➡️cuda::stream_ref{cudaStreamLegacy}rmm::cuda_stream_per_thread➡️cuda::stream_ref{cudaStreamPerThread}.get()when calling an API that requires a rawcudaStream_t, including CUDA runtime, library, CUB, and legacy API boundaries (previouslyrmm::cuda_stream_viewusedvalue()).sync()when synchronizing acuda::stream_ref(previouslyrmm::cuda_stream_viewusedsynchronize())CUDA context initialization
Includes the context fix from #24147 so it lands together with the stream migration.
cuda::stream_ref::sync()calls the CUDA Driver API, which does not implicitly establish a current context on fresh execution threads. Empty-column/metadata operations can therefore fail withCUDA_ERROR_INVALID_CONTEXTin Dask-cuDF and cuStreamz.pylibcudf.utils._get_stream()now calls_ensure_cuda_context(), which usescuda.bindings.cyruntime.cudaFree(NULL)to establish a current context when needed. RMM is unchanged, initialization does not occur at import time, and no profiler-header dependency is needed. RuntimecudaStream_tobjects are accepted directly, including older bindings versions without__cuda_stream__support.The runtime-only
_get_stream()measured 126 ns per warm call, 7.11 µs for first use on a fresh thread, and 192 ms for first-process CUDA initialization. It returned while GPU work remained pending in all four tested stream modes; no forced GPU synchronization was observed. The current implementation releases the GIL on every call, which can introduce scheduling delays under Python-thread contention. See updated benchmark methodology, comparisons, and validation limitations.Checklist