From db3ba19052629332091ccea91001023a361f8fbb Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Tue, 4 Aug 2026 02:59:00 -0700 Subject: [PATCH 1/2] refactor: replace rmm::device_scalar with cudf::detail::device_scalar Introduce a self-contained cudf::detail::device_scalar wrapper (size-1 rmm::device_uvector storage) and migrate all libcudf source, tests, and libcudf_streaming to use it. Updates the scalar hierarchy ctor signatures from rmm::device_scalar to cudf::detail::device_scalar (public API change), refreshes DEVELOPER_GUIDE.md guidance. This completes plan drop-rmm-device-scalar.md. A follow-up change removes the final rmm::device_scalar usage in compute_single_pass_aggs.cuh. --- .../developer_guide/DEVELOPER_GUIDE.md | 39 ++++++++----------- cpp/include/cudf/detail/device_scalar.hpp | 33 ++++++++++------ cpp/include/cudf/scalar/scalar.hpp | 15 ++++--- .../tests/streaming/test_bloom_filter.cu | 5 +-- cpp/src/copying/get_element.cu | 4 +- cpp/src/dictionary/encode.cu | 3 +- cpp/src/scalar/scalar.cpp | 15 +++---- cpp/src/transform/transform.cu | 5 ++- .../device_atomics/device_atomics_test.cu | 7 ++-- 9 files changed, 64 insertions(+), 62 deletions(-) diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index d8758dcacbf5..d6589f283ed5 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/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 -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` of length 1. +#### cudf::detail::device_scalar +A self-contained device scalar that owns a size-1 `rmm::device_uvector` 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 _storage{1, stream, mr}`. +- Uses a `cudf::detail::host_vector` 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_scalar{42, stream, mr}; +cudf::detail::device_scalar 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 -Acts as a drop-in replacement for `rmm::device_scalar`, 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 above -cudf::detail::device_scalar 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 diff --git a/cpp/include/cudf/detail/device_scalar.hpp b/cpp/include/cudf/detail/device_scalar.hpp index e2380ef8cb1d..5b9757e75757 100644 --- a/cpp/include/cudf/detail/device_scalar.hpp +++ b/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 #include -#include +#include #include +#include +#include + namespace CUDF_EXPORT cudf { namespace detail { template -class device_scalar : public rmm::device_scalar { +class device_scalar { public: + static_assert(std::is_trivially_copyable_v, + "cudf::detail::device_scalar requires T to be trivially copyable"); + using value_type = T; + #ifdef __CUDACC__ #pragma nv_exec_check_disable #endif @@ -35,7 +42,7 @@ class device_scalar : public rmm::device_scalar { explicit device_scalar( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) - : rmm::device_scalar(stream, mr), bounce_buffer{make_pinned_vector(1, stream)} + : _storage{1, stream, std::move(mr)}, bounce_buffer{make_pinned_vector(1, stream)} { } @@ -43,40 +50,44 @@ class device_scalar : public rmm::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(stream, mr), bounce_buffer{make_pinned_vector(1, stream)} + : _storage{1, stream, std::move(mr)}, bounce_buffer{make_pinned_vector(1, stream)} { - bounce_buffer[0] = initial_value; - cuda_memcpy_async(device_span{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(other, stream, mr), bounce_buffer{make_pinned_vector(1, stream)} + : _storage{other._storage, stream, mr}, bounce_buffer{make_pinned_vector(1, stream)} { } [[nodiscard]] T value(rmm::cuda_stream_view stream) const { - cuda_memcpy(bounce_buffer, device_span(this->data(), 1), stream); + cuda_memcpy(bounce_buffer, device_span{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(device_span(this->data(), 1), bounce_buffer, stream); + cuda_memcpy_async(device_span{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(device_span{this->data(), 1}, bounce_buffer, stream); + cuda_memcpy_async(device_span{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 _storage; mutable cudf::detail::host_vector bounce_buffer; }; diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 12d35eec400d..1edf778cc46f 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -14,7 +14,6 @@ #include #include -#include #include #include @@ -189,7 +188,7 @@ class fixed_width_scalar : public scalar { [[nodiscard]] T const* data() const; protected: - rmm::device_scalar _data; ///< device memory containing the value + cudf::detail::device_scalar _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&& data, + fixed_width_scalar(cudf::detail::device_scalar&& 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 { * @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&& data, + numeric_scalar(cudf::detail::device_scalar&& 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&& data, + fixed_point_scalar(cudf::detail::device_scalar&& 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 _data; ///< device memory containing the value + cudf::detail::device_scalar _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& data, + string_scalar(cudf::detail::device_scalar& 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 { * @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&& data, + chrono_scalar(cudf::detail::device_scalar&& 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()); diff --git a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu index b8fc5b2e978a..c6d04dd8df21 100644 --- a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu +++ b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu @@ -5,14 +5,13 @@ #include +#include #include #include #include #include -#include - #include #include #include @@ -40,7 +39,7 @@ TEST(BloomFilterPolicyTest, UsesBlocksBeyondFormerArrowLimit) constexpr auto num_blocks = arrow_max_blocks + 1; constexpr auto upper_hash = std::numeric_limits::max(); auto const stream = cudf::get_default_stream(); - rmm::device_scalar index{0, stream}; + cudf::detail::device_scalar index{0, stream}; block_index_kernel<<<1, 1, 0, stream.value()>>>(upper_hash, num_blocks, index.data()); CUDF_CHECK_CUDA(stream.value()); diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 1d989f7c9fb5..85bdf99f4941 100644 --- a/cpp/src/copying/get_element.cu +++ b/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 temp_data(stream, mr); + cudf::detail::device_scalar temp_data(stream, mr); cudf::detail::device_scalar temp_valid(stream, mr); device_single_thread( diff --git a/cpp/src/dictionary/encode.cu b/cpp/src/dictionary/encode.cu index a43d18c080a4..5062f4835ed0 100644 --- a/cpp/src/dictionary/encode.cu +++ b/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 #include -#include #include #include #include diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index cd437c360496..6f3887585edf 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -4,6 +4,7 @@ */ #include +#include #include #include #include @@ -65,7 +66,7 @@ string_scalar::string_scalar(string_scalar const& other, { } -string_scalar::string_scalar(rmm::device_scalar& data, +string_scalar::string_scalar(cudf::detail::device_scalar& data, bool is_valid, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) @@ -139,7 +140,7 @@ fixed_point_scalar::fixed_point_scalar(T value, } template -fixed_point_scalar::fixed_point_scalar(rmm::device_scalar&& data, +fixed_point_scalar::fixed_point_scalar(cudf::detail::device_scalar&& data, numeric::scale_type scale, bool is_valid, rmm::cuda_stream_view stream, @@ -206,7 +207,7 @@ fixed_width_scalar::fixed_width_scalar(T value, } template -fixed_width_scalar::fixed_width_scalar(rmm::device_scalar&& data, +fixed_width_scalar::fixed_width_scalar(cudf::detail::device_scalar&& data, bool is_valid, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) @@ -290,11 +291,11 @@ numeric_scalar::numeric_scalar(T value, } template -numeric_scalar::numeric_scalar(rmm::device_scalar&& data, +numeric_scalar::numeric_scalar(cudf::detail::device_scalar&& data, bool is_valid, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) - : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) + : detail::fixed_width_scalar(std::move(data), is_valid, stream, mr) { } @@ -337,11 +338,11 @@ chrono_scalar::chrono_scalar(T value, } template -chrono_scalar::chrono_scalar(rmm::device_scalar&& data, +chrono_scalar::chrono_scalar(cudf::detail::device_scalar&& data, bool is_valid, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) - : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) + : detail::fixed_width_scalar(std::move(data), is_valid, stream, mr) { } diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu index 9e08a0950557..7175cebd8473 100644 --- a/cpp/src/transform/transform.cu +++ b/cpp/src/transform/transform.cu @@ -6,6 +6,7 @@ #include #include #include +#include #include #include #include @@ -1062,7 +1063,7 @@ std::unique_ptr 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 d_max_error(static_cast(errc::SUCCESS), stream, mr); + cudf::detail::device_scalar d_max_error(static_cast(errc::SUCCESS), stream, mr); jit_transform::run(is_null_aware == null_aware::YES, user_data.has_value(), @@ -1253,7 +1254,7 @@ std::unique_ptr
transform_lto(std::span 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 d_max_error(static_cast(errc::SUCCESS), stream, mr); + cudf::detail::device_scalar d_max_error(static_cast(errc::SUCCESS), stream, mr); jit_transform::run_lto(precompiled_kernel_fragment, is_null_aware == null_aware::YES, diff --git a/cpp/tests/device_atomics/device_atomics_test.cu b/cpp/tests/device_atomics/device_atomics_test.cu index 90ef2f17533d..50a5c81be6ad 100644 --- a/cpp/tests/device_atomics/device_atomics_test.cu +++ b/cpp/tests/device_atomics/device_atomics_test.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,6 +9,7 @@ #include #include +#include #include #include #include @@ -17,8 +18,6 @@ #include #include -#include - #include #include @@ -264,7 +263,7 @@ class Atomic128Test : public cudf::test::BaseFixture { __int128_t add_value, __int128_t expected_result) { - rmm::device_scalar<__int128_t> d_target(initial_value, cudf::get_default_stream()); + cudf::detail::device_scalar<__int128_t> d_target(initial_value, cudf::get_default_stream()); test_single_atomic_add_kernel<<<32, 256, 0, cudf::get_default_stream().value()>>>( d_target.data(), add_value); CUDF_CHECK_CUDA(cudf::get_default_stream().value()); From 307032e3c42139a61ff49cc8efcba5a4cd456932 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Tue, 4 Aug 2026 03:33:21 -0700 Subject: [PATCH 2/2] refactor: drop last rmm::device_scalar usage in groupby hash aggs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace the exemption rmm::device_scalar in compute_single_pass_aggs.cuh with rmm::device_uvector of size 1 — semantically identical (device_scalar is a size-1 device_uvector under the hood) but eliminates the last direct rmm::device_scalar reference from libcudf source. atomic_flag cannot use cudf::detail::device_scalar because it is not trivially copyable, but device_uvector has no such requirement. Completes plan drop-final-rmm-device-scalar-usage.md. --- cpp/src/groupby/hash/compute_single_pass_aggs.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh index 39fd7611e317..b3217b6a430d 100644 --- a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh +++ b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh @@ -93,7 +93,7 @@ std::pair, bool> compute_single_pass_aggs( rmm::device_uvector block_cardinality(grid_size, stream); // Flag indicating whether a global memory aggregation fallback is required or not. - rmm::device_scalar needs_global_memory_fallback(stream); + rmm::device_uvector 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, 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),