distinct_count_estimator aka HyperLogLog++

examples/CMakeLists.txt

@@ -41,3 +41,4 @@ ConfigureExample(STATIC_MAP_DEVICE_SIDE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/sta
 ConfigureExample(STATIC_MAP_CUSTOM_TYPE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/custom_type_example.cu")
 ConfigureExample(STATIC_MAP_COUNT_BY_KEY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/count_by_key_example.cu")
 ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multimap/host_bulk_example.cu")
+ConfigureExample(DISTINCT_COUNT_ESTIMATOR_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/distinct_count_estimator/host_bulk_example.cu")

examples/distinct_count_estimator/host_bulk_example.cu

@@ -0,0 +1,61 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include <cuco/distinct_count_estimator.cuh>
+
+#include <thrust/device_vector.h>
+#include <thrust/sequence.h>
+
+#include <cstddef>
+#include <iostream>
+
+int main()
+{
+  using T                         = int;
+  std::size_t constexpr num_items = 1ull << 30;  // 4GB
+
+  thrust::device_vector<T> items(num_items);
+  // create a vector of distinct items
+  thrust::sequence(items.begin(), items.end(), 0);
+
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+
+  cuco::distinct_count_estimator<T> estimator;
+  cudaEventRecord(start);
+  // add all items to the estimator
+  estimator.add(items.begin(), items.end());
+  // after the estimator has seen all items, we can calculate the cardinality
+  std::size_t const estimated_cardinality = estimator.estimate();
+  cudaEventRecord(stop);
+  cudaEventSynchronize(stop);
+
+  float milliseconds = 0;
+  cudaEventElapsedTime(&milliseconds, start, stop);
+  float input_size_gb = num_items * sizeof(T) / 1073741824.0f;
+  float throughput    = input_size_gb / (milliseconds / 1000.0f);
+
+  std::cout << "True cardinality:\t" << num_items << "\nEstimated cardinality:\t"
+            << estimated_cardinality << "\nRelative error:\t"
+            << abs(static_cast<double>(num_items) - static_cast<double>(estimated_cardinality)) /
+                 num_items
+            << "\nData size:\t" << input_size_gb << "GB"
+            << "\nElapsed time:\t" << milliseconds << "ms"
+            << "\nMemory throughput\t" << throughput << "GB/s" << std::endl;
+
+  cudaEventDestroy(start);
+  cudaEventDestroy(stop);
+}

include/cuco/detail/distinct_count_estimator/distinct_count_estimator.inl

@@ -0,0 +1,106 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+namespace cuco {
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+constexpr distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::distinct_count_estimator(
+  cuco::cuda_thread_scope<Scope> scope,
+  Hash const& hash,
+  Allocator const& alloc,
+  cuco::cuda_stream_ref stream)
+  : impl_{std::make_unique<impl_type>(scope, hash, alloc, stream)}
+{
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::clear_async(
+  cuco::cuda_stream_ref stream) noexcept
+{
+  this->impl_->clear_async(stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::clear(
+  cuco::cuda_stream_ref stream)
+{
+  this->impl_->clear(stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <class InputIt>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::add_async(
+  InputIt first, InputIt last, cuco::cuda_stream_ref stream) noexcept
+{
+  this->impl_->add_async(first, last, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <class InputIt>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::add(
+  InputIt first, InputIt last, cuco::cuda_stream_ref stream)
+{
+  this->impl_->add(first, last, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <cuda::thread_scope OtherScope, class OtherAllocator>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::merge_async(
+  distinct_count_estimator<T, Precision, OtherScope, Hash, OtherAllocator> const& other,
+  cuco::cuda_stream_ref stream) noexcept
+{
+  this->impl_->merge_async(other, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <cuda::thread_scope OtherScope, class OtherAllocator>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::merge(
+  distinct_count_estimator<T, Precision, OtherScope, Hash, OtherAllocator> const& other,
+  cuco::cuda_stream_ref stream)
+{
+  this->impl_->merge(other, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <cuda::thread_scope OtherScope>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::merge_async(
+  ref_type<OtherScope> const& other, cuco::cuda_stream_ref stream) noexcept
+{
+  this->impl_->merge_async(other, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+template <cuda::thread_scope OtherScope>
+void distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::merge(
+  ref_type<OtherScope> const& other, cuco::cuda_stream_ref stream)
+{
+  this->impl_->merge(other, stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+std::size_t distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::estimate(
+  cuco::cuda_stream_ref stream) const
+{
+  return this->impl_->estimate(stream);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash, class Allocator>
+typename distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::ref_type<>
+distinct_count_estimator<T, Precision, Scope, Hash, Allocator>::ref() const noexcept
+{
+  return this->impl_->ref();
+}
+}  // namespace cuco

include/cuco/detail/distinct_count_estimator/distinct_count_estimator_ref.inl

@@ -0,0 +1,57 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+namespace cuco {
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash>
+__host__ __device__ constexpr distinct_count_estimator_ref<T, Precision, Scope, Hash>::
+  distinct_count_estimator_ref(storage_type& storage,
+                               cuco::cuda_thread_scope<Scope> scope,
+                               Hash const& hash) noexcept
+  : impl_{storage, scope, hash}
+{
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash>
+template <class CG>
+__device__ void distinct_count_estimator_ref<T, Precision, Scope, Hash>::clear(
+  CG const& group) noexcept
+{
+  this->impl_.clear(group);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash>
+__device__ void distinct_count_estimator_ref<T, Precision, Scope, Hash>::add(T const& item) noexcept
+{
+  this->impl_.add(item);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash>
+template <class CG, cuda::thread_scope OtherScope>
+__device__ void distinct_count_estimator_ref<T, Precision, Scope, Hash>::merge(
+  CG const& group,
+  distinct_count_estimator_ref<T, Precision, OtherScope, Hash> const& other) noexcept
+{
+  this->impl_.merge(group, other);
+}
+
+template <class T, int32_t Precision, cuda::thread_scope Scope, class Hash>
+__device__ std::size_t distinct_count_estimator_ref<T, Precision, Scope, Hash>::estimate(
+  cooperative_groups::thread_block const& group) const noexcept
+{
+  this->impl_.estimate(group);
+}
+}  // namespace cuco

include/cuco/detail/hyperloglog/finalizer.cuh

@@ -0,0 +1,79 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
+
+#include <cuco/detail/hyperloglog/tuning.cuh>
+
+#include <cuda/std/cmath>
+
+namespace cuco::hyperloglog_ns::detail {
+template <int32_t Precision>
+class finalizer {
+  // this minimum number of registers is required by HLL++
+  static_assert(Precision >= 4, "Precision must be greater or equal to 4");
+
+ public:
+  __host__ __device__ static double constexpr finalize(double z, int v) noexcept
+  {
+    auto e = alpha_mm() / z;
+    // TODO remove test code
+    // printf("raw e: %lf\n", e);
+
+    if (v > 0) {
+      // Use linear counting for small cardinality estimates.
+      double const h = m * log(static_cast<double>(m) / v);
+      // HLL++ is defined only when p < 19, otherwise we need to fallback to HLL.
+      // The threshold `2.5 * m` is from the original HLL algorithm.
+      if ((Precision < 19 and h <= thresholds[Precision - 4]) or e <= 2.5 * m) {
+        e = h;
+      } else {
+        e = bias_corrected_estimate(e);
+      }
+    } else {
+      e = bias_corrected_estimate(e);
+    }
+
+    return cuda::std::round(e);
+  }
+
+ private:
+  static auto constexpr m = (1 << Precision);
+
+  __host__ __device__ static double constexpr alpha_mm() noexcept
+  {
+    if constexpr (m == 16) {
+      return 0.673 * m * m;
+    } else if constexpr (m == 32) {
+      return 0.697 * m * m;
+    } else if constexpr (m == 64) {
+      return 0.709 * m * m;
+    } else {
+      return (0.7213 / (1.0 + 1.079 / m)) * m * m;
+    }
+  }
+
+  __host__ __device__ static double constexpr bias_corrected_estimate(double e) noexcept
+  {
+    if constexpr (Precision < 19) {
+      if (e < 5.0 * m) { return e - bias(e); }
+    }
+    return e;
+  }
+
+  // TODO implement HLL++ bias correction
+  __host__ __device__ static double constexpr bias(double e) noexcept { return e * 0; }
+};
+}  // namespace cuco::hyperloglog_ns::detail