Skip to content

Commit

Permalink
Merge remote-tracking branch 'upstream/dev' into bloom-filter-new
Browse files Browse the repository at this point in the history
  • Loading branch information
sleeepyjack committed Oct 1, 2024
2 parents 70de880 + 760ade8 commit d11c050
Show file tree
Hide file tree
Showing 18 changed files with 348 additions and 693 deletions.
8 changes: 4 additions & 4 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -234,13 +234,13 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
#### Examples:
- [Host-bulk APIs (TODO)]()

### `distinct_count_estimator`
### `hyperloglog`

`cuco::distinct_count_estimator` implements the well-established [HyperLogLog++ algorithm](https://static.googleusercontent.com/media/research.google.com/de//pubs/archive/40671.pdf) for approximating the count of distinct items in a multiset/stream.
`cuco::hyperloglog` implements the well-established [HyperLogLog++ algorithm](https://static.googleusercontent.com/media/research.google.com/de//pubs/archive/40671.pdf) for approximating the count of distinct items in a multiset/stream.

#### Examples:
- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/distinct_count_estimator/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/z/sMfofM6qd))
- [Device-ref APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/distinct_count_estimator/device_ref_example.cu) (see [live example in godbolt](https://godbolt.org/z/156T9ox7h))
- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/hyperloglog/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/z/G4qdcTezE))
- [Device-ref APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/hyperloglog/device_ref_example.cu) (see [live example in godbolt](https://godbolt.org/z/n88713o4n))

### `bloom_filter`

Expand Down
6 changes: 3 additions & 3 deletions benchmarks/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -95,9 +95,9 @@ ConfigureBench(HASH_FUNCTION_BENCH
hash_function/hash_function_bench.cu)

###################################################################################################
# - distinct_count_estimator benchmarks -----------------------------------------------------------
ConfigureBench(DISTINCT_COUNT_ESTIMATOR_BENCH
distinct_count_estimator/distinct_count_estimator_bench.cu)
# - hyperloglog benchmarks -----------------------------------------------------------
ConfigureBench(HYPERLOGLOG_BENCH
hyperloglog/hyperloglog_bench.cu)

###################################################################################################
# - bloom_filter benchmarks -----------------------------------------------------------------------
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#include <benchmark_defaults.hpp>
#include <benchmark_utils.hpp>

#include <cuco/distinct_count_estimator.cuh>
#include <cuco/hyperloglog.cuh>
#include <cuco/static_set.cuh>
#include <cuco/utility/key_generator.cuh>

Expand Down Expand Up @@ -74,12 +74,12 @@ template <class Estimator, class Dist>
}

/**
* @brief A benchmark evaluating `cuco::distinct_count_estimator` end-to-end performance
* @brief A benchmark evaluating `cuco::hyperloglog` end-to-end performance
*/
template <typename T, typename Dist>
void distinct_count_estimator_e2e(nvbench::state& state, nvbench::type_list<T, Dist>)
void hyperloglog_e2e(nvbench::state& state, nvbench::type_list<T, Dist>)
{
using estimator_type = cuco::distinct_count_estimator<T>;
using estimator_type = cuco::hyperloglog<T>;

auto const num_items = state.get_int64("NumInputs");
auto const sketch_size_kb = state.get_int64("SketchSizeKB");
Expand Down Expand Up @@ -114,12 +114,12 @@ void distinct_count_estimator_e2e(nvbench::state& state, nvbench::type_list<T, D
}

/**
* @brief A benchmark evaluating `cuco::distinct_count_estimator::add_async` performance
* @brief A benchmark evaluating `cuco::hyperloglog::add_async` performance
*/
template <typename T, typename Dist>
void distinct_count_estimator_add(nvbench::state& state, nvbench::type_list<T, Dist>)
void hyperloglog_add(nvbench::state& state, nvbench::type_list<T, Dist>)
{
using estimator_type = cuco::distinct_count_estimator<T>;
using estimator_type = cuco::hyperloglog<T>;

auto const num_items = state.get_int64("NumInputs");
auto const sketch_size_kb = state.get_int64("SketchSizeKB");
Expand All @@ -144,18 +144,18 @@ void distinct_count_estimator_add(nvbench::state& state, nvbench::type_list<T, D

using TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t, __int128_t>;

NVBENCH_BENCH_TYPES(distinct_count_estimator_e2e,
NVBENCH_BENCH_TYPES(hyperloglog_e2e,
NVBENCH_TYPE_AXES(TYPE_RANGE, nvbench::type_list<distribution::uniform>))
.set_name("distinct_count_estimator_e2e_uniform")
.set_name("hyperloglog_e2e_uniform")
.set_type_axes_names({"T", "Distribution"})
.add_int64_power_of_two_axis("NumInputs", {28, 29, 30})
.add_int64_axis("SketchSizeKB", {8, 16, 32, 64, 128, 256}) // 256KB uses gmem fallback kernel
.add_int64_axis("Multiplicity", {1})
.set_max_noise(defaults::MAX_NOISE);

NVBENCH_BENCH_TYPES(distinct_count_estimator_add,
NVBENCH_BENCH_TYPES(hyperloglog_add,
NVBENCH_TYPE_AXES(TYPE_RANGE, nvbench::type_list<distribution::uniform>))
.set_name("distinct_count_estimator_add_uniform")
.set_name("hyperloglog_add_uniform")
.set_type_axes_names({"T", "Distribution"})
.add_int64_power_of_two_axis("NumInputs", {28, 29, 30})
.add_int64_axis("SketchSizeKB", {8, 16, 32, 64, 128, 256})
Expand Down
4 changes: 2 additions & 2 deletions examples/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,6 @@ 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")
ConfigureExample(DISTINCT_COUNT_ESTIMATOR_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/distinct_count_estimator/device_ref_example.cu")
ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/host_bulk_example.cu")
ConfigureExample(HYPERLOGLOG_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/device_ref_example.cu")
ConfigureExample(BLOOM_FILTER_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/host_bulk_example.cu")
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <cuco/distinct_count_estimator.cuh>
#include <cuco/hyperloglog.cuh>

#include <cuda/std/cstddef>
#include <thrust/device_vector.h>
Expand All @@ -23,9 +23,9 @@

/**
* @file device_ref_example.cu
* @brief Demonstrates usage of `cuco::distinct_count_estimator` device-side APIs.
* @brief Demonstrates usage of `cuco::hyperloglog` device-side APIs.
*
* This example demonstrates how the non-owning reference type `cuco::distinct_count_estimator_ref`
* This example demonstrates how the non-owning reference type `cuco::hyperloglog_ref`
* can be used to implement a custom kernel that fuses the cardinality estimation step with any
* other workload that traverses the input data.
*/
Expand Down Expand Up @@ -119,7 +119,7 @@ __global__ void device_estimate_kernel(cuco::sketch_size_kb sketch_size_kb,
int main(void)
{
using T = int;
using estimator_type = cuco::distinct_count_estimator<T>;
using estimator_type = cuco::hyperloglog<T>;
constexpr std::size_t num_items = 1ull << 28; // 1GB
auto const sketch_size_kb = 32_KB;

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <cuco/distinct_count_estimator.cuh>
#include <cuco/hyperloglog.cuh>

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
Expand All @@ -24,7 +24,7 @@

/**
* @file host_bulk_example.cu
* @brief Demonstrates usage of `cuco::distinct_count_estimator` "bulk" host APIs.
* @brief Demonstrates usage of `cuco::hyperloglog` "bulk" host APIs.
*/
int main(void)
{
Expand All @@ -41,7 +41,7 @@ int main(void)
auto const sd = cuco::standard_deviation{0.0122197};

// Initialize the estimator
cuco::distinct_count_estimator<T> estimator{sd};
cuco::hyperloglog<T> estimator{sd};

// Add all items to the estimator
estimator.add(items.begin(), items.end());
Expand Down

This file was deleted.

Loading

0 comments on commit d11c050

Please sign in to comment.