Forward-merge release/26.10 into main

CONTRIBUTING.md

@@ -32,7 +32,7 @@ conda install cudf -c rapidsai-nightly -c conda-forge
 ```
 
 3. Build and view the docs locally following the instructions in the [Building
-documentation docs](https://docs.rapids.ai/api/cudf/stable/developer_guide/documentation/#building-documentation)
+documentation docs](https://docs.nvidia.com/cudf/latest/cudf/developer_guide/documentation/#building-documentation)
 4. Follow steps 7-10 in the section [Your first issue](#your-first-issue)
 
 ## Code contributions
@@ -325,9 +325,9 @@ This will bring up an interactive prompt to select which spelling fixes to apply
 
 ## Developer Guidelines
 
-The [C++ Developer Guide](cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md) includes details on contributing to libcudf C++ code.
+The [C++ Developer Guide](https://docs.nvidia.com/cudf/latest/libcudf/developer_guide/) includes details on contributing to libcudf C++ code.
 
-The [Python Developer Guide](https://docs.rapids.ai/api/cudf/stable/cudf/developer_guide/) includes details on contributing to cuDF Python code.
+The [Python Developer Guide](https://docs.nvidia.com/cudf/latest/developer_guide/) includes details on contributing to cuDF Python code.
 
 
 ## Attribution

README.md

@@ -1,21 +1,22 @@
-# <div align="left"><img src="img/rapids_logo.png" width="90px"/>&nbsp;cuDF - A GPU-accelerated DataFrame library for tabular data processing</div>
+# NVIDIA cuDF: A GPU-accelerated DataFrame library for tabular data processing
 
-cuDF (pronounced "KOO-dee-eff") is an [Apache 2.0 licensed](LICENSE), GPU-accelerated DataFrame library
-for tabular data processing. The cuDF library is one part of the [RAPIDS](https://rapids.ai/) GPU
-Accelerated Data Science suite of libraries.
+NVIDIA cuDF (pronounced "KOO-dee-eff") is an [Apache 2.0 licensed](LICENSE), GPU-accelerated DataFrame library
+for tabular data processing. The cuDF library is one part of the [NVIDIA
+CUDA-X](https://developer.nvidia.com/cuda/cuda-x-libraries) suite of GPU
+Accelerated libraries.
 
 ## About
 
 cuDF is composed of multiple libraries including:
 
-* [libcudf](https://docs.rapids.ai/api/libcudf/stable/): A CUDA C++ library with [Apache Arrow](https://arrow.apache.org/) compliant
+* [libcudf](https://docs.nvidia.com/cudf/latest/libcudf/): A CUDA C++ library with [Apache Arrow](https://arrow.apache.org/) compliant
 data structures and fundamental algorithms for tabular data.
-* [pylibcudf](https://docs.rapids.ai/api/cudf/stable/pylibcudf/): A Python library providing [Cython](https://cython.org/) bindings for libcudf.
-* [cudf](https://docs.rapids.ai/api/cudf/stable/cudf/): A Python library providing
+* [pylibcudf](https://docs.nvidia.com/cudf/latest/pylibcudf/): A Python library providing [Cython](https://cython.org/) bindings for libcudf.
+* [cudf](https://docs.nvidia.com/cudf/latest/cudf/): A Python library providing
     - A DataFrame library mirroring the [pandas](https://pandas.pydata.org/) API
-    - A zero-code change accelerator, [cudf.pandas](https://docs.rapids.ai/api/cudf/stable/cudf_pandas/), for existing pandas code.
-* [cudf-polars](https://docs.rapids.ai/api/cudf/stable/cudf_polars/): A Python library providing a GPU engine for [Polars](https://pola.rs/)
-* [dask-cudf](https://docs.rapids.ai/api/dask-cudf/stable/): A Python library providing a GPU backend for [Dask](https://www.dask.org/) DataFrames
+    - A zero-code change accelerator, [cudf.pandas](https://docs.nvidia.com/cudf/latest/cudf_pandas/), for existing pandas code.
+* [cudf-polars](https://docs.nvidia.com/cudf/latest/cudf_polars/): A Python library providing a GPU engine for [Polars](https://pola.rs/)
+* [dask-cudf](https://docs.nvidia.com/dask-cudf/latest/): A Python library providing a GPU backend for [Dask](https://www.dask.org/) DataFrames
 
 Notable projects that use cuDF include:
 

conda/recipes/libcudf/recipe.yaml

@@ -209,7 +209,7 @@ outputs:
           - libzlib
           - libnvcomp
     about:
-      homepage: https://rapids.ai/
+      homepage: https://docs.nvidia.com/cudf/
       license: Apache-2.0
       summary: libcudf_kafka library
 
@@ -253,7 +253,7 @@ outputs:
           - libzlib
           - libnvcomp
     about:
-      homepage: https://rapids.ai/
+      homepage: https://docs.nvidia.com/cudf/
       license: Apache-2.0
       summary: libcudf-streaming library
 
@@ -431,6 +431,6 @@ outputs:
           - librmm
           - libnvcomp
     about:
-      homepage: https://rapids.ai/
+      homepage: https://docs.nvidia.com/cudf/
       license: Apache-2.0
       summary: libcudf-streaming test & benchmark executables

cpp/cmake/Modules/ConfigureCUDA.cmake

@@ -1,12 +1,12 @@
 # =============================================================================
 # cmake-format: off
-# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: Apache-2.0
 # cmake-format: on
 # =============================================================================
 
 if(CMAKE_COMPILER_IS_GNUCXX)
-  list(APPEND CUDF_CXX_FLAGS -Wall -Werror -Wno-unknown-pragmas -Wno-error=deprecated-declarations)
+  list(APPEND CUDF_CXX_FLAGS -Wall -Werror -Wno-unknown-pragmas)
 endif()
 
 list(APPEND CUDF_CUDA_FLAGS --expt-extended-lambda)
@@ -17,7 +17,7 @@ if(CUDA_WARNINGS_AS_ERRORS)
 else()
   list(APPEND CUDF_CUDA_FLAGS -Werror=cross-execution-space-call)
 endif()
-list(APPEND CUDF_CUDA_FLAGS -Xcompiler=-Wall,-Werror,-Wno-error=deprecated-declarations)
+list(APPEND CUDF_CUDA_FLAGS -Xcompiler=-Wall,-Werror)
 # This warning needs to be suppressed because some parts of cudf instantiate templated CCCL
 # functions in contexts where the resulting instantiations would have internal linkage (e.g. in
 # anonymous namespaces). In such contexts, the visibility attribute on the template is ignored, and

cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md

@@ -26,7 +26,7 @@ A column is an array of data of a single type. Along with Tables, columns are th
 structures used in libcudf. Most libcudf algorithms operate on columns. Columns may have a validity
 mask representing whether each element is valid or null (invalid). Columns of nested types are
 supported, meaning that a column may have child columns. A column is the C++ equivalent to a cuDF
-Python [Series](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.series/).
+Python [Series](https://docs.nvidia.com/cudf/latest/cudf/api_docs/series/).
 
 ### Element
 
@@ -41,7 +41,7 @@ A type representing a single element of a data type.
 A table is a collection of columns that all have the same number of elements (rows). A table may
 also have zero columns while still carrying a row count, mirroring an `(N, 0)` DataFrame. A table is
 the C++ equivalent to a cuDF Python
-[DataFrame](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.dataframe/).
+[DataFrame](https://docs.nvidia.com/cudf/latest/cudf/api_docs/dataframe/).
 
 ### View
 
@@ -674,7 +674,7 @@ cudf::detail::copy_if(
 
 ## Memory Allocation
 
-Device [memory resources](#rmmdevice_memory_resource) are used in libcudf to abstract and control
+Device [memory resources](#memory-resources) are used in libcudf to abstract and control
 how device memory is allocated.
 
 ### Output Memory

cpp/doxygen/developer_guide/DOCUMENTATION.md

@@ -1,6 +1,6 @@
 # libcudf C++ Documentation Guide
 
-These guidelines apply to documenting all libcudf C++ source files using doxygen style formatting although only public APIs and classes are actually [published](https://docs.rapids.ai/api/libcudf/stable/index.html).
+These guidelines apply to documenting all libcudf C++ source files using doxygen style formatting although only public APIs and classes are actually [published](https://docs.nvidia.com/cudf/latest/libcudf/api_docs/).
 
 ## Copyright License
 
@@ -438,4 +438,5 @@ Then open `<IP address>:8000` in your local web browser, inserting the IP addres
 
 The doxygen output is intended for building documentation only for the public APIs and classes.
 For example, the output should not include documentation for `detail` or `/src` files, and these directories are excluded in the `Doxyfile` configuration.
-When published by the build/CI system, the doxygen output will appear on our external [RAPIDS web site](https://docs.rapids.ai/api/libcudf/stable/index.html).
+When published by the build/CI system, the doxygen output will appear as
+part of the [cuDF documentation](https://docs.nvidia.com/cudf/latest/libcudf/).

cpp/include/cudf/groupby.hpp

@@ -519,6 +519,13 @@ class streaming_groupby {
    * are updated atomically. The input `data` table is not referenced after this
    * call returns.
    *
+   * This function may be called concurrently from multiple host threads on the same object,
+   * and each call may supply a different stream. Callers do not need to serialize the calls or
+   * synchronize between them. Key insertion is serialized internally, on the host and across
+   * streams, because a batch's newly discovered keys are held in a transient encoding that is
+   * only valid while that one insertion is in flight. The aggregation that follows each
+   * insertion updates every group atomically, so those phases overlap freely across streams.
+   *
    * @param data Table containing both key and value columns
    * @param stream CUDA stream used for device memory operations and kernel launches
    *
@@ -532,6 +539,10 @@ class streaming_groupby {
    *
    * Extracts the other object's accumulated intermediate state and merges it into this
    * object's persistent hash table. The other object is not modified.
+   *
+   * This function shares the insertion path with `aggregate()` and is serialized against it, so
+   * it is safe to call while other host threads are calling `aggregate()` on this object. The
+   * source object must not be mutated concurrently.
    * Both objects must have been constructed with compatible aggregation requests,
    * and this object must have had at least one `aggregate()` call.
    *

cpp/src/groupby/streaming_groupby/aggregate.cu

@@ -17,15 +17,14 @@
 #include <thrust/for_each.h>
 
 #include <limits>
+#include <mutex>
 #include <string>
 
 namespace cudf::groupby {
 
 void streaming_groupby::impl::do_aggregate(table_view const& data, cuda::stream_ref stream)
 {
-  CUDF_EXPECTS(!_invalidated,
-               "streaming_groupby is in an invalidated state from a prior failure; "
-               "no further aggregate()/merge() is allowed.  finalize() may still be called.");
+  ensure_not_invalidated();
 
   auto const batch_size = data.num_rows();
   if (batch_size == 0) { return; }
@@ -40,15 +39,29 @@ void streaming_groupby::impl::do_aggregate(table_view const& data, cuda::stream_
                "Transient key encoding (max_distinct_keys + batch_size) would overflow size_type.",
                std::invalid_argument);
 
-  if (!_initialized) { initialize(data, stream); }
+  // The transient key encoding is only valid while a single insertion is in flight, so
+  // insertion is serialized across concurrent callers on the host and, via the event, on the
+  // device.  The aggregation below is per-group atomic and runs unserialized.
+  auto const result = [&] {
+    std::lock_guard const lock{_insert_mutex};
 
-  auto const batch_keys = data.select(_key_indices);
+    // Re-check under the lock: another caller may have invalidated the object since the
+    // fail-fast check above.
+    ensure_not_invalidated();
 
-  update_nullable_state(batch_keys);
+    if (!_initialized) { initialize(data, stream); }
 
-  if (!_key_set) { create_key_set(stream); }
+    auto const batch_keys = data.select(_key_indices);
 
-  auto result = probe_and_insert(batch_keys, stream);
+    update_nullable_state(batch_keys);
+
+    if (!_key_set) { create_key_set(stream); }
+
+    _insert_done.wait(stream);
+    auto inserted = probe_and_insert(batch_keys, stream);
+    _insert_done.record(stream);
+    return inserted;
+  }();
 
   auto const values_view = data.select(_value_col_indices);
   auto const d_values    = table_device_view::create(values_view, stream);

cpp/src/groupby/streaming_groupby/common.cuh

@@ -26,11 +26,40 @@
 #include <cuda/std/utility>
 #include <cuda/stream>
 
+#include <atomic>
 #include <memory>
+#include <mutex>
 #include <vector>
 
 namespace cudf::groupby {
 
+/*
+ * Minimal owning wrapper around a CUDA event, used to order the insertion phase of
+ * `aggregate()` / `merge()` calls that overlap on different streams.  Only the insertion
+ * phase needs this ordering; the aggregation phase updates each group with atomics and is
+ * safe to overlap.
+ */
+class insert_order_event {
+ public:
+  insert_order_event() { CUDF_CUDA_TRY(cudaEventCreateWithFlags(&_event, cudaEventDisableTiming)); }
+  ~insert_order_event() { cudaEventDestroy(_event); }
+  insert_order_event(insert_order_event const&)            = delete;
+  insert_order_event& operator=(insert_order_event const&) = delete;
+
+  /// Makes `stream` wait for the most recently recorded insertion.  No-op before the first
+  /// `record()`, which is exactly the behavior the first call needs.
+  void wait(cuda::stream_ref stream) const
+  {
+    CUDF_CUDA_TRY(cudaStreamWaitEvent(stream.get(), _event));
+  }
+
+  /// Records completion of the insertion just enqueued on `stream`.
+  void record(cuda::stream_ref stream) { CUDF_CUDA_TRY(cudaEventRecord(_event, stream.get())); }
+
+ private:
+  cudaEvent_t _event{};
+};
+
 /*
  * Companion location for a stored dense ID: which compacted batch table the key
  * lives in (`first`) and the row index within that table (`second`).  Packed into
@@ -250,17 +279,27 @@ struct streaming_groupby::impl {
   null_policy _null_handling;
   cuda::mr::any_resource<cuda::mr::device_accessible> _mr;
 
+  /*
+   * Serializes the insertion phase of `aggregate()` and `merge()`.  Callers may invoke those
+   * from multiple host threads; everything they mutate on the host, and the transient key
+   * encoding they place in the hash set, is guarded here.
+   */
+  std::mutex _insert_mutex;
+  /// Orders the insertion phase across calls that supply different streams.
+  insert_order_event _insert_done;
+
   bool _initialized{false};
   /// Set true once an `aggregate()` / `merge()` call has thrown after touching the
   /// hash set.  Subsequent `aggregate()` / `merge()` calls fail fast; only
-  /// `finalize()` may still be called to recover partial results.
-  bool _invalidated{false};
+  /// `finalize()` may still be called to recover partial results.  Atomic so the
+  /// fail-fast check in `do_aggregate` can run ahead of `_insert_mutex`.
+  std::atomic<bool> _invalidated{false};
   /*
    * Number of distinct keys accumulated so far.  Also serves as the high-water
    * mark of dense IDs in the persistent hash set: stored slot values are in
    * [0, _distinct_keys).
    */
-  size_type _distinct_keys{0};
+  std::atomic<size_type> _distinct_keys{0};
   bool _has_nullable_keys{false};
   bool _has_nested_keys{false};
 
@@ -301,7 +340,12 @@ struct streaming_groupby::impl {
   std::unique_ptr<streaming_set_t> _key_set;
 
   [[nodiscard]] size_type num_keys() const { return static_cast<size_type>(_key_indices.size()); }
-  [[nodiscard]] bool has_state() const { return _initialized && _distinct_keys > 0; }
+  void ensure_not_invalidated() const
+  {
+    CUDF_EXPECTS(!_invalidated.load(std::memory_order_relaxed),
+                 "streaming_groupby is in an invalidated state from a prior failure; "
+                 "no further aggregate()/merge() is allowed.  finalize() may still be called.");
+  }
 
   impl(host_span<size_type const> key_indices,
        host_span<streaming_aggregation_request const> requests,

cpp/src/groupby/streaming_groupby/impl.cu

@@ -243,7 +243,9 @@ std::unique_ptr<table> streaming_groupby::impl::gather_agg_results(
   // The results we care about are dense in `[0, _distinct_keys)` and can be extracted by
   // slice+copy.
   auto const sliced =
-    cudf::detail::slice(_agg_results->view(), {0, _distinct_keys}, stream).front();
+    cudf::detail::slice(
+      _agg_results->view(), {0, _distinct_keys.load(std::memory_order_relaxed)}, stream)
+      .front();
   return std::make_unique<table>(sliced, stream, mr);
 }
 
@@ -378,7 +380,10 @@ std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> streaming_gro
   return _impl->do_finalize(stream, mr);
 }
 
-size_type streaming_groupby::distinct_keys() const noexcept { return _impl->_distinct_keys; }
+size_type streaming_groupby::distinct_keys() const noexcept
+{
+  return _impl->_distinct_keys.load(std::memory_order_relaxed);
+}
 
 bool is_streaming_groupby_supported(data_type values_type, aggregation::Kind kind)
 {

cpp/src/groupby/streaming_groupby/insert.cuh

@@ -96,9 +96,10 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_
     // Bound check: the hash set has already been written above (transient slot values),
     // so on failure the object is left invalidated; further aggregate()/merge() calls
     // will throw immediately while finalize() can still recover partial results.
-    if (_distinct_keys + new_distinct_keys > _max_distinct_keys) {
+    auto const distinct_so_far = _distinct_keys.load(std::memory_order_relaxed);
+    if (distinct_so_far + new_distinct_keys > _max_distinct_keys) {
       _invalidated = true;
-      CUDF_FAIL("Distinct key count (" + std::to_string(_distinct_keys + new_distinct_keys) +
+      CUDF_FAIL("Distinct key count (" + std::to_string(distinct_so_far + new_distinct_keys) +
                 ") would exceed max_distinct_keys (" + std::to_string(_max_distinct_keys) + ").");
     }
 
@@ -115,7 +116,7 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_
 
     // Store the compacted batch.
     auto const new_batch_id    = static_cast<size_type>(_compacted_batches.size());
-    auto const dense_id_offset = _distinct_keys;
+    auto const dense_id_offset = distinct_so_far;
     _compacted_batches.push_back(std::move(compacted));
     _preprocessed_batches.push_back(preprocessed_compacted);
 
@@ -138,7 +139,7 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_
                        update_transient_target_indices_fn{
                          base, slot_offsets.data(), _max_distinct_keys, target_indices.data()});
 
-    _distinct_keys += new_distinct_keys;
+    _distinct_keys.fetch_add(new_distinct_keys, std::memory_order_relaxed);
   }
   // If new_distinct_keys == 0, target_indices is already final from Pass 1 — every
   // slot held a dense ID at probe time, so *iter was already the correct dense ID.