refactor: replace rmm::device_scalar with cudf::detail::device_scalar

cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md

@@ -749,37 +749,30 @@ custom_memory_resource *mr...;
 rmm::device_buffer custom_buff(100, mr, stream);
 ```
 
-#### rmm::device_scalar<T>
-Allocates a single element of the specified type initialized to the specified value. Use this for
-scalar input/outputs into device kernels, e.g., reduction results, null count, etc. This is
-effectively a convenience wrapper around a `rmm::device_vector<T>` of length 1.
+#### cudf::detail::device_scalar<T>
+A self-contained device scalar that owns a size-1 `rmm::device_uvector<T>` internally.
+All host<->device transfers go through a pinned-host bounce buffer, avoiding the implicit
+stream synchronization overhead that pageable memory copies require.
+
+Use this for scalar input/outputs into device kernels, e.g., reduction results, null count, etc.
+
+Key properties:
+- Owns `rmm::device_uvector<T> _storage{1, stream, mr}`.
+- Uses a `cudf::detail::host_vector<T>` bounce buffer (pinned host memory) for `value()` and
+  `set_value_async()`, enabling fully async host<->device transfers.
+- Requires `T` to be trivially copyable (enforced via `static_assert`).
+- Exposes `data()`, `value()`, `set_value_async()`, etc.
 
 ```c++
 // Allocates device memory for a single int using the specified resource and stream
 // and initializes the value to 42
-rmm::device_scalar<int> int_scalar{42, stream, mr};
+cudf::detail::device_scalar<int> int_scalar{42, stream, mr};
 
 // scalar.data() returns pointer to value in device memory
 kernel<<<...>>>(int_scalar.data(),...);
 
-// scalar.value() synchronizes the scalar's stream and copies the
-// value from device to host and returns the value
-int host_value = int_scalar.value();
-```
-
-##### cudf::detail::device_scalar<T>
-Acts as a drop-in replacement for `rmm::device_scalar<T>`, with the key difference
-being the use of pinned host memory as a bounce buffer for data transfers.
-It is recommended for internal use to avoid the implicit synchronization overhead caused by
-memcpy operations on pageable host memory.
-
-```c++
-// Same as the case with rmm::device_scalar<T> above
-cudf::detail::device_scalar<int> int_scalar{42, stream, mr};
-kernel<<<...>>>(int_scalar.data(),...);
-
-// Note: This device-to-host transfer uses host-pinned bounce buffer for efficient memcpy
-int host_value = int_scalar.value();
+// value() copies device->host via pinned bounce buffer; no implicit stream sync
+int host_value = int_scalar.value(stream);
 ```
 
 #### rmm::device_vector<T>

cpp/include/cudf/detail/device_scalar.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,15 +10,22 @@
 #include <cudf/detail/utilities/vector_factories.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_scalar.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <type_traits>
+#include <utility>
+
 namespace CUDF_EXPORT cudf {
 namespace detail {
 
 template <typename T>
-class device_scalar : public rmm::device_scalar<T> {
+class device_scalar {
  public:
+  static_assert(std::is_trivially_copyable_v<T>,
+                "cudf::detail::device_scalar<T> requires T to be trivially copyable");
+  using value_type = T;
+
 #ifdef __CUDACC__
 #pragma nv_exec_check_disable
 #endif
@@ -35,48 +42,52 @@ class device_scalar : public rmm::device_scalar<T> {
   explicit device_scalar(
     rmm::cuda_stream_view stream,
     rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
-    : rmm::device_scalar<T>(stream, mr), bounce_buffer{make_pinned_vector<T>(1, stream)}
+    : _storage{1, stream, std::move(mr)}, bounce_buffer{make_pinned_vector<T>(1, stream)}
   {
   }
 
   explicit device_scalar(
     T const& initial_value,
     rmm::cuda_stream_view stream,
     rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
-    : rmm::device_scalar<T>(stream, mr), bounce_buffer{make_pinned_vector<T>(1, stream)}
+    : _storage{1, stream, std::move(mr)}, bounce_buffer{make_pinned_vector<T>(1, stream)}
   {
-    bounce_buffer[0] = initial_value;
-    cuda_memcpy_async<T>(device_span<T>{this->data(), 1}, bounce_buffer, stream);
+    set_value_async(initial_value, stream);
   }
 
   device_scalar(device_scalar const& other,
                 rmm::cuda_stream_view stream,
                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
-    : rmm::device_scalar<T>(other, stream, mr), bounce_buffer{make_pinned_vector<T>(1, stream)}
+    : _storage{other._storage, stream, mr}, bounce_buffer{make_pinned_vector<T>(1, stream)}
   {
   }
 
   [[nodiscard]] T value(rmm::cuda_stream_view stream) const
   {
-    cuda_memcpy<T>(bounce_buffer, device_span<T const>(this->data(), 1), stream);
+    cuda_memcpy<T>(bounce_buffer, device_span<T const>{data(), 1}, stream);
     return std::move(bounce_buffer[0]);
   }
 
   void set_value_async(T const& value, rmm::cuda_stream_view stream)
   {
     bounce_buffer[0] = value;
-    cuda_memcpy_async<T>(device_span<T>(this->data(), 1), bounce_buffer, stream);
+    cuda_memcpy_async<T>(device_span<T>{data(), 1}, bounce_buffer, stream);
   }
 
   void set_value_async(T&& value, rmm::cuda_stream_view stream)
   {
     bounce_buffer[0] = std::move(value);
-    cuda_memcpy_async<T>(device_span<T>{this->data(), 1}, bounce_buffer, stream);
+    cuda_memcpy_async<T>(device_span<T>{data(), 1}, bounce_buffer, stream);
   }
 
   void set_value_to_zero_async(rmm::cuda_stream_view stream) { set_value_async(T{}, stream); }
 
+  [[nodiscard]] T* data() noexcept { return _storage.data(); }
+
+  [[nodiscard]] T const* data() const noexcept { return _storage.data(); }
+
  private:
+  rmm::device_uvector<T> _storage;
   mutable cudf::detail::host_vector<T> bounce_buffer;
 };
 

cpp/include/cudf/scalar/scalar.hpp

@@ -14,7 +14,6 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
-#include <rmm/device_scalar.hpp>
 
 #include <span>
 #include <string_view>
@@ -189,7 +188,7 @@ class fixed_width_scalar : public scalar {
   [[nodiscard]] T const* data() const;
 
  protected:
-  rmm::device_scalar<T> _data;  ///< device memory containing the value
+  cudf::detail::device_scalar<T> _data;  ///< device memory containing the value
 
   /**
    * @brief Construct a new fixed width scalar object.
@@ -212,7 +211,7 @@ class fixed_width_scalar : public scalar {
    * @param stream CUDA stream used for device memory operations.
    * @param mr Device memory resource to use for device memory allocation.
    */
-  fixed_width_scalar(rmm::device_scalar<T>&& data,
+  fixed_width_scalar(cudf::detail::device_scalar<T>&& data,
                      bool is_valid                     = true,
                      rmm::cuda_stream_view stream      = cudf::get_default_stream(),
                      rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
@@ -274,7 +273,7 @@ class numeric_scalar : public detail::fixed_width_scalar<T> {
    * @param stream CUDA stream used for device memory operations.
    * @param mr Device memory resource to use for device memory allocation.
    */
-  numeric_scalar(rmm::device_scalar<T>&& data,
+  numeric_scalar(cudf::detail::device_scalar<T>&& data,
                  bool is_valid                     = true,
                  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
                  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
@@ -366,7 +365,7 @@ class fixed_point_scalar : public scalar {
    * @param stream CUDA stream used for device memory operations.
    * @param mr Device memory resource to use for device memory allocation.
    */
-  fixed_point_scalar(rmm::device_scalar<rep_type>&& data,
+  fixed_point_scalar(cudf::detail::device_scalar<rep_type>&& data,
                      numeric::scale_type scale,
                      bool is_valid                     = true,
                      rmm::cuda_stream_view stream      = cudf::get_default_stream(),
@@ -402,7 +401,7 @@ class fixed_point_scalar : public scalar {
   [[nodiscard]] rep_type const* data() const;
 
  protected:
-  rmm::device_scalar<rep_type> _data;  ///< device memory containing the value
+  cudf::detail::device_scalar<rep_type> _data;  ///< device memory containing the value
 };
 
 /**
@@ -476,7 +475,7 @@ class string_scalar : public scalar {
    * @param stream CUDA stream used for device memory operations.
    * @param mr Device memory resource to use for device memory allocation.
    */
-  string_scalar(rmm::device_scalar<value_type>& data,
+  string_scalar(cudf::detail::device_scalar<value_type>& data,
                 bool is_valid                     = true,
                 rmm::cuda_stream_view stream      = cudf::get_default_stream(),
                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
@@ -585,7 +584,7 @@ class chrono_scalar : public detail::fixed_width_scalar<T> {
    * @param stream CUDA stream used for device memory operations.
    * @param mr Device memory resource to use for device memory allocation.
    */
-  chrono_scalar(rmm::device_scalar<T>&& data,
+  chrono_scalar(cudf::detail::device_scalar<T>&& data,
                 bool is_valid                     = true,
                 rmm::cuda_stream_view stream      = cudf::get_default_stream(),
                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());

cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu

@@ -5,14 +5,13 @@
 
 #include <cudf_test/cudf_gtest.hpp>
 
+#include <cudf/detail/device_scalar.hpp>
 #include <cudf/reduction/bloom_filter.cuh>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
 #include <cudf_streaming/detail/device_bloom_filter.hpp>
 
-#include <rmm/device_scalar.hpp>
-
 #include <cuco/bloom_filter_ref.cuh>
 #include <cuco/extent.cuh>
 #include <cuco/hash_functions.cuh>
@@ -40,7 +39,7 @@ TEST(BloomFilterPolicyTest, UsesBlocksBeyondFormerArrowLimit)
   constexpr auto num_blocks       = arrow_max_blocks + 1;
   constexpr auto upper_hash       = std::numeric_limits<std::uint32_t>::max();
   auto const stream               = cudf::get_default_stream();
-  rmm::device_scalar<std::uint32_t> index{0, stream};
+  cudf::detail::device_scalar<std::uint32_t> index{0, stream};
 
   block_index_kernel<<<1, 1, 0, stream.value()>>>(upper_hash, num_blocks, index.data());
   CUDF_CHECK_CUDA(stream.value());

cpp/src/copying/get_element.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -60,7 +60,7 @@ struct get_element_functor {
   {
     auto device_col = column_device_view::create(input, stream);
 
-    rmm::device_scalar<string_view> temp_data(stream, mr);
+    cudf::detail::device_scalar<string_view> temp_data(stream, mr);
     cudf::detail::device_scalar<bool> temp_valid(stream, mr);
 
     device_single_thread(

cpp/src/dictionary/encode.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -24,7 +24,6 @@
 #include <cudf/utilities/memory_resource.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/polymorphic_allocator.hpp>

cpp/src/groupby/hash/compute_single_pass_aggs.cuh

@@ -93,7 +93,7 @@ std::pair<rmm::device_uvector<size_type>, bool> compute_single_pass_aggs(
   rmm::device_uvector<size_type> block_cardinality(grid_size, stream);
 
   // Flag indicating whether a global memory aggregation fallback is required or not.
-  rmm::device_scalar<cuda::std::atomic_flag> needs_global_memory_fallback(stream);
+  rmm::device_uvector<cuda::std::atomic_flag> needs_global_memory_fallback(1, stream);
   CUDF_CUDA_TRY(cudaMemsetAsync(
     needs_global_memory_fallback.data(), 0, sizeof(cuda::std::atomic_flag), stream.value()));
 
@@ -110,8 +110,8 @@ std::pair<rmm::device_uvector<size_type>, bool> compute_single_pass_aggs(
 
   auto const needs_fallback = [&] {
     cuda::std::atomic_flag h_needs_fallback;
-    // Cannot use `device_scalar::value` as it requires a copy constructor, which
-    // `atomic_flag` doesn't have.
+    // Cannot use a value-returning helper because atomic_flag is not copy-constructible;
+    // copy the raw bytes back to host instead.
     CUDF_CUDA_TRY(cudf::detail::memcpy_async(&h_needs_fallback,
                                              needs_global_memory_fallback.data(),
                                              sizeof(cuda::std::atomic_flag),

cpp/src/scalar/scalar.cpp

@@ -4,6 +4,7 @@
  */
 
 #include <cudf/column/column.hpp>
+#include <cudf/detail/device_scalar.hpp>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/structs/utilities.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
@@ -65,7 +66,7 @@ string_scalar::string_scalar(string_scalar const& other,
 {
 }
 
-string_scalar::string_scalar(rmm::device_scalar<value_type>& data,
+string_scalar::string_scalar(cudf::detail::device_scalar<value_type>& data,
                              bool is_valid,
                              rmm::cuda_stream_view stream,
                              rmm::device_async_resource_ref mr)
@@ -139,7 +140,7 @@ fixed_point_scalar<T>::fixed_point_scalar(T value,
 }
 
 template <typename T>
-fixed_point_scalar<T>::fixed_point_scalar(rmm::device_scalar<rep_type>&& data,
+fixed_point_scalar<T>::fixed_point_scalar(cudf::detail::device_scalar<rep_type>&& data,
                                           numeric::scale_type scale,
                                           bool is_valid,
                                           rmm::cuda_stream_view stream,
@@ -206,7 +207,7 @@ fixed_width_scalar<T>::fixed_width_scalar(T value,
 }
 
 template <typename T>
-fixed_width_scalar<T>::fixed_width_scalar(rmm::device_scalar<T>&& data,
+fixed_width_scalar<T>::fixed_width_scalar(cudf::detail::device_scalar<T>&& data,
                                           bool is_valid,
                                           rmm::cuda_stream_view stream,
                                           rmm::device_async_resource_ref mr)
@@ -290,11 +291,11 @@ numeric_scalar<T>::numeric_scalar(T value,
 }
 
 template <typename T>
-numeric_scalar<T>::numeric_scalar(rmm::device_scalar<T>&& data,
+numeric_scalar<T>::numeric_scalar(cudf::detail::device_scalar<T>&& data,
                                   bool is_valid,
                                   rmm::cuda_stream_view stream,
                                   rmm::device_async_resource_ref mr)
-  : detail::fixed_width_scalar<T>(std::forward<rmm::device_scalar<T>>(data), is_valid, stream, mr)
+  : detail::fixed_width_scalar<T>(std::move(data), is_valid, stream, mr)
 {
 }
 
@@ -337,11 +338,11 @@ chrono_scalar<T>::chrono_scalar(T value,
 }
 
 template <typename T>
-chrono_scalar<T>::chrono_scalar(rmm::device_scalar<T>&& data,
+chrono_scalar<T>::chrono_scalar(cudf::detail::device_scalar<T>&& data,
                                 bool is_valid,
                                 rmm::cuda_stream_view stream,
                                 rmm::device_async_resource_ref mr)
-  : detail::fixed_width_scalar<T>(std::forward<rmm::device_scalar<T>>(data), is_valid, stream, mr)
+  : detail::fixed_width_scalar<T>(std::move(data), is_valid, stream, mr)
 {
 }
 

cpp/src/transform/transform.cu

@@ -6,6 +6,7 @@
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/device_scalar.hpp>
 #include <cudf/detail/null_mask.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
@@ -1062,7 +1063,7 @@ std::unique_ptr<table> execute_transform(std::string const& udf,
   auto stencil_arg       = stencil.has_value() ? stencil->first : nullptr;
   auto stencil_has_nulls = stencil.has_value() ? (stencil->second > 0) : false;
 
-  rmm::device_scalar<int32_t> d_max_error(static_cast<int32_t>(errc::SUCCESS), stream, mr);
+  cudf::detail::device_scalar<int32_t> d_max_error(static_cast<int32_t>(errc::SUCCESS), stream, mr);
 
   jit_transform::run(is_null_aware == null_aware::YES,
                      user_data.has_value(),
@@ -1253,7 +1254,7 @@ std::unique_ptr<table> transform_lto(std::span<uint8_t const> udf,
   auto precompiled_kernel_fragment = dispatch_lto_kernel_fragment(
     is_null_aware == null_aware::YES, user_data.has_value(), inputs, output_columns);
 
-  rmm::device_scalar<int32_t> d_max_error(static_cast<int32_t>(errc::SUCCESS), stream, mr);
+  cudf::detail::device_scalar<int32_t> d_max_error(static_cast<int32_t>(errc::SUCCESS), stream, mr);
 
   jit_transform::run_lto(precompiled_kernel_fragment,
                          is_null_aware == null_aware::YES,