Skip to content

Make stream and memory resource explicit in hyperloglog APIs #9720

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
PointKernel wants to merge 8 commits into
base: main
Choose a base branch
from
Open

Make stream and memory resource explicit in hyperloglog APIs #9720

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
ae3b6ea
Make stream and memory resource explicit in hyperloglog APIs
PointKernel Jul 6, 2026
84e48f4
Apply cudax cuco conventions to hyperloglog headers
PointKernel Jul 6, 2026
653cbea
Address review feedback: stream-ordered ctor docs, granular no_init i…
PointKernel Jul 7, 2026
e981a78
Merge remote-tracking branch 'upstream/main' into hll-explicit-stream-mr
PointKernel Jul 7, 2026
03f8474
Add #else/#endif comments and hide defaulted special members from ABI
PointKernel Jul 8, 2026
669032e
Merge remote-tracking branch 'upstream/main' into hll-explicit-stream-mr
PointKernel Jul 8, 2026
412bbd9
Use to_address instead of subscripting the iterator in host code
PointKernel Jul 8, 2026
333eb34
Merge remote-tracking branch 'upstream/main' into hll-explicit-stream-mr
PointKernel Jul 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion cudax/include/cuda/experimental/__cuco/detail/hyperloglog/default_policy.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
2 changes: 1 addition & 1 deletion cudax/include/cuda/experimental/__cuco/detail/hyperloglog/finalizer.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
5 changes: 2 additions & 3 deletions cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,6 @@
#include <cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh>
#include <cuda/experimental/__cuco/detail/utility/strong_type.cuh>
#include <cuda/experimental/__cuco/hash_functions.cuh>
#include <cuda/experimental/memory_resource.cuh>

#include <cooperative_groups.h>

Expand Down Expand Up @@ -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.
Expand Down Expand Up @@ -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
Expand Down
10 changes: 7 additions & 3 deletions cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,8 @@

#include <cuda/std/__cccl/prologue.h>

#if _CCCL_CUDA_COMPILATION()

_CCCL_DIAG_PUSH
_CCCL_DIAG_SUPPRESS_GCC("-Wattributes")

Expand Down Expand Up @@ -97,15 +99,15 @@ __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)
{
__local_ref.__add(*(__first + __n - __i - 1));
}
});
#else
# else
Comment thread
PointKernel marked this conversation as resolved.
Outdated
if (__grid.thread_rank() == 0)
{
const auto __remainder = __n % _VectorSize;
Expand All @@ -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();

Expand Down Expand Up @@ -173,6 +175,8 @@ _CCCL_KERNEL_ATTRIBUTES void __merge(_OtherRefType __other_ref, _RefType __ref)

_CCCL_DIAG_POP

#endif // _CCCL_CUDA_COMPILATION()

#include <cuda/std/__cccl/epilogue.h>

#endif // _CUDAX___CUCO_DETAIL_HYPERLOGLOG_KERNELS_CUH
Loading
Loading