From ae3b6eaf5f6828c44c1a17710761efdae7edc1fe Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 6 Jul 2026 21:48:59 +0000 Subject: [PATCH 1/5] Make stream and memory resource explicit in hyperloglog APIs --- .../cuda/experimental/__cuco/hyperloglog.cuh | 140 +++++------------- .../experimental/__cuco/hyperloglog_ref.cuh | 39 ++--- .../test/cuco/hyperloglog/test_hyperloglog.cu | 61 +++++--- 3 files changed, 100 insertions(+), 140 deletions(-) diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh index 94bac39c116..68e1762fda5 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh @@ -100,71 +100,56 @@ public: //! //! @note This function synchronizes the given stream. //! + //! @param __stream CUDA stream used to initialize the object //! @param __memory_resource A memory resource used for allocating device storage //! @param __sketch_size_kb Maximum sketch size in KB //! @param __policy The policy used to hash items and finalize the estimate - //! @param __stream CUDA stream used to initialize the object //! //! @throw If sketch size implies precision outside [4, 18]. template - constexpr hyperloglog(_MemoryResource_&& __memory_resource, + constexpr hyperloglog(::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, sketch_size_kb __sketch_size_kb = sketch_size_kb{32.0}, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) - : hyperloglog{::cuda::std::forward<_MemoryResource_>(__memory_resource), + const _Policy& __policy = {}) + : hyperloglog{__stream, + ::cuda::std::forward<_MemoryResource_>(__memory_resource), __to_precision(__sketch_size_kb), - __policy, - __stream} + __policy} {} //! @brief Constructs a `hyperloglog` host object. //! //! @note This function synchronizes the given stream. //! - //! @param __sketch_size_kb Maximum sketch size in KB - //! @param __policy The policy used to hash items and finalize the estimate //! @param __stream CUDA stream used to initialize the object - //! - //! @throw If sketch size implies precision outside [4, 18]. - constexpr hyperloglog(sketch_size_kb __sketch_size_kb = sketch_size_kb{32.0}, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) - : hyperloglog{__to_precision(__sketch_size_kb), __policy, __stream} - {} - - //! @brief Constructs a `hyperloglog` host object. - //! - //! @note This function synchronizes the given stream. - //! //! @param __memory_resource A memory resource used for allocating device storage //! @param __sd Desired standard deviation for the approximation error //! @param __policy The policy used to hash items and finalize the estimate - //! @param __stream CUDA stream used to initialize the object //! //! @throw If standard deviation implies precision outside [4, 18]. template - constexpr hyperloglog(_MemoryResource_&& __memory_resource, + constexpr hyperloglog(::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, standard_deviation __sd, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) - : hyperloglog{::cuda::std::forward<_MemoryResource_>(__memory_resource), __to_precision(__sd), __policy, __stream} + const _Policy& __policy = {}) + : hyperloglog{__stream, ::cuda::std::forward<_MemoryResource_>(__memory_resource), __to_precision(__sd), __policy} {} //! @brief Constructs a `hyperloglog` host object. //! //! @note This function synchronizes the given stream. //! + //! @param __stream CUDA stream used to initialize the object //! @param __memory_resource A memory resource used for allocating device storage //! @param __precision HyperLogLog precision parameter (determines number of registers as 2^precision) //! @param __policy The policy used to hash items and finalize the estimate - //! @param __stream CUDA stream used to initialize the object //! //! @throw If precision is outside [4, 18]. template - constexpr hyperloglog(_MemoryResource_&& __memory_resource, + constexpr hyperloglog(::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, precision __precision, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + const _Policy& __policy = {}) : __sketch_buffer{__stream, ::cuda::std::forward<_MemoryResource_>(__memory_resource), ref_type<>::sketch_bytes( @@ -177,45 +162,6 @@ public: clear_async(__stream); } - //! @brief Constructs a `hyperloglog` host object. - //! - //! @note This function synchronizes the given stream. - //! - //! @param __sd Desired standard deviation for the approximation error - //! @param __policy The policy used to hash items and finalize the estimate - //! @param __stream CUDA stream used to initialize the object - //! - //! @throw If standard deviation implies precision outside [4, 18]. - constexpr hyperloglog(standard_deviation __sd, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) - : hyperloglog{__to_precision(__sd), __policy, __stream} - {} - - //! @brief Constructs a `hyperloglog` host object. - //! - //! @note This function synchronizes the given stream. - //! - //! @param __precision HyperLogLog precision parameter (determines number of registers as 2^precision) - //! @param __policy The policy used to hash items and finalize the estimate - //! @param __stream CUDA stream used to initialize the object - //! - //! @throw If precision is outside [4, 18]. - constexpr hyperloglog(precision __precision, - const _Policy& __policy = {}, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) - : __sketch_buffer{__stream, - ::cuda::device_default_memory_pool(::cuda::device_ref{0}), - ref_type<>::sketch_bytes( - __precision_in_bounds(__precision, "HyperLogLog precision must be in [4, 18]")) - / sizeof(register_type), - ::cuda::no_init} - , __ref{::cuda::std::as_writable_bytes(::cuda::std::span{__sketch_buffer.data(), __sketch_buffer.size()}), - __policy} - { - clear_async(__stream); - } - ~hyperloglog() = default; hyperloglog(const hyperloglog&) = delete; @@ -230,7 +176,7 @@ public: //! @brief Asynchronously resets the estimator, i.e., clears the current count estimate. //! //! @param __stream CUDA stream this operation is executed in - constexpr void clear_async(::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) noexcept + constexpr void clear_async(::cuda::stream_ref __stream) noexcept { __ref.clear_async(__stream); } @@ -241,7 +187,7 @@ public: //! `clear_async`. //! //! @param __stream CUDA stream this operation is executed in - constexpr void clear(::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void clear(::cuda::stream_ref __stream) { __ref.clear(__stream); } @@ -252,14 +198,13 @@ public: //! std::is_convertible::value_type, //! _Tp> is `true` //! + //! @param __stream CUDA stream this operation is executed in //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items - //! @param __stream CUDA stream this operation is executed in template - constexpr void - add_async(_InputIt __first, _InputIt __last, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void add_async(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { - __ref.add_async(__first, __last, __stream); + __ref.add_async(__stream, __first, __last); } //! @brief Adds to be counted items to the estimator. @@ -271,14 +216,13 @@ public: //! std::is_convertible::value_type, //! _Tp> is `true` //! + //! @param __stream CUDA stream this operation is executed in //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items - //! @param __stream CUDA stream this operation is executed in template - constexpr void - add(_InputIt __first, _InputIt __last, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void add(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { - __ref.add(__first, __last, __stream); + __ref.add(__stream, __first, __last); } //! @brief Asynchronously merges the result of `other` estimator into `*this` estimator. @@ -288,13 +232,13 @@ public: //! @tparam _OtherScope Thread scope of `other` estimator //! @tparam _OtherMemoryResource Memory resource type of `other` estimator //! - //! @param __other Other estimator to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other Other estimator to be merged into `*this` template <::cuda::thread_scope _OtherScope, class _OtherMemoryResource> - constexpr void merge_async(const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void merge_async(::cuda::stream_ref __stream, + const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) { - __ref.merge_async(__other.__ref, __stream); + __ref.merge_async(__stream, __other.__ref); } //! @brief Merges the result of `other` estimator into `*this` estimator. @@ -307,13 +251,13 @@ public: //! @tparam _OtherScope Thread scope of `other` estimator //! @tparam _OtherMemoryResource Memory resource type of `other` estimator //! - //! @param __other Other estimator to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other Other estimator to be merged into `*this` template <::cuda::thread_scope _OtherScope, class _OtherMemoryResource> - constexpr void merge(const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void merge(::cuda::stream_ref __stream, + const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) { - __ref.merge(__other.__ref, __stream); + __ref.merge(__stream, __other.__ref); } //! @brief Asynchronously merges the result of `other` estimator reference into `*this` estimator. @@ -322,13 +266,12 @@ public: //! //! @tparam _OtherScope Thread scope of `other` estimator //! - //! @param __other_ref Other estimator reference to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other_ref Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - constexpr void merge_async(const ref_type<_OtherScope>& __other_ref, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void merge_async(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) { - __ref.merge_async(__other_ref, __stream); + __ref.merge_async(__stream, __other_ref); } //! @brief Merges the result of `other` estimator reference into `*this` estimator. @@ -340,13 +283,12 @@ public: //! //! @tparam _OtherScope Thread scope of `other` estimator //! - //! @param __other_ref Other estimator reference to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other_ref Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - constexpr void merge(const ref_type<_OtherScope>& __other_ref, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + constexpr void merge(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) { - __ref.merge(__other_ref, __stream); + __ref.merge(__stream, __other_ref); } //! @brief Compute the estimated distinct items count. @@ -356,15 +298,15 @@ public: //! @tparam _MemoryResource Host memory resource used for allocating the host buffer required to //! compute the final estimate by copying the sketch from device to host //! - //! @param __host_mr Host memory resource used for copying the sketch //! @param __stream CUDA stream this operation is executed in + //! @param __host_mr Host memory resource used for copying the sketch //! //! @return Approximate distinct items count template - [[nodiscard]] constexpr ::cuda::std::size_t estimate( - _HostMemoryResource __host_mr = {}, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) const + [[nodiscard]] constexpr ::cuda::std::size_t + estimate(::cuda::stream_ref __stream, _HostMemoryResource __host_mr = {}) const { - return __ref.estimate(__host_mr, __stream); + return __ref.estimate(__stream, __host_mr); } //! @brief Get device ref. diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh index ebfa9f99a6a..b5769186918 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh @@ -147,8 +147,7 @@ public: //! @brief Asynchronously resets the estimator, i.e., clears the current count estimate. //! //! @param __stream CUDA stream this operation is executed in - _CCCL_HOST_API constexpr void - clear_async(::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) noexcept + _CCCL_HOST_API constexpr void clear_async(::cuda::stream_ref __stream) noexcept { __impl.__clear_async(__stream); } @@ -159,7 +158,7 @@ public: //! `clear_async`. //! //! @param __stream CUDA stream this operation is executed in - _CCCL_HOST_API constexpr void clear(::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + _CCCL_HOST_API constexpr void clear(::cuda::stream_ref __stream) { __impl.__clear(__stream); } @@ -178,12 +177,11 @@ public: //! std::is_convertible::value_type, //! _Tp> is `true` //! + //! @param __stream CUDA stream this operation is executed in //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items - //! @param __stream CUDA stream this operation is executed in template - _CCCL_HOST_API constexpr void - add_async(_InputIt __first, _InputIt __last, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + _CCCL_HOST_API constexpr void add_async(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { __impl.__add_async(__first, __last, __stream); } @@ -197,12 +195,11 @@ public: //! std::is_convertible::value_type, //! _Tp> is `true` //! + //! @param __stream CUDA stream this operation is executed in //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items - //! @param __stream CUDA stream this operation is executed in template - _CCCL_HOST_API constexpr void - add(_InputIt __first, _InputIt __last, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + _CCCL_HOST_API constexpr void add(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { __impl.__add(__first, __last, __stream); } @@ -217,8 +214,12 @@ public: //! @param __group CUDA Cooperative group this operation is executed in //! @param __other Other estimator reference to be merged into `*this` template - _CCCL_DEVICE_API constexpr void merge(_CG __group, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) + _CCCL_DEVICE_API constexpr ::cuda::std::enable_if_t> + merge(_CG __group, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) { + // The enable_if above works around the same host/device overload preference issue as + // documented in `clear(_CG)`: `merge(_CG, ...)` would otherwise conflict with + // `merge(::cuda::stream_ref, ...)` when called from `hyperloglog::merge(::cuda::stream_ref, ...)`. __impl.__merge(__group, __other.__impl); } @@ -230,11 +231,11 @@ public: // Review: For host-side merges, consider validating `__other` is on the same device. //! @tparam _OtherScope Thread scope of `other` estimator //! - //! @param __other Other estimator reference to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - _CCCL_HOST_API constexpr void merge_async(const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + _CCCL_HOST_API constexpr void + merge_async(::cuda::stream_ref __stream, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) { __impl.__merge_async(__other.__impl, __stream); } @@ -248,11 +249,11 @@ public: //! //! @tparam _OtherScope Thread scope of `other` estimator //! - //! @param __other Other estimator reference to be merged into `*this` //! @param __stream CUDA stream this operation is executed in + //! @param __other Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - _CCCL_HOST_API constexpr void merge(const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other, - ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) + _CCCL_HOST_API constexpr void + merge(::cuda::stream_ref __stream, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) { __impl.__merge(__other.__impl, __stream); } @@ -275,13 +276,13 @@ public: //! @tparam _HostMemoryResource Host memory resource used for allocating the host buffer required to //! compute the final estimate by copying the sketch from device to host //! - //! @param __host_mr Host memory resource used for copying the sketch //! @param __stream CUDA stream this operation is executed in + //! @param __host_mr Host memory resource used for copying the sketch //! //! @return Approximate distinct items count template - [[nodiscard]] _CCCL_HOST_API constexpr ::cuda::std::size_t estimate( - _HostMemoryResource __host_mr = {}, ::cuda::stream_ref __stream = ::cuda::stream_ref{cudaStream_t{nullptr}}) const + [[nodiscard]] _CCCL_HOST_API constexpr ::cuda::std::size_t + estimate(::cuda::stream_ref __stream, _HostMemoryResource __host_mr = {}) const { return __impl.__estimate(__host_mr, __stream); } diff --git a/cudax/test/cuco/hyperloglog/test_hyperloglog.cu b/cudax/test/cuco/hyperloglog/test_hyperloglog.cu index 91220309018..f45cd5e4fea 100644 --- a/cudax/test/cuco/hyperloglog/test_hyperloglog.cu +++ b/cudax/test/cuco/hyperloglog/test_hyperloglog.cu @@ -12,8 +12,10 @@ #include #include +#include #include #include +#include #include #include @@ -74,13 +76,16 @@ C2H_TEST("HyperLogLog device ref", "[hyperloglog]", test_types) // Generate `num_items` distinct items thrust::sequence(items.begin(), items.end(), T{0}); + ::cuda::stream stream{::cuda::device_ref{0}}; + auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + // Initialize the estimator - estimator_type estimator{sketch_size_kb}; + estimator_type estimator{stream, mr, sketch_size_kb}; // Add all items to the estimator - estimator.add(items.begin(), items.end()); + estimator.add(stream, items.begin(), items.end()); - const auto host_estimate = estimator.estimate(); + const auto host_estimate = estimator.estimate(stream); thrust::device_vector device_estimate(1); estimate_kernel> @@ -114,28 +119,31 @@ C2H_TEST("HyperLogLog unique sequence", "[hyperloglog]", test_types) // Generate `num_items` distinct items thrust::sequence(items.begin(), items.end(), T{0}); + ::cuda::stream stream{::cuda::device_ref{0}}; + auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + // Initialize the estimator - estimator_type estimator{sketch_size_kb}; + estimator_type estimator{stream, mr, sketch_size_kb}; - REQUIRE(estimator.estimate() == 0); + REQUIRE(estimator.estimate(stream) == 0); // Add all items to the estimator - estimator.add(items.begin(), items.end()); + estimator.add(stream, items.begin(), items.end()); - const auto estimate = estimator.estimate(); + const auto estimate = estimator.estimate(stream); // Adding the same items again should not affect the result - estimator.add(items.begin(), items.begin() + num_items / 2); - REQUIRE(estimator.estimate() == estimate); + estimator.add(stream, items.begin(), items.begin() + num_items / 2); + REQUIRE(estimator.estimate(stream) == estimate); // Adding the same items again (might use shared memory code path) should not affect the result auto* ptr = thrust::raw_pointer_cast(items.data()); - estimator.add(ptr, ptr + num_items / 2); - REQUIRE(estimator.estimate() == estimate); + estimator.add(stream, ptr, ptr + num_items / 2); + REQUIRE(estimator.estimate(stream) == estimate); // Clearing the estimator should reset the estimate - estimator.clear(); - REQUIRE(estimator.estimate() == 0); + estimator.clear(stream); + REQUIRE(estimator.estimate(stream) == 0); const double relative_error = std::abs((static_cast(estimate) / static_cast(num_items)) - 1.0); @@ -178,14 +186,17 @@ C2H_TEST("HyperLogLog Spark parity deterministic", "[hyperloglog]") return static_cast(i / repeats); })); - estimator_type estimator{sd}; + ::cuda::stream stream{::cuda::device_ref{0}}; + auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + + estimator_type estimator{stream, mr, sd}; - REQUIRE(estimator.estimate() == 0); + REQUIRE(estimator.estimate(stream) == 0); // Add all items to the estimator - estimator.add(items_begin, items_begin + num_items); + estimator.add(stream, items_begin, items_begin + num_items); - const auto estimate = estimator.estimate(); + const auto estimate = estimator.estimate(stream); const double expected_count = static_cast(num_items) / static_cast(repeats); const double relative_error = std::abs((static_cast(estimate) / expected_count) - 1.0); @@ -210,10 +221,13 @@ C2H_TEST("HyperLogLog precision constructor", "[hyperloglog]") REQUIRE(estimator_type::sketch_bytes(precision) == expected_sketch_bytes); - estimator_type estimator{precision}; + ::cuda::stream stream{::cuda::device_ref{0}}; + auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + + estimator_type estimator{stream, mr, precision}; REQUIRE(estimator.sketch_bytes() == expected_sketch_bytes); - REQUIRE(estimator.estimate() == 0); + REQUIRE(estimator.estimate(stream) == 0); } #if _CCCL_CTK_AT_LEAST(12, 9) // Pinned memory resource is only supported with CTK 12.9 and later @@ -234,11 +248,14 @@ C2H_TEST("Hyperloglog estimate works with pinned memory pool", "[hyperloglog]") thrust::device_vector items(num_items); thrust::sequence(items.begin(), items.end(), T{0}); - estimator_type estimator{sketch_size_kb}; - estimator.add(items.begin(), items.end()); + ::cuda::stream stream{::cuda::device_ref{0}}; + auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + + estimator_type estimator{stream, mr, sketch_size_kb}; + estimator.add(stream, items.begin(), items.end()); auto host_mr = ::cuda::pinned_default_memory_pool(); - const auto estimate = estimator.estimate(host_mr); + const auto estimate = estimator.estimate(stream, host_mr); const double relative_error = std::abs((static_cast(estimate) / static_cast(num_items)) - 1.0); From 84e48f497534c0d4a6fc385f7b19a7cda5bba429 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 6 Jul 2026 22:00:44 +0000 Subject: [PATCH 2/5] Apply cudax cuco conventions to hyperloglog headers --- .../detail/hyperloglog/default_policy.cuh | 2 +- .../__cuco/detail/hyperloglog/finalizer.cuh | 2 +- .../detail/hyperloglog/hyperloglog_impl.cuh | 5 +- .../__cuco/detail/hyperloglog/kernels.cuh | 10 ++- .../cuda/experimental/__cuco/hyperloglog.cuh | 87 +++++++++++-------- .../experimental/__cuco/hyperloglog_ref.cuh | 20 ++--- 6 files changed, 70 insertions(+), 56 deletions(-) diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/default_policy.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/default_policy.cuh index 4a9eee4ad0b..e3d773a0c50 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/default_policy.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/default_policy.cuh @@ -74,7 +74,7 @@ struct default_hll_policy //! @return The hash value of `__k`. [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr hash_result_type hash(const _Key& __k) const noexcept { - return this->hasher_(__k); + return hasher_(__k); } //! @brief Extracts the register index from the hash. diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/finalizer.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/finalizer.cuh index 3ac8a900c38..53609adb868 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/finalizer.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/finalizer.cuh @@ -126,7 +126,7 @@ private: auto __high = ::cuda::std::min(__low + __k, __n); // Keep moving bounds as long as the (exclusive) high bound is closer to the estimate than // the lower (inclusive) bound. - while (__high < __n && this->__distance(__e, __high) < this->__distance(__e, __low)) + while (__high < __n && __distance(__e, __high) < __distance(__e, __low)) { __low += 1; __high += 1; diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh index c0adc92fd5b..7eb232fe1a7 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh @@ -42,7 +42,6 @@ #include #include #include -#include #include @@ -163,7 +162,7 @@ public: _CCCL_DEVICE_API constexpr void __add(const _Tp& __item) noexcept { const auto __h = __policy.hash(__item); - this->__update_max(__policy.register_index(__h, __precision), __policy.register_value(__h, __precision)); + __update_max(__policy.register_index(__h, __precision), __policy.register_value(__h, __precision)); } //! @brief Asynchronously adds to be counted items to the estimator. @@ -527,7 +526,7 @@ public: // implementation taken from // https://github.com/apache/spark/blob/6a27789ad7d59cd133653a49be0bb49729542abe/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/util/HyperLogLogPlusPlusHelper.scala#L43 - auto const __precision_from_sd = + const auto __precision_from_sd = static_cast<::cuda::std::int32_t>(::cuda::std::ceil(2.0 * ::cuda::std::log2(1.106 / __standard_deviation))); // minimum precision is 4 or 64 bytes diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh index aff1bd6a62b..d0bd93d3911 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh @@ -30,6 +30,8 @@ #include +#if _CCCL_CUDA_COMPILATION() + _CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_GCC("-Wattributes") @@ -97,7 +99,7 @@ __add_shmem_vectorized(const typename _RefType::__value_type* __first, ::cuda::s __idx += __loop_stride; } // a single thread processes the remaining items -#if _CCCL_CTK_AT_LEAST(12, 1) +# if _CCCL_CTK_AT_LEAST(12, 1) ::cooperative_groups::invoke_one(__grid, [&]() { const auto __remainder = __n % _VectorSize; for (int __i = 0; __i < __remainder; ++__i) @@ -105,7 +107,7 @@ __add_shmem_vectorized(const typename _RefType::__value_type* __first, ::cuda::s __local_ref.__add(*(__first + __n - __i - 1)); } }); -#else +# else if (__grid.thread_rank() == 0) { const auto __remainder = __n % _VectorSize; @@ -114,7 +116,7 @@ __add_shmem_vectorized(const typename _RefType::__value_type* __first, ::cuda::s __local_ref.__add(*(__first + __n - __i - 1)); } } -#endif +# endif __block.sync(); @@ -173,6 +175,8 @@ _CCCL_KERNEL_ATTRIBUTES void __merge(_OtherRefType __other_ref, _RefType __ref) _CCCL_DIAG_POP +#endif // _CCCL_CUDA_COMPILATION() + #include #endif // _CUDAX___CUCO_DETAIL_HYPERLOGLOG_KERNELS_CUH diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh index 68e1762fda5..de9f2acf833 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh @@ -21,6 +21,9 @@ # pragma system_header #endif // no system header +#include +#include +#include #include #include #include @@ -28,15 +31,15 @@ #include #include #include +#include #include #include -#include -#include -#include #include +#if !_CCCL_COMPILER(NVRTC) + namespace cuda::experimental::cuco { //! @brief A GPU-accelerated utility for approximating the number of distinct items in a multiset. @@ -107,10 +110,11 @@ public: //! //! @throw If sketch size implies precision outside [4, 18]. template - constexpr hyperloglog(::cuda::stream_ref __stream, - _MemoryResource_&& __memory_resource, - sketch_size_kb __sketch_size_kb = sketch_size_kb{32.0}, - const _Policy& __policy = {}) + _CCCL_HOST_API constexpr hyperloglog( + ::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, + sketch_size_kb __sketch_size_kb = sketch_size_kb{32.0}, + const _Policy& __policy = {}) : hyperloglog{__stream, ::cuda::std::forward<_MemoryResource_>(__memory_resource), __to_precision(__sketch_size_kb), @@ -128,10 +132,11 @@ public: //! //! @throw If standard deviation implies precision outside [4, 18]. template - constexpr hyperloglog(::cuda::stream_ref __stream, - _MemoryResource_&& __memory_resource, - standard_deviation __sd, - const _Policy& __policy = {}) + _CCCL_HOST_API constexpr hyperloglog( + ::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, + standard_deviation __sd, + const _Policy& __policy = {}) : hyperloglog{__stream, ::cuda::std::forward<_MemoryResource_>(__memory_resource), __to_precision(__sd), __policy} {} @@ -146,10 +151,11 @@ public: //! //! @throw If precision is outside [4, 18]. template - constexpr hyperloglog(::cuda::stream_ref __stream, - _MemoryResource_&& __memory_resource, - precision __precision, - const _Policy& __policy = {}) + _CCCL_HOST_API constexpr hyperloglog( + ::cuda::stream_ref __stream, + _MemoryResource_&& __memory_resource, + precision __precision, + const _Policy& __policy = {}) : __sketch_buffer{__stream, ::cuda::std::forward<_MemoryResource_>(__memory_resource), ref_type<>::sketch_bytes( @@ -176,7 +182,7 @@ public: //! @brief Asynchronously resets the estimator, i.e., clears the current count estimate. //! //! @param __stream CUDA stream this operation is executed in - constexpr void clear_async(::cuda::stream_ref __stream) noexcept + _CCCL_HOST_API constexpr void clear_async(::cuda::stream_ref __stream) noexcept { __ref.clear_async(__stream); } @@ -187,7 +193,7 @@ public: //! `clear_async`. //! //! @param __stream CUDA stream this operation is executed in - constexpr void clear(::cuda::stream_ref __stream) + _CCCL_HOST_API constexpr void clear(::cuda::stream_ref __stream) { __ref.clear(__stream); } @@ -202,7 +208,7 @@ public: //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items template - constexpr void add_async(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) + _CCCL_HOST_API constexpr void add_async(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { __ref.add_async(__stream, __first, __last); } @@ -220,7 +226,7 @@ public: //! @param __first Beginning of the sequence of items //! @param __last End of the sequence of items template - constexpr void add(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) + _CCCL_HOST_API constexpr void add(::cuda::stream_ref __stream, _InputIt __first, _InputIt __last) { __ref.add(__stream, __first, __last); } @@ -235,8 +241,8 @@ public: //! @param __stream CUDA stream this operation is executed in //! @param __other Other estimator to be merged into `*this` template <::cuda::thread_scope _OtherScope, class _OtherMemoryResource> - constexpr void merge_async(::cuda::stream_ref __stream, - const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) + _CCCL_HOST_API constexpr void + merge_async(::cuda::stream_ref __stream, const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) { __ref.merge_async(__stream, __other.__ref); } @@ -254,8 +260,8 @@ public: //! @param __stream CUDA stream this operation is executed in //! @param __other Other estimator to be merged into `*this` template <::cuda::thread_scope _OtherScope, class _OtherMemoryResource> - constexpr void merge(::cuda::stream_ref __stream, - const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) + _CCCL_HOST_API constexpr void + merge(::cuda::stream_ref __stream, const hyperloglog<_Tp, _OtherMemoryResource, _OtherScope, _Policy>& __other) { __ref.merge(__stream, __other.__ref); } @@ -269,7 +275,7 @@ public: //! @param __stream CUDA stream this operation is executed in //! @param __other_ref Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - constexpr void merge_async(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) + _CCCL_HOST_API constexpr void merge_async(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) { __ref.merge_async(__stream, __other_ref); } @@ -286,7 +292,7 @@ public: //! @param __stream CUDA stream this operation is executed in //! @param __other_ref Other estimator reference to be merged into `*this` template <::cuda::thread_scope _OtherScope> - constexpr void merge(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) + _CCCL_HOST_API constexpr void merge(::cuda::stream_ref __stream, const ref_type<_OtherScope>& __other_ref) { __ref.merge(__stream, __other_ref); } @@ -303,7 +309,7 @@ public: //! //! @return Approximate distinct items count template - [[nodiscard]] constexpr ::cuda::std::size_t + [[nodiscard]] _CCCL_HOST_API constexpr ::cuda::std::size_t estimate(::cuda::stream_ref __stream, _HostMemoryResource __host_mr = {}) const { return __ref.estimate(__stream, __host_mr); @@ -312,7 +318,7 @@ public: //! @brief Get device ref. //! //! @return Device ref object of the current `hyperloglog` host object - [[nodiscard]] constexpr ref_type<> ref() const noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr ref_type<> ref() const noexcept { return {sketch(), policy()}; } @@ -320,7 +326,7 @@ public: //! @brief Get hash function. //! //! @return The hash function - [[nodiscard]] constexpr auto hash_function() const noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto hash_function() const noexcept { return __ref.hash_function(); } @@ -328,7 +334,7 @@ public: //! @brief Get the policy. //! //! @return The policy - [[nodiscard]] constexpr const _Policy& policy() const noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr const _Policy& policy() const noexcept { return __ref.policy(); } @@ -336,7 +342,7 @@ public: //! @brief Gets the span of the sketch. //! //! @return The ::cuda::std::span of the sketch - [[nodiscard]] constexpr ::cuda::std::span<::cuda::std::byte> sketch() const noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr ::cuda::std::span<::cuda::std::byte> sketch() const noexcept { return __ref.sketch(); } @@ -344,7 +350,7 @@ public: //! @brief Gets the number of bytes required for the sketch storage. //! //! @return The number of bytes required for the sketch - [[nodiscard]] constexpr ::cuda::std::size_t sketch_bytes() const noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr ::cuda::std::size_t sketch_bytes() const noexcept { return __ref.sketch_bytes(); } @@ -354,7 +360,8 @@ public: //! @param __sketch_size_kb Upper bound sketch size in KB //! //! @return The number of bytes required for the sketch - [[nodiscard]] static constexpr ::cuda::std::size_t sketch_bytes(sketch_size_kb __sketch_size_kb) noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr ::cuda::std::size_t + sketch_bytes(sketch_size_kb __sketch_size_kb) noexcept { return ref_type<>::sketch_bytes(__sketch_size_kb); } @@ -364,7 +371,8 @@ public: //! @param __standard_deviation Upper bound standard deviation for approximation error //! //! @return The number of bytes required for the sketch - [[nodiscard]] static constexpr ::cuda::std::size_t sketch_bytes(standard_deviation __standard_deviation) noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr ::cuda::std::size_t + sketch_bytes(standard_deviation __standard_deviation) noexcept { return ref_type<>::sketch_bytes(__standard_deviation); } @@ -374,7 +382,7 @@ public: //! @param __precision HyperLogLog precision parameter //! //! @return The number of bytes required for the sketch - [[nodiscard]] static constexpr ::cuda::std::size_t sketch_bytes(precision __precision) noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr ::cuda::std::size_t sketch_bytes(precision __precision) noexcept { return ref_type<>::sketch_bytes(__precision); } @@ -382,13 +390,14 @@ public: //! @brief Gets the alignment required for the sketch storage. //! //! @return The required alignment - [[nodiscard]] static constexpr ::cuda::std::size_t sketch_alignment() noexcept + [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr ::cuda::std::size_t sketch_alignment() noexcept { return ref_type<>::sketch_alignment(); } private: - [[nodiscard]] static constexpr precision __precision_in_bounds(precision __precision, const char* __message) + [[nodiscard]] _CCCL_HOST_API static constexpr precision + __precision_in_bounds(precision __precision, const char* __message) { const auto __value = static_cast<::cuda::std::int32_t>(__precision); const auto __in_range = ::cuda::in_range(__value, 4, 18); @@ -399,7 +408,7 @@ private: return __precision; } - [[nodiscard]] static constexpr precision __to_precision(sketch_size_kb __sketch_size_kb) + [[nodiscard]] _CCCL_HOST_API static constexpr precision __to_precision(sketch_size_kb __sketch_size_kb) { const auto __bytes = ref_type<>::sketch_bytes(__sketch_size_kb) / sizeof(register_type); const auto __precision = static_cast(::cuda::std::countr_zero(static_cast<::cuda::std::size_t>(__bytes))); @@ -407,7 +416,7 @@ private: precision{__precision}, "HyperLogLog sketch size must be in range [0.0625 KB, 1024 KB]"); } - [[nodiscard]] static constexpr precision __to_precision(standard_deviation __standard_deviation) + [[nodiscard]] _CCCL_HOST_API static constexpr precision __to_precision(standard_deviation __standard_deviation) { const auto __bytes = ref_type<>::sketch_bytes(__standard_deviation) / sizeof(register_type); const auto __precision = static_cast(::cuda::std::countr_zero(static_cast<::cuda::std::size_t>(__bytes))); @@ -417,6 +426,8 @@ private: }; } // namespace cuda::experimental::cuco +#endif // !_CCCL_COMPILER(NVRTC) + #include #endif // _CUDAX___CUCO_HYPERLOGLOG_CUH diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh index b5769186918..732dffb391b 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh @@ -21,11 +21,11 @@ # pragma system_header #endif // no system header +#include +#include #include -#include #include #include -#include #include #include @@ -108,11 +108,11 @@ public: //! @tparam _CG CUDA Cooperative Group type //! //! @param __group CUDA Cooperative group this operation is executed in - template - _CCCL_DEVICE_API constexpr ::cuda::std::enable_if_t> - clear(_CG __group) noexcept + _CCCL_TEMPLATE(class _CG) + _CCCL_REQUIRES((!::cuda::std::is_convertible_v<_CG, ::cuda::stream_ref>) ) + _CCCL_DEVICE_API constexpr void clear(_CG __group) noexcept { - // The enable_if above is to work around an incompatibility between host and device + // The constraint above is to work around an incompatibility between host and device // overload preference for clang and NVCC. See // https://llvm.org/docs/CompileCudaWithLLVM.html#overloading-based-on-host-and-device-attributes // for further reading, but the bottom line is when: @@ -213,11 +213,11 @@ public: //! //! @param __group CUDA Cooperative group this operation is executed in //! @param __other Other estimator reference to be merged into `*this` - template - _CCCL_DEVICE_API constexpr ::cuda::std::enable_if_t> - merge(_CG __group, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) + _CCCL_TEMPLATE(class _CG, ::cuda::thread_scope _OtherScope) + _CCCL_REQUIRES((!::cuda::std::is_convertible_v<_CG, ::cuda::stream_ref>) ) + _CCCL_DEVICE_API constexpr void merge(_CG __group, const hyperloglog_ref<_Tp, _OtherScope, _Policy>& __other) { - // The enable_if above works around the same host/device overload preference issue as + // The constraint above works around the same host/device overload preference issue as // documented in `clear(_CG)`: `merge(_CG, ...)` would otherwise conflict with // `merge(::cuda::stream_ref, ...)` when called from `hyperloglog::merge(::cuda::stream_ref, ...)`. __impl.__merge(__group, __other.__impl); From 653cbea50d53da5357b35c5851cb733b97c4c543 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 7 Jul 2026 18:25:15 +0000 Subject: [PATCH 3/5] Address review feedback: stream-ordered ctor docs, granular no_init include, buffer-based tests --- .../cuda/experimental/__cuco/hyperloglog.cuh | 10 ++- .../experimental/__cuco/hyperloglog_ref.cuh | 1 - .../test/cuco/hyperloglog/test_hyperloglog.cu | 62 +++++++++++-------- 3 files changed, 43 insertions(+), 30 deletions(-) diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh index de9f2acf833..1d02a885bfd 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -101,7 +102,8 @@ public: // TODO enable CTAD //! @brief Constructs a `hyperloglog` host object. //! - //! @note This function synchronizes the given stream. + //! @note Construction is stream-ordered: the initial clear is enqueued on `__stream` without + //! synchronizing it. //! //! @param __stream CUDA stream used to initialize the object //! @param __memory_resource A memory resource used for allocating device storage @@ -123,7 +125,8 @@ public: //! @brief Constructs a `hyperloglog` host object. //! - //! @note This function synchronizes the given stream. + //! @note Construction is stream-ordered: the initial clear is enqueued on `__stream` without + //! synchronizing it. //! //! @param __stream CUDA stream used to initialize the object //! @param __memory_resource A memory resource used for allocating device storage @@ -142,7 +145,8 @@ public: //! @brief Constructs a `hyperloglog` host object. //! - //! @note This function synchronizes the given stream. + //! @note Construction is stream-ordered: the initial clear is enqueued on `__stream` without + //! synchronizing it. //! //! @param __stream CUDA stream used to initialize the object //! @param __memory_resource A memory resource used for allocating device storage diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh index 732dffb391b..5e1a9afae7d 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog_ref.cuh @@ -228,7 +228,6 @@ public: //! //! @throw If sketch_bytes() != __other.sketch_bytes() //! - // Review: For host-side merges, consider validating `__other` is on the same device. //! @tparam _OtherScope Thread scope of `other` estimator //! //! @param __stream CUDA stream this operation is executed in diff --git a/cudax/test/cuco/hyperloglog/test_hyperloglog.cu b/cudax/test/cuco/hyperloglog/test_hyperloglog.cu index f45cd5e4fea..2879f73df89 100644 --- a/cudax/test/cuco/hyperloglog/test_hyperloglog.cu +++ b/cudax/test/cuco/hyperloglog/test_hyperloglog.cu @@ -8,10 +8,11 @@ // //===----------------------------------------------------------------------===// -#include +#include #include -#include +#include +#include #include #include #include @@ -58,6 +59,17 @@ __global__ void estimate_kernel(typename Ref::sketch_size_kb sketch_size_kb, Inp using test_types = c2h::type_list; +// Maps index i to i / repeats, yielding `repeats` duplicates of each value +struct scaled_index +{ + std::size_t repeats; + + __device__ int operator()(std::size_t i) const noexcept + { + return static_cast(i / repeats); + } +}; + C2H_TEST("HyperLogLog device ref", "[hyperloglog]", test_types) { using T = c2h::get<0, TestType>; @@ -71,14 +83,13 @@ C2H_TEST("HyperLogLog device ref", "[hyperloglog]", test_types) CAPTURE(num_items, hll_precision, sketch_size_kb); - thrust::device_vector items(num_items); - - // Generate `num_items` distinct items - thrust::sequence(items.begin(), items.end(), T{0}); - ::cuda::stream stream{::cuda::device_ref{0}}; auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + // Generate `num_items` distinct items + auto items = ::cuda::make_buffer(stream, mr, num_items, ::cuda::no_init); + thrust::sequence(thrust::cuda::par_nosync.on(stream.get()), items.begin(), items.end(), T{0}); + // Initialize the estimator estimator_type estimator{stream, mr, sketch_size_kb}; @@ -87,13 +98,16 @@ C2H_TEST("HyperLogLog device ref", "[hyperloglog]", test_types) const auto host_estimate = estimator.estimate(stream); - thrust::device_vector device_estimate(1); + auto device_estimate = ::cuda::make_buffer(stream, mr, 1, ::cuda::no_init); estimate_kernel> - <<<1, 512, estimator.sketch_bytes()>>>(sketch_size_kb, items.begin(), num_items, device_estimate.begin()); - - REQUIRE_CUDART(cudaDeviceSynchronize()); - - std::size_t device_estimate_value = device_estimate[0]; + <<<1, 512, estimator.sketch_bytes(), stream.get()>>>( + sketch_size_kb, items.begin(), num_items, device_estimate.begin()); + REQUIRE(cudaGetLastError() == cudaSuccess); + + std::size_t device_estimate_value{}; + REQUIRE_CUDART(cudaMemcpyAsync( + &device_estimate_value, device_estimate.data(), sizeof(std::size_t), cudaMemcpyDeviceToHost, stream.get())); + stream.sync(); REQUIRE(device_estimate_value == host_estimate); } @@ -114,14 +128,13 @@ C2H_TEST("HyperLogLog unique sequence", "[hyperloglog]", test_types) // RSD for a given precision is given by the following formula const double relative_standard_deviation = 1.04 / std::sqrt(static_cast(1ull << hll_precision)); - thrust::device_vector items(num_items); - - // Generate `num_items` distinct items - thrust::sequence(items.begin(), items.end(), T{0}); - ::cuda::stream stream{::cuda::device_ref{0}}; auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + // Generate `num_items` distinct items + auto items = ::cuda::make_buffer(stream, mr, num_items, ::cuda::no_init); + thrust::sequence(thrust::cuda::par_nosync.on(stream.get()), items.begin(), items.end(), T{0}); + // Initialize the estimator estimator_type estimator{stream, mr, sketch_size_kb}; @@ -137,7 +150,7 @@ C2H_TEST("HyperLogLog unique sequence", "[hyperloglog]", test_types) REQUIRE(estimator.estimate(stream) == estimate); // Adding the same items again (might use shared memory code path) should not affect the result - auto* ptr = thrust::raw_pointer_cast(items.data()); + auto* ptr = items.data(); estimator.add(stream, ptr, ptr + num_items / 2); REQUIRE(estimator.estimate(stream) == estimate); @@ -181,10 +194,7 @@ C2H_TEST("HyperLogLog Spark parity deterministic", "[hyperloglog]") REQUIRE(estimator_type::sketch_bytes(sd) == expected_sketch_bytes); REQUIRE(estimator_type::sketch_bytes(sd) == estimator_type::sketch_bytes(sb)); - auto items_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), cuda::proclaim_return_type([repeats] __device__(auto i) { - return static_cast(i / repeats); - })); + auto items_begin = cuda::transform_iterator(cuda::counting_iterator{0}, scaled_index{repeats}); ::cuda::stream stream{::cuda::device_ref{0}}; auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); @@ -245,12 +255,12 @@ C2H_TEST("Hyperloglog estimate works with pinned memory pool", "[hyperloglog]") constexpr double tolerance_factor = 2.5; const double relative_standard_deviation = 1.04 / std::sqrt(static_cast(1ull << hll_precision)); - thrust::device_vector items(num_items); - thrust::sequence(items.begin(), items.end(), T{0}); - ::cuda::stream stream{::cuda::device_ref{0}}; auto mr = ::cuda::device_default_memory_pool(::cuda::device_ref{0}); + auto items = ::cuda::make_buffer(stream, mr, num_items, ::cuda::no_init); + thrust::sequence(thrust::cuda::par_nosync.on(stream.get()), items.begin(), items.end(), T{0}); + estimator_type estimator{stream, mr, sketch_size_kb}; estimator.add(stream, items.begin(), items.end()); From 03f847456b4fd705c2466ab874b08c228ac860cc Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 8 Jul 2026 17:07:20 +0000 Subject: [PATCH 4/5] Add #else/#endif comments and hide defaulted special members from ABI --- .../experimental/__cuco/detail/hyperloglog/kernels.cuh | 4 ++-- cudax/include/cuda/experimental/__cuco/hyperloglog.cuh | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh index d0bd93d3911..a33e2f47455 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh @@ -107,7 +107,7 @@ __add_shmem_vectorized(const typename _RefType::__value_type* __first, ::cuda::s __local_ref.__add(*(__first + __n - __i - 1)); } }); -# else +# else // ^^^ _CCCL_CTK_AT_LEAST(12, 1) ^^^ / vvv _CCCL_CTK_BELOW(12, 1) vvv if (__grid.thread_rank() == 0) { const auto __remainder = __n % _VectorSize; @@ -116,7 +116,7 @@ __add_shmem_vectorized(const typename _RefType::__value_type* __first, ::cuda::s __local_ref.__add(*(__first + __n - __i - 1)); } } -# endif +# endif // ^^^ _CCCL_CTK_BELOW(12, 1) ^^^ __block.sync(); diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh index 1d02a885bfd..8f3dafbd626 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh @@ -172,16 +172,16 @@ public: clear_async(__stream); } - ~hyperloglog() = default; + _CCCL_HIDE_FROM_ABI ~hyperloglog() = default; hyperloglog(const hyperloglog&) = delete; //! @brief Copy-assignment operator. //! //! @return Copy of `*this` - hyperloglog& operator=(const hyperloglog&) = delete; - hyperloglog(hyperloglog&&) = default; ///< Move constructor + hyperloglog& operator=(const hyperloglog&) = delete; + _CCCL_HIDE_FROM_ABI hyperloglog(hyperloglog&&) = default; ///< Move constructor - hyperloglog& operator=(hyperloglog&&) = default; + _CCCL_HIDE_FROM_ABI hyperloglog& operator=(hyperloglog&&) = default; //! @brief Asynchronously resets the estimator, i.e., clears the current count estimate. //! From 412bbd977844af14bf41fa119d40ccd6d4fa3aba Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 8 Jul 2026 19:25:55 +0000 Subject: [PATCH 5/5] Use to_address instead of subscripting the iterator in host code --- .../__cuco/detail/hyperloglog/hyperloglog_impl.cuh | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh index 7eb232fe1a7..a14a698571f 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh @@ -34,7 +34,6 @@ #include #include #include -#include #include #include @@ -231,7 +230,7 @@ public: __kernel, __shmem_bytes); - const auto __ptr = ::cuda::std::addressof(__first[0]); + const auto __ptr = ::cuda::std::to_address(__first); void* __kernel_args[] = {const_cast(reinterpret_cast(&__ptr)), const_cast(reinterpret_cast(&__num_items)), reinterpret_cast(this)};