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

cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md

@@ -749,37 +749,26 @@ 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 for internal libcudf code that needs a single trivially copyable
+value in device memory, such as a reduction result, temporary counter, or kernel status value.
+
+Use this for internal scalar input/output with device kernels. Public libcudf APIs should use
+`cudf::scalar` and derived public scalar classes instead of this detail type.
+
+It exposes `data()` for kernels and `value()`/`set_value_async()` for stream-ordered host/device
+transfers.
 
 ```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(),...);
+kernel<<<..., stream>>>(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 the device value to the host on the specified stream
+int host_value = int_scalar.value(stream);
 ```
 
 #### rmm::device_vector<T>

cpp/include/cudf/detail/device_scalar.hpp

@@ -9,17 +9,24 @@
 #include <cudf/detail/utilities/host_vector.hpp>
 #include <cudf/detail/utilities/vector_factories.hpp>
 
-#include <rmm/device_scalar.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/stream>
 
+#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
@@ -36,48 +43,52 @@ class device_scalar : public rmm::device_scalar<T> {
   explicit device_scalar(
     cuda::stream_ref 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,
     cuda::stream_ref 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,
                 cuda::stream_ref 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(cuda::stream_ref 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, cuda::stream_ref 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, cuda::stream_ref 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(cuda::stream_ref 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/reduction/detail/reduction.cuh

@@ -139,7 +139,7 @@ std::unique_ptr<scalar> reduce(InputIterator d_in,
                             initial_value,
                             stream.get());
 
-  return std::make_unique<cudf::string_scalar>(dev_result, true, stream, mr);
+  return std::make_unique<cudf::string_scalar>(dev_result.value(stream), true, stream, mr);
 }
 
 /**

cpp/include/cudf/scalar/scalar.hpp

@@ -213,7 +213,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,
                      cuda::stream_ref stream           = cudf::get_default_stream(),
                      rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
@@ -267,18 +267,37 @@ class numeric_scalar : public detail::fixed_width_scalar<T> {
                  cuda::stream_ref stream           = cudf::get_default_stream(),
                  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new numeric scalar object from another scalar.
+   *
+   * The input scalar's type must exactly match this scalar's type.
+   *
+   * @throws cudf::data_type_error if the input scalar type does not match.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit numeric_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
   /**
    * @brief Construct a new numeric scalar object from existing device memory.
    *
+   * @deprecated Use the cudf::scalar constructor instead.
+   *
    * @param data The scalar's data in device memory.
    * @param is_valid Whether the value held by the scalar is valid.
    * @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,
-                 bool is_valid                     = true,
-                 cuda::stream_ref stream           = cudf::get_default_stream(),
-                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+  [[deprecated("Use the cudf::scalar constructor instead.")]] numeric_scalar(
+    rmm::device_scalar<T>&& data,
+    bool is_valid                     = true,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 };
 
 /**
@@ -358,20 +377,40 @@ class fixed_point_scalar : public scalar {
                      cuda::stream_ref stream           = cudf::get_default_stream(),
                      rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new fixed_point scalar object from another scalar.
+   *
+   * The input scalar's type ID must match this scalar's type. The input scalar's scale is
+   * preserved.
+   *
+   * @throws cudf::data_type_error if the input scalar type ID does not match.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit fixed_point_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
   /**
    * @brief Construct a new fixed_point scalar object from existing device memory.
    *
+   * @deprecated Use the cudf::scalar constructor instead.
+   *
    * @param data The scalar's data in device memory.
    * @param scale The scale of the fixed_point scalar.
    * @param is_valid Whether the value held by the scalar is valid.
    * @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,
-                     numeric::scale_type scale,
-                     bool is_valid                     = true,
-                     cuda::stream_ref stream           = cudf::get_default_stream(),
-                     rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+  [[deprecated("Use the cudf::scalar constructor instead.")]] fixed_point_scalar(
+    rmm::device_scalar<rep_type>&& data,
+    numeric::scale_type scale,
+    bool is_valid                     = true,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
   /**
    * @brief Get the value of the scalar.
@@ -466,20 +505,39 @@ class string_scalar : public scalar {
                 cuda::stream_ref stream           = cudf::get_default_stream(),
                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new string scalar object from another scalar.
+   *
+   * The input scalar's type must be STRING.
+   *
+   * @throws cudf::data_type_error if the input scalar type is not STRING.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit string_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
   /**
    * @brief Construct a new string scalar object from string_view in device memory.
    *
    * Note that this function copies the data pointed by string_view.
    *
+   * @deprecated Use the cudf::scalar constructor instead.
+   *
    * @param data The device_scalar of string_view pointing to the string value to copy.
    * @param is_valid Whether the value held by the scalar is valid.
    * @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,
-                bool is_valid                     = true,
-                cuda::stream_ref stream           = cudf::get_default_stream(),
-                rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+  [[deprecated("Use the cudf::scalar constructor instead.")]] string_scalar(
+    rmm::device_scalar<value_type>& data,
+    bool is_valid                     = true,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
   /**
    * @brief Construct a new string scalar object by moving an existing string data buffer.
@@ -576,18 +634,37 @@ class chrono_scalar : public detail::fixed_width_scalar<T> {
                 cuda::stream_ref stream           = cudf::get_default_stream(),
                 rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new chrono scalar object from another scalar.
+   *
+   * The input scalar's type must exactly match this scalar's type.
+   *
+   * @throws cudf::data_type_error if the input scalar type does not match.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit chrono_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
   /**
    * @brief Construct a new chrono scalar object from existing device memory.
    *
+   * @deprecated Use the cudf::scalar constructor instead.
+   *
    * @param data The scalar's data in device memory.
    * @param is_valid Whether the value held by the scalar is valid.
    * @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,
-                bool is_valid                     = true,
-                cuda::stream_ref stream           = cudf::get_default_stream(),
-                rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+  [[deprecated("Use the cudf::scalar constructor instead.")]] chrono_scalar(
+    rmm::device_scalar<T>&& data,
+    bool is_valid                     = true,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 };
 
 /**
@@ -622,6 +699,25 @@ class timestamp_scalar : public chrono_scalar<T> {
                    cuda::stream_ref stream           = cudf::get_default_stream(),
                    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new timestamp scalar object from another scalar.
+   *
+   * The input scalar's type must exactly match this scalar's type.
+   *
+   * @throws cudf::data_type_error if the input scalar type does not match.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit timestamp_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
+    : chrono_scalar<T>(data, stream, mr)
+  {
+  }
+
   /**
    * @brief Construct a new timestamp scalar object from a duration that is
    * convertible to T::duration
@@ -678,6 +774,25 @@ class duration_scalar : public chrono_scalar<T> {
                   cuda::stream_ref stream           = cudf::get_default_stream(),
                   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+  /**
+   * @brief Construct a new duration scalar object from another scalar.
+   *
+   * The input scalar's type must exactly match this scalar's type.
+   *
+   * @throws cudf::data_type_error if the input scalar type does not match.
+   *
+   * @param data The scalar to copy.
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation.
+   */
+  explicit duration_scalar(
+    scalar const& data,
+    cuda::stream_ref stream           = cudf::get_default_stream(),
+    rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
+    : chrono_scalar<T>(data, stream, mr)
+  {
+  }
+
   /**
    * @brief Construct a new duration scalar object from tick counts.
    *