From 90cc686427e04b8d7a55768d9f9b3d4a3371fd3c Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 25 Oct 2024 01:11:22 +0000 Subject: [PATCH 01/29] WIP. Add Arrow BF policy --- examples/CMakeLists.txt | 1 + .../arrow_bf_host_bulk_example.cu | 76 +++++++++++ include/cuco/arrow_bf_policy.cuh | 123 ++++++++++++++++++ .../detail/bloom_filter/arrow_bf_policy.inl | 52 ++++++++ .../bloom_filter/arrow_bf_policy_impl.cuh | 94 +++++++++++++ 5 files changed, 346 insertions(+) create mode 100644 examples/bloom_filter/arrow_bf_host_bulk_example.cu create mode 100644 include/cuco/arrow_bf_policy.cuh create mode 100644 include/cuco/detail/bloom_filter/arrow_bf_policy.inl create mode 100644 include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 8e7be947a..5065003b8 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -47,3 +47,4 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/ 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") +ConfigureExample(ARROW_BLOOM_FILTER_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/arrow_bf_host_bulk_example.cu") diff --git a/examples/bloom_filter/arrow_bf_host_bulk_example.cu b/examples/bloom_filter/arrow_bf_host_bulk_example.cu new file mode 100644 index 000000000..c791d7ef2 --- /dev/null +++ b/examples/bloom_filter/arrow_bf_host_bulk_example.cu @@ -0,0 +1,76 @@ +/* + * 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 +#include + +#include +#include +#include +#include + +#include + +int main(void) +{ + // Generate 10'000 keys and insert the first 5'000 into the filter. + int constexpr num_keys = 10'000; + int constexpr num_tp = num_keys * 0.5; + int constexpr num_tn = num_keys - num_tp; + + using policy_type = cuco::arrow_bf_policy; + using filter_type = + cuco::bloom_filter, cuda::thread_scope_device, policy_type>; + + // Create an instance of the Arrow BF policy + policy_type policy{200}; + + // Spawn a filter with 200 sub-filters. + filter_type filter{200, {}, policy}; + + thrust::device_vector keys(num_keys); + thrust::sequence(keys.begin(), keys.end(), 1); + + auto tp_begin = keys.begin(); + auto tp_end = tp_begin + num_tp; + auto tn_begin = tp_end; + auto tn_end = keys.end(); + + // Insert the first half of the keys. + filter.add(tp_begin, tp_end); + + thrust::device_vector tp_result(num_tp, false); + thrust::device_vector tn_result(num_keys - num_tp, false); + + // Query the filter for the previously inserted keys. + // This should result in a true-positive rate of TPR=1. + filter.contains(tp_begin, tp_end, tp_result.begin()); + + // Query the filter for the keys that are not present in the filter. + // Since bloom filters are probalistic data structures, the filter + // exhibits a false-positive rate FPR>0 depending on the number of bits in + // the filter and the number of hashes used per key. + filter.contains(tn_begin, tn_end, tn_result.begin()); + + float tp_rate = + float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); + float fp_rate = + float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); + + std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; + + return 0; +} \ No newline at end of file diff --git a/include/cuco/arrow_bf_policy.cuh b/include/cuco/arrow_bf_policy.cuh new file mode 100644 index 000000000..b81cd7da5 --- /dev/null +++ b/include/cuco/arrow_bf_policy.cuh @@ -0,0 +1,123 @@ +/* + * 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 + +#include + +namespace cuco { + +/** + * @brief A policy that defines how a Blocked Bloom Filter generates and stores a key's fingerprint. + * + * @note `Word` type must be an atomically updatable integral type. `WordsPerBlock` must + * be a power-of-two. + * + * @tparam Hash Hash function used to generate a key's fingerprint + * @tparam Word Underlying word/segment type of a filter block + * @tparam WordsPerBlock Number of words/segments in each block + */ +template +class arrow_bf_policy { + using impl_type = cuco::detail::arrow_bf_policy_impl; + + public: + using hasher = typename impl_type::hasher; ///< Type of the hash function + using hash_argument_type = typename impl_type::hash_argument_type; ///< Hash function input type + using hash_result_type = typename impl_type::hash_result_type; ///< hash function output type + using word_type = + typename impl_type::word_type; ///< Underlying word/segment type of a filter block + + static constexpr std::uint32_t words_per_block = + impl_type::words_per_block; ///< Number of words/segments in each filter block + + static constexpr std::uint32_t bits_set_per_block = + impl_type::bits_set_per_block; ///< Number of words/segments in each filter block + + public: + /** + * @brief Constructs the `arrow_bf_policy` object. + * + * @MH: Fix this doc. + * + * @throws Compile-time error if the specified number of words in a filter block is not a + * power-of-two or is larger than 32. If called from host: throws exception; If called from + * device: Traps the kernel. + * + * @throws If the `hash_result_type` is too narrow to generate the requested number of + * `pattern_bits`. If called from host: throws exception; If called from device: Traps the kernel. + * + * @throws If `pattern_bits` is smaller than the number of words in a filter block or larger than + * the total number of bits in a filter block. If called from host: throws exception; If called + * from device: Traps the kernel. + * + * @param num_blocks Number of bloom filter blocks + * @param hash Hash function used to generate a key's fingerprint + */ + __host__ __device__ constexpr arrow_bf_policy(std::uint32_t num_blocks, hasher hash = {}); + + /** + * @brief Generates the hash value for a given key. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @param key The key to hash + * + * @return The hash value of the key + */ + __device__ constexpr hash_result_type hash(hash_argument_type const& key) const; + + /** + * @brief Determines the filter block a key is added into. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @tparam Extent Size type that is used to determine the number of blocks in the filter + * + * @param hash Hash value of the key + * @param num_blocks Number of block in the filter + * + * @return The block index for the given key's hash value + */ + template + __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const; + + /** + * @brief Determines the fingerprint pattern for a word/segment within the filter block for a + * given key's hash value. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @param hash Hash value of the key + * @param word_index Target word/segment within the filter block + * + * @return The bit pattern for the word/segment in the filter block + */ + __device__ constexpr word_type word_pattern(hash_result_type hash, + std::uint32_t word_index) const; + + private: + impl_type impl_; ///< Policy implementation +}; + +} // namespace cuco + +#include \ No newline at end of file diff --git a/include/cuco/detail/bloom_filter/arrow_bf_policy.inl b/include/cuco/detail/bloom_filter/arrow_bf_policy.inl new file mode 100644 index 000000000..cfdf3c84b --- /dev/null +++ b/include/cuco/detail/bloom_filter/arrow_bf_policy.inl @@ -0,0 +1,52 @@ +/* + * 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 + +namespace cuco { + +template +__host__ __device__ constexpr arrow_bf_policy::arrow_bf_policy(std::uint32_t num_blocks, + hasher hash) + : impl_{num_blocks, hash} +{ +} + +template +__device__ constexpr typename arrow_bf_policy::hash_result_type arrow_bf_policy::hash( + typename arrow_bf_policy::hash_argument_type const& key) const +{ + return impl_.hash(key); +} + +template +template +__device__ constexpr auto arrow_bf_policy::block_index( + typename arrow_bf_policy::hash_result_type hash, Extent num_blocks) const +{ + return impl_.block_index(hash, num_blocks); +} + +template +__device__ constexpr typename arrow_bf_policy::word_type arrow_bf_policy::word_pattern( + arrow_bf_policy::hash_result_type hash, std::uint32_t word_index) const +{ + return impl_.word_pattern(hash, word_index); +} + +} // namespace cuco \ No newline at end of file diff --git a/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh b/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh new file mode 100644 index 000000000..eb15f7c3a --- /dev/null +++ b/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh @@ -0,0 +1,94 @@ +/* + * 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 +#include + +#include +#include +#include +#include + +#include +#include + +namespace cuco::detail { + +template +class arrow_bf_policy_impl { + public: + using hasher = cuco::xxhash_64; + using word_type = std::uint32_t; + using hash_argument_type = typename hasher::argument_type; + using hash_result_type = decltype(std::declval()(std::declval())); + + // Bytes in a tiny BF block. + static constexpr int bytes_per_filter_block = 32; // hardcoded values from Arrow BF + static constexpr uint32_t bits_set_per_block = 8; // hardcoded values from Arrow BF + static constexpr uint32_t words_per_block = 8; // hardcoded values from Arrow BF + static constexpr std::uint32_t min_arrow_bf_bytes = 32; + static constexpr std::uint32_t max_arrow_bf_bytes = 128 * 1024 * 1024; + + __host__ __device__ explicit constexpr arrow_bf_policy_impl(std::uint32_t num_blocks, hasher hash) + : hash_{hash} + { + NV_DISPATCH_TARGET( + NV_IS_HOST, + (CUCO_EXPECTS(num_blocks >= 1 and num_blocks <= (max_arrow_bf_bytes / bytes_per_filter_block), + "`num_blocks` must be in the range of [1, 4194304]");), + NV_IS_DEVICE, + (if (num_blocks < 1 or num_blocks > (max_arrow_bf_bytes / bytes_per_filter_block)) { + __trap(); // TODO this kills the kernel and corrupts the CUDA context. Not ideal. + })); + } + + __device__ constexpr hash_result_type hash(hash_argument_type const& key) const + { + return hash_(key); + } + + template + __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const + { + constexpr auto hash_bits = cuda::std::numeric_limits::digits; + return static_cast(((hash >> hash_bits) * num_blocks) >> hash_bits); + } + + __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const + { + // The block-based algorithm needs eight odd SALT values to calculate eight indexes of bit to + // set, one bit in each 32-bit word. + + // @MH: Fix this + constexpr std::uint32_t SALT[bits_set_per_block] = {0x47b6137bU, + 0x44974d91U, + 0x8824ad5bU, + 0xa2b7289dU, + 0x705495c7U, + 0x2df1424bU, + 0x9efc4947U, + 0x5c6bfb31U}; + word_type const key = static_cast(hash); + return word_type{1} << ((key * SALT[word_index]) >> 27); + } + + private: + hasher hash_; +}; + +} // namespace cuco::detail \ No newline at end of file From 8c614d07e928a9bcfdafc6c7a84ccde5667aae06 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 25 Oct 2024 20:23:16 +0000 Subject: [PATCH 02/29] Add test and benchmark --- benchmarks/bloom_filter/contains_bench.cu | 62 ++++++++++++++++++- .../arrow_bf_host_bulk_example.cu | 1 - include/cuco/arrow_bf_policy.cuh | 15 +---- include/cuco/bloom_filter.cuh | 1 + .../bloom_filter/arrow_bf_policy_impl.cuh | 25 ++++---- tests/bloom_filter/unique_sequence_test.cu | 15 +++++ 6 files changed, 92 insertions(+), 27 deletions(-) diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index 67ba80d95..8060cbf4a 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/benchmarks/bloom_filter/contains_bench.cu @@ -88,6 +88,57 @@ void bloom_filter_contains( }); } +/** + * @brief A benchmark evaluating `cuco::bloom_filter::contains_async` performance with + * `arrow_bf_policy` + */ +template +void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list) +{ + // cudaDeviceSetLimit(cudaLimitMaxL2FetchGranularity, 32); // slightly improves peformance if + // filter block fits into a 32B sector + using policy_type = cuco::arrow_bf_policy; + using filter_type = + cuco::bloom_filter, cuda::thread_scope_device, policy_type>; + + auto const num_keys = state.get_int64("NumInputs"); + auto const filter_size_mb = state.get_int64("FilterSizeMB"); + + std::size_t const num_sub_filters = + (filter_size_mb * 1024 * 1024) / + (sizeof(typename filter_type::word_type) * filter_type::words_per_block); + + try { + auto const policy = policy_type{static_cast(num_sub_filters)}; + } catch (std::exception const& e) { + state.skip(e.what()); // skip invalid configurations + } + + thrust::device_vector keys(num_keys); + thrust::device_vector result(num_keys, false); + + key_generator gen; + gen.generate(dist_from_state(state), keys.begin(), keys.end()); + + state.add_element_count(num_keys); + + filter_type filter{num_sub_filters, {}, {static_cast(num_sub_filters)}}; + + state.collect_dram_throughput(); + state.collect_l1_hit_rates(); + state.collect_l2_hit_rates(); + state.collect_loads_efficiency(); + state.collect_stores_efficiency(); + + add_fpr_summary(state, filter); + + filter.add(keys.begin(), keys.end()); + + state.exec([&](nvbench::launch& launch) { + filter.contains_async(keys.begin(), keys.end(), result.begin(), {launch.get_stream()}); + }); +} + NVBENCH_BENCH_TYPES(bloom_filter_contains, NVBENCH_TYPE_AXES(nvbench::type_list, nvbench::type_list, @@ -122,4 +173,13 @@ NVBENCH_BENCH_TYPES(bloom_filter_contains, .set_type_axes_names({"Key", "Hash", "Word", "WordsPerBlock", "Distribution"}) .set_max_noise(defaults::MAX_NOISE) .add_int64_axis("NumInputs", {defaults::BF_N}) - .add_int64_axis("FilterSizeMB", {defaults::BF_SIZE_MB}); \ No newline at end of file + .add_int64_axis("FilterSizeMB", {defaults::BF_SIZE_MB}); + +NVBENCH_BENCH_TYPES(arrow_bloom_filter_contains, + NVBENCH_TYPE_AXES(nvbench::type_list, + nvbench::type_list)) + .set_name("arrow_bloom_filter_contains_unique_size") + .set_type_axes_names({"Key", "Distribution"}) + .set_max_noise(defaults::MAX_NOISE) + .add_int64_axis("NumInputs", {defaults::BF_N / 2}) + .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); \ No newline at end of file diff --git a/examples/bloom_filter/arrow_bf_host_bulk_example.cu b/examples/bloom_filter/arrow_bf_host_bulk_example.cu index c791d7ef2..bf7530360 100644 --- a/examples/bloom_filter/arrow_bf_host_bulk_example.cu +++ b/examples/bloom_filter/arrow_bf_host_bulk_example.cu @@ -14,7 +14,6 @@ * limitations under the License. */ -#include #include #include diff --git a/include/cuco/arrow_bf_policy.cuh b/include/cuco/arrow_bf_policy.cuh index b81cd7da5..1af5df917 100644 --- a/include/cuco/arrow_bf_policy.cuh +++ b/include/cuco/arrow_bf_policy.cuh @@ -53,18 +53,9 @@ class arrow_bf_policy { /** * @brief Constructs the `arrow_bf_policy` object. * - * @MH: Fix this doc. - * - * @throws Compile-time error if the specified number of words in a filter block is not a - * power-of-two or is larger than 32. If called from host: throws exception; If called from - * device: Traps the kernel. - * - * @throws If the `hash_result_type` is too narrow to generate the requested number of - * `pattern_bits`. If called from host: throws exception; If called from device: Traps the kernel. - * - * @throws If `pattern_bits` is smaller than the number of words in a filter block or larger than - * the total number of bits in a filter block. If called from host: throws exception; If called - * from device: Traps the kernel. + * @throws If number of filter blocks (`num_blocks`) is smaller than 1 + * or larger than 4194304. If called from host: throws exception; + * If called from device: Traps the kernel. * * @param num_blocks Number of bloom filter blocks * @param hash Hash function used to generate a key's fingerprint diff --git a/include/cuco/bloom_filter.cuh b/include/cuco/bloom_filter.cuh index 16642ab57..6f26ab947 100644 --- a/include/cuco/bloom_filter.cuh +++ b/include/cuco/bloom_filter.cuh @@ -16,6 +16,7 @@ #pragma once +#include #include #include #include diff --git a/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh b/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh index eb15f7c3a..1964bc1f3 100644 --- a/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh +++ b/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh @@ -32,21 +32,21 @@ namespace cuco::detail { template class arrow_bf_policy_impl { public: - using hasher = cuco::xxhash_64; - using word_type = std::uint32_t; + using hasher = cuco::xxhash_64; ///< xxhash_64 hasher for Arrow bloom filter policy + using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy using hash_argument_type = typename hasher::argument_type; using hash_result_type = decltype(std::declval()(std::declval())); - // Bytes in a tiny BF block. - static constexpr int bytes_per_filter_block = 32; // hardcoded values from Arrow BF - static constexpr uint32_t bits_set_per_block = 8; // hardcoded values from Arrow BF - static constexpr uint32_t words_per_block = 8; // hardcoded values from Arrow BF - static constexpr std::uint32_t min_arrow_bf_bytes = 32; - static constexpr std::uint32_t max_arrow_bf_bytes = 128 * 1024 * 1024; + static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow bf block + static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow bf block __host__ __device__ explicit constexpr arrow_bf_policy_impl(std::uint32_t num_blocks, hasher hash) : hash_{hash} { + constexpr std::uint32_t bytes_per_filter_block = 32; ///< Number of bytes in one Arrow bf block + constexpr std::uint32_t max_arrow_bf_bytes = + 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter + NV_DISPATCH_TARGET( NV_IS_HOST, (CUCO_EXPECTS(num_blocks >= 1 and num_blocks <= (max_arrow_bf_bytes / bytes_per_filter_block), @@ -71,10 +71,8 @@ class arrow_bf_policy_impl { __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const { - // The block-based algorithm needs eight odd SALT values to calculate eight indexes of bit to - // set, one bit in each 32-bit word. - - // @MH: Fix this + // Arrow's block-based bloom filter algorithm needs these eight odd SALT values to calculate + // eight indexes of bit to set, one bit in each 32-bit (uint32_t) word. constexpr std::uint32_t SALT[bits_set_per_block] = {0x47b6137bU, 0x44974d91U, 0x8824ad5bU, @@ -83,7 +81,8 @@ class arrow_bf_policy_impl { 0x2df1424bU, 0x9efc4947U, 0x5c6bfb31U}; - word_type const key = static_cast(hash); + + word_type const key = static_cast(hash); return word_type{1} << ((key * SALT[word_index]) >> 27); } diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 3919e77bf..428d6e3f0 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -103,3 +103,18 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence", test_unique_sequence(filter, num_keys); } + +TEMPLATE_TEST_CASE_SIG("Unique sequence Arrow BF", + "", + ((class Key, class Policy), Key, Policy), + (int32_t, cuco::arrow_bf_policy), + (float, cuco::arrow_bf_policy)) +{ + using filter_type = + cuco::bloom_filter, cuda::thread_scope_device, Policy>; + constexpr size_type num_keys{400}; + + auto filter = filter_type{1000, {}, {1000}}; + + test_unique_sequence(filter, num_keys); +} From f4d67d8896384ea864fa33650666c6aa1aeba16c Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 25 Oct 2024 20:40:40 +0000 Subject: [PATCH 03/29] Minor improvements to example --- .../arrow_bf_host_bulk_example.cu | 24 +++++++++++-------- 1 file changed, 14 insertions(+), 10 deletions(-) diff --git a/examples/bloom_filter/arrow_bf_host_bulk_example.cu b/examples/bloom_filter/arrow_bf_host_bulk_example.cu index bf7530360..abe1dac84 100644 --- a/examples/bloom_filter/arrow_bf_host_bulk_example.cu +++ b/examples/bloom_filter/arrow_bf_host_bulk_example.cu @@ -25,22 +25,26 @@ int main(void) { - // Generate 10'000 keys and insert the first 5'000 into the filter. - int constexpr num_keys = 10'000; - int constexpr num_tp = num_keys * 0.5; - int constexpr num_tn = num_keys - num_tp; + using key_type = int; + using extent_type = size_t; - using policy_type = cuco::arrow_bf_policy; + // Generate 10'000 keys and insert the first 5'000 into the filter containing of 200 sub-filters. + int constexpr num_keys = 10'000; + int constexpr num_tp = num_keys * 0.5; + int constexpr num_tn = num_keys - num_tp; + int constexpr sub_filters = 200; + + using policy_type = cuco::arrow_bf_policy; using filter_type = - cuco::bloom_filter, cuda::thread_scope_device, policy_type>; + cuco::bloom_filter, cuda::thread_scope_device, policy_type>; - // Create an instance of the Arrow BF policy - policy_type policy{200}; + // Create an instance of the Arrow BF policy with 200 sub-filters + policy_type policy{sub_filters}; // Spawn a filter with 200 sub-filters. - filter_type filter{200, {}, policy}; + filter_type filter{sub_filters, {}, policy}; - thrust::device_vector keys(num_keys); + thrust::device_vector keys(num_keys); thrust::sequence(keys.begin(), keys.end(), 1); auto tp_begin = keys.begin(); From a5e625ff15b3ac97d0ae85ed80632636d489a949 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 25 Oct 2024 21:08:44 +0000 Subject: [PATCH 04/29] Update the example to insert and evaluate both policies in one --- .../arrow_bf_host_bulk_example.cu | 79 ------------------ examples/bloom_filter/host_bulk_example.cu | 82 ++++++++++++------- 2 files changed, 53 insertions(+), 108 deletions(-) delete mode 100644 examples/bloom_filter/arrow_bf_host_bulk_example.cu diff --git a/examples/bloom_filter/arrow_bf_host_bulk_example.cu b/examples/bloom_filter/arrow_bf_host_bulk_example.cu deleted file mode 100644 index abe1dac84..000000000 --- a/examples/bloom_filter/arrow_bf_host_bulk_example.cu +++ /dev/null @@ -1,79 +0,0 @@ -/* - * 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 - -#include -#include -#include -#include - -#include - -int main(void) -{ - using key_type = int; - using extent_type = size_t; - - // Generate 10'000 keys and insert the first 5'000 into the filter containing of 200 sub-filters. - int constexpr num_keys = 10'000; - int constexpr num_tp = num_keys * 0.5; - int constexpr num_tn = num_keys - num_tp; - int constexpr sub_filters = 200; - - using policy_type = cuco::arrow_bf_policy; - using filter_type = - cuco::bloom_filter, cuda::thread_scope_device, policy_type>; - - // Create an instance of the Arrow BF policy with 200 sub-filters - policy_type policy{sub_filters}; - - // Spawn a filter with 200 sub-filters. - filter_type filter{sub_filters, {}, policy}; - - thrust::device_vector keys(num_keys); - thrust::sequence(keys.begin(), keys.end(), 1); - - auto tp_begin = keys.begin(); - auto tp_end = tp_begin + num_tp; - auto tn_begin = tp_end; - auto tn_end = keys.end(); - - // Insert the first half of the keys. - filter.add(tp_begin, tp_end); - - thrust::device_vector tp_result(num_tp, false); - thrust::device_vector tn_result(num_keys - num_tp, false); - - // Query the filter for the previously inserted keys. - // This should result in a true-positive rate of TPR=1. - filter.contains(tp_begin, tp_end, tp_result.begin()); - - // Query the filter for the keys that are not present in the filter. - // Since bloom filters are probalistic data structures, the filter - // exhibits a false-positive rate FPR>0 depending on the number of bits in - // the filter and the number of hashes used per key. - filter.contains(tn_begin, tn_end, tn_result.begin()); - - float tp_rate = - float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); - float fp_rate = - float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); - - std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; - - return 0; -} \ No newline at end of file diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu index 14a2a5cfa..9a920aa32 100644 --- a/examples/bloom_filter/host_bulk_example.cu +++ b/examples/bloom_filter/host_bulk_example.cu @@ -25,44 +25,68 @@ int main(void) { - // Generate 10'000 keys and insert the first 5'000 into the filter. - int constexpr num_keys = 10'000; - int constexpr num_tp = num_keys * 0.5; - int constexpr num_tn = num_keys - num_tp; + using key_type = int; + using extent_type = size_t; - // Spawn a filter with 200 sub-filters. - cuco::bloom_filter filter{200}; + // Generate 10'000 keys and insert the first 5'000 into the filter containing of 200 sub-filters. + int constexpr num_keys = 10'000; + int constexpr num_tp = num_keys * 0.5; + int constexpr num_tn = num_keys - num_tp; + int constexpr sub_filters = 200; - thrust::device_vector keys(num_keys); - thrust::sequence(keys.begin(), keys.end(), 1); + // Lambda function to bulk insert to the provided bloom filter and evaluate FPR. + auto bulk_insert_and_evaluate_bloom_filter = [&](auto& filter) { + thrust::device_vector keys(num_keys); + thrust::sequence(keys.begin(), keys.end(), 1); - auto tp_begin = keys.begin(); - auto tp_end = tp_begin + num_tp; - auto tn_begin = tp_end; - auto tn_end = keys.end(); + auto tp_begin = keys.begin(); + auto tp_end = tp_begin + num_tp; + auto tn_begin = tp_end; + auto tn_end = keys.end(); - // Insert the first half of the keys. - filter.add(tp_begin, tp_end); + // Insert the first half of the keys. + filter.add(tp_begin, tp_end); - thrust::device_vector tp_result(num_tp, false); - thrust::device_vector tn_result(num_keys - num_tp, false); + thrust::device_vector tp_result(num_tp, false); + thrust::device_vector tn_result(num_keys - num_tp, false); - // Query the filter for the previously inserted keys. - // This should result in a true-positive rate of TPR=1. - filter.contains(tp_begin, tp_end, tp_result.begin()); + // Query the filter for the previously inserted keys. + // This should result in a true-positive rate of TPR=1. + filter.contains(tp_begin, tp_end, tp_result.begin()); - // Query the filter for the keys that are not present in the filter. - // Since bloom filters are probalistic data structures, the filter - // exhibits a false-positive rate FPR>0 depending on the number of bits in - // the filter and the number of hashes used per key. - filter.contains(tn_begin, tn_end, tn_result.begin()); + // Query the filter for the keys that are not present in the filter. + // Since bloom filters are probalistic data structures, the filter + // exhibits a false-positive rate FPR>0 depending on the number of bits in + // the filter and the number of hashes used per key. + filter.contains(tn_begin, tn_end, tn_result.begin()); - float tp_rate = - float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); - float fp_rate = - float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); + float tp_rate = float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / + float(num_tp); + float fp_rate = float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / + float(num_tn); - std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; + std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; + }; + + // Spawn a bloom filter with default policy and 200 sub-filters. + cuco::bloom_filter def_filter{sub_filters}; + + // bulk insert to the bloom filter and evaluate + std::cout << "Bulk insert and evaluate bloom filter with default policy: " << std::endl; + bulk_insert_and_evaluate_bloom_filter(def_filter); + + // Arrow bloom filter policy type + using arrow_policy_type = cuco::arrow_bf_policy; + // bloom filter with arrow policy type + using arrow_policy_filter_type = cuco:: + bloom_filter, cuda::thread_scope_device, arrow_policy_type>; + + // Spawn a bloom filter with arrow policy and 200 sub-filters. + arrow_policy_filter_type arrow_filter{sub_filters, {}, arrow_policy_type{sub_filters}}; + + // bulk insert to the bloom filter and evaluate + std::cout << "Bulk insert and evaluate bloom filter with arrow policy: " << std::endl; + bulk_insert_and_evaluate_bloom_filter(arrow_filter); return 0; } \ No newline at end of file From 7bbb36a1f8ffe729d29c6bc75b51f3f1dc9072fa Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Fri, 25 Oct 2024 21:08:54 +0000 Subject: [PATCH 05/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index a11f5afdb..0dfae136f 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVm1v20YM_iuE9mF2I78FCwo4L4CXpJuxwsnstEUxD8HpRFuHynfavdgxgvz38U6SLSduMSwBEptHPnz4kEfpOTJojFDSRMO_niORRsNBHOVMLh1bYjSMuEtZFEdGOc399967uYR3cK2KrRbLzEKLt-G0f_pLDJPP45vxCK7vpvd309HD-G7S9b7B_6PgKA2m4GSKGmyGMCoYp3_VSQyfUXsicNrtQ8s7zKPqbB61zwPKVjlYsS1IZcEZJBhhYCFyBHziWFgQErhaFblgkiNshM1Cqgon0IGvFYhKLCN_RhEFfVs0PYHZHXX_k1lbDHu9zWbTZYF2V-llLy-dTe_j-Pp2MrvtEPVd2CeZk7Kg8R8nNBWebIEVxIyzhPjmbANKA1tqpDOrPPONFlbIZQxGLeyGaQw4qTBWi8TZA_FqnlR_04HkY5KEG81gPJtH8OtoNp7FAefL-OH3u08P8GU0nY4mD-PbGdxNqVmTm7FvFX37AKPJV_hjPLmJAUk6SoVPhfZVEFXhZcW01HCGeEBjoUpapkAuFoJDPUGwVGvUksqCAvVKlLNGJNOAk4uVsMwG25viQqreXM7lT0Ly3KUIF9xx1UtypVaP1HeLustddnXoYzPtjO1x5aTt-sM3RymuKcXjGrlV-rgLPiF3nthjoahp2-NehrqLNGrd1xyEoqYgWwWzkJYmTsjWWom0PZfPVBj0evAbStTMIgz6P_f7ffiG2yANDYNBbYMWC6GNhbNwTjiqMobSPYzH5iSf9a0C6VaPAeWywjw_7mMLGurLvfs76HfPvucrD307FcK5Ly3UMSvYxl-kklZ57U6Jr3FJp7SZwNU3bzhsdu-C0l1Vcc8U8lKBlvIOhweNKp09h1ZNJuyFnXfdjZY_6ia4JMnbcYjookz950G7ysCc17J4DF5UXjPkvOlAgaH-ne_Jvv7aS-5gyoCDoxpgT2Ov3Ph1ozOWL-pdFAK8Y9VulqatmkVcZWr_ULBEqfzKe9IldrltlcRjWLDc4KF2RwNlM_Cg-Q2MqpQ_HeptYzp3O4E2yFooZ_JtNde0p3alUeCDX-ImUy5PoUwHYS1b7bBTKENLcY0Q7gkJ83A_vRw0VaFh9XvcvJEm3hdeN_a_0A1l2oxZoA0cnjR-BaIMrF5dPj_7wj9pwkxXJyYEFlolLKfVTMswZZYB7QPHrSOsuAFToeBTJhJhjb9EXtdXdX-4n171IcWCyvKrVJVMqBUJMSdVQqyoa2tU5bfJoWvGTIbGPz5Tv5J9vcfllDs5ZSWn_I6ci1yRXF5tT_bSmypjqx6vsIxbh8N2pEFNU3Vj_Ri029CrAMvxK2e3zLv4H3lfV9I0_SCvrCs2Ng3YFi4u6HHrh5KetvS5FiGYfd8q-2JvD7GUJK-wNNJUSPCr-oXetPz7C73R6P0LWSTXnA9Oz9yAjlVhy7e1qENAl_zkZPAeOkzz7NKsHt_3odOhzW3pj6UcmHZytkrCK1wukgYm5zwn47p86SID1Su_RS9xfU67-uCctIte_g6__wINDYAL)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV2tv2zYU_SsXGtDZrfxIsKKA8wDcxzZjhdPZaYuiLgRKomyiMqnxYccL8t93SUqWZDtYin1Y8iEx74Pnnnt4Sd8HiirFBFfB6Ot9wNJgdBYGOeFLQ5Y0GAWJSUkQBkoYmdjPg-cLDs_hjSh2ki1XGjpJF86H57-EMP00eTsZw5ub2Yeb2fh2cjPtW1_n_54llCuaguEplaBXFMYFSfBPaQnhE5UWCJz3h9CxDougtC2C7oXLshMG1mQHXGgwimIapiBjOQV6l9BCA-OQiHWRM8ITClumV26rMo-DA1_KJCLWBP0JRhT4KWt6AtF76PZnpXUxGgy2222fONh9IZeD3DurwfvJm3fT-bseQt-HfeQ5MguS_mWYxMLjHZACkSUkRrw52YKQQJaSok0Li3wrmWZ8GYISmd4SSV2elCktWWx0i7wKJ9bfdED6CEfixnOYzBcBvB7PJ_PQ5fk8uf395uMtfB7PZuPp7eTdHG5m2Kzp24ltFX76FcbTL_DHZPo2BIrU4Vb0rpC2CoTKLK009RzOKW3ByISHpQqasIwlUCkIlmJDJceyoKByzbzWEGTq8uRszTTRbu2oOLfVYMEX_CfGk9ykFC4Tk4hBnAuxjrDvmsp-YlbXbR-9kkbpQSIM131rPDKldINbRBuaaCFPu9A7mhgLLCoENm132kthdylKrX-IgQlsCiVrt8y4RsUx3tkIlnYX_B4LQ_1aUr7TXaR3BbUiu0IR6IvaRu805dqbr0Cxv2lkzdZhMIDfKKeSaApnw5-Hw6HN5HjFJIpK7YjMmFQaXjo75hblouUNVc-t_u1GqP1z9FAm7nmjstTbCOultFUBcLOO3B4Oqd_04rSbLtyhQbd90HMY9l8-5s6P3HtlnhMRiLJsvcKIcweipOQ9WccpgczwxHbOHqvY5N_3jPj6Cyk2LLUn0sqoosMyRzckN5bSXz_MHAPElCkinyJCr6jyipoyRChfn33r2IBnZcouuD4DeK2MRi3VXVadv3ad61Slu0nXCKoU1rHGfkyXKKNu6GL6lKf2_7NuSUEJWBeR80NMzaCLtgsGO8733i-anFd-fJ_KhxwYqyQ1nD0U7MfkUIkrkmfVpHUh3rU8ySRNOxWasNyvzneaxViI_Nr64pgyue74EkLISK7oIZcnQ3kztCW-RpZ9SX8aKnfNY1RNPpyTGyaMynel3FBgjRIx9NZeVmolTJ6C3xLc9aOlob1CKBz-GwruSCNFtx9mV2dtfsoTq45ICmsCqmY_DbQrV6-IBrxt3K1qxz0OHYus9q9rmDN7rzYPjnKheKZikuNFhKM_JZoATj-TaIPZwkaifR56t2Ix0xjsOT6oH8_f9RBSWmBxbkB5NNiWGNEjOy6W1RU2arMHue28ImpFlX0wpPYSslU_RizfE8tLYvmjxGa5QOIs8xbylf_cqbTmbp9OW3kn-tRcKo-z1UO3C4NFORbbP34XL9BK3x5J9mNIDgtrLv0oEl6zonTqttRweYkPEStjfIfg_xVRbtl2uFzP6nUXi3vnrq6HerDPC7K1R6U1sd3rLqUZsSfJX9Ku-aeuMvtsGI2aA7sxfzFHuXjfuFwa25-4RB69O2zIAQevG-Gta-bfyhmBp6jNy5Puo05dVLcuZCyl2La3LYmzTNQvD2L9yodP9f7wFHpLnJXGmsWLiqujmlzIE_bxMe3tvKZO9i0sXfwr6bLxWLq2ppSMRqh5StJI4QOfRpXwj2q7forOWkU8prJHy_GGY42FcP9wAlFLhv-LDpvl_hcVNguvdCgp3gsc7LvtAb9Z2u9r-A1O1l9AA75JkrPzl-YMzaLQ_ttp0EMAV8mLF2evoEdksrpS6-jVEHo9fBbqnus_vup6uXsEYmTO4kbOJElyXNz4L5m4gMONfw8ewsqOYmrZUS_Bwzf3-w_h-DkJ)) \ No newline at end of file From d31d147bd846c606666915b80644c2b0bc96aa68 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 25 Oct 2024 21:37:53 +0000 Subject: [PATCH 06/29] Convert templated lambda function into normal one --- examples/CMakeLists.txt | 1 - examples/bloom_filter/host_bulk_example.cu | 81 +++++++++++----------- 2 files changed, 42 insertions(+), 40 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 5065003b8..8e7be947a 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -47,4 +47,3 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/ 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") -ConfigureExample(ARROW_BLOOM_FILTER_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/arrow_bf_host_bulk_example.cu") diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu index 9a920aa32..b3881c29e 100644 --- a/examples/bloom_filter/host_bulk_example.cu +++ b/examples/bloom_filter/host_bulk_example.cu @@ -23,70 +23,73 @@ #include -int main(void) +// key type for bloom filter +using key_type = int; + +// Function to bulk insert to the provided bloom filter and evaluate FPR. +template +void bulk_insert_and_eval_bloom_filter(FilterType& filter, int num_keys) { - using key_type = int; - using extent_type = size_t; + // Insert the first half keys into the filter. + int const num_tp = num_keys * 0.5; + int const num_tn = num_keys - num_tp; - // Generate 10'000 keys and insert the first 5'000 into the filter containing of 200 sub-filters. - int constexpr num_keys = 10'000; - int constexpr num_tp = num_keys * 0.5; - int constexpr num_tn = num_keys - num_tp; - int constexpr sub_filters = 200; + thrust::device_vector keys(num_keys); + thrust::sequence(keys.begin(), keys.end(), 1); - // Lambda function to bulk insert to the provided bloom filter and evaluate FPR. - auto bulk_insert_and_evaluate_bloom_filter = [&](auto& filter) { - thrust::device_vector keys(num_keys); - thrust::sequence(keys.begin(), keys.end(), 1); + auto tp_begin = keys.begin(); + auto tp_end = tp_begin + num_tp; + auto tn_begin = tp_end; + auto tn_end = keys.end(); - auto tp_begin = keys.begin(); - auto tp_end = tp_begin + num_tp; - auto tn_begin = tp_end; - auto tn_end = keys.end(); + // Insert the first half of the keys. + filter.add(tp_begin, tp_end); - // Insert the first half of the keys. - filter.add(tp_begin, tp_end); + thrust::device_vector tp_result(num_tp, false); + thrust::device_vector tn_result(num_keys - num_tp, false); - thrust::device_vector tp_result(num_tp, false); - thrust::device_vector tn_result(num_keys - num_tp, false); + // Query the filter for the previously inserted keys. + // This should result in a true-positive rate of TPR=1. + filter.contains(tp_begin, tp_end, tp_result.begin()); - // Query the filter for the previously inserted keys. - // This should result in a true-positive rate of TPR=1. - filter.contains(tp_begin, tp_end, tp_result.begin()); + // Query the filter for the keys that are not present in the filter. + // Since bloom filters are probalistic data structures, the filter + // exhibits a false-positive rate FPR>0 depending on the number of bits in + // the filter and the number of hashes used per key. + filter.contains(tn_begin, tn_end, tn_result.begin()); - // Query the filter for the keys that are not present in the filter. - // Since bloom filters are probalistic data structures, the filter - // exhibits a false-positive rate FPR>0 depending on the number of bits in - // the filter and the number of hashes used per key. - filter.contains(tn_begin, tn_end, tn_result.begin()); + float tp_rate = + float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); + float fp_rate = + float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); - float tp_rate = float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / - float(num_tp); - float fp_rate = float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / - float(num_tn); + std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; +} - std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; - }; +int main(void) +{ + int constexpr num_keys = 10'000; ///< Generate 10'000 keys + int constexpr sub_filters = 200; ///< 200 sub-filters per bloom filter // Spawn a bloom filter with default policy and 200 sub-filters. - cuco::bloom_filter def_filter{sub_filters}; + cuco::bloom_filter filter{sub_filters}; // bulk insert to the bloom filter and evaluate std::cout << "Bulk insert and evaluate bloom filter with default policy: " << std::endl; - bulk_insert_and_evaluate_bloom_filter(def_filter); + bulk_insert_and_eval_bloom_filter(filter, num_keys); // Arrow bloom filter policy type using arrow_policy_type = cuco::arrow_bf_policy; // bloom filter with arrow policy type using arrow_policy_filter_type = cuco:: - bloom_filter, cuda::thread_scope_device, arrow_policy_type>; + bloom_filter, cuda::thread_scope_device, arrow_policy_type>; // Spawn a bloom filter with arrow policy and 200 sub-filters. - arrow_policy_filter_type arrow_filter{sub_filters, {}, arrow_policy_type{sub_filters}}; + arrow_policy_filter_type filter_arrow_policy{sub_filters, {}, arrow_policy_type{sub_filters}}; // bulk insert to the bloom filter and evaluate std::cout << "Bulk insert and evaluate bloom filter with arrow policy: " << std::endl; - bulk_insert_and_evaluate_bloom_filter(arrow_filter); + bulk_insert_and_eval_bloom_filter(filter_arrow_policy, num_keys); return 0; } \ No newline at end of file From 863a508731d7bdab470033808bd91b7241ce0e1d Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Fri, 25 Oct 2024 21:38:03 +0000 Subject: [PATCH 07/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 0dfae136f..39990b323 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV2tv2zYU_SsXGtDZrfxIsKKA8wDcxzZjhdPZaYuiLgRKomyiMqnxYccL8t93SUqWZDtYin1Y8iEx74Pnnnt4Sd8HiirFBFfB6Ot9wNJgdBYGOeFLQ5Y0GAWJSUkQBkoYmdjPg-cLDs_hjSh2ki1XGjpJF86H57-EMP00eTsZw5ub2Yeb2fh2cjPtW1_n_54llCuaguEplaBXFMYFSfBPaQnhE5UWCJz3h9CxDougtC2C7oXLshMG1mQHXGgwimIapiBjOQV6l9BCA-OQiHWRM8ITClumV26rMo-DA1_KJCLWBP0JRhT4KWt6AtF76PZnpXUxGgy2222fONh9IZeD3DurwfvJm3fT-bseQt-HfeQ5MguS_mWYxMLjHZACkSUkRrw52YKQQJaSok0Li3wrmWZ8GYISmd4SSV2elCktWWx0i7wKJ9bfdED6CEfixnOYzBcBvB7PJ_PQ5fk8uf395uMtfB7PZuPp7eTdHG5m2Kzp24ltFX76FcbTL_DHZPo2BIrU4Vb0rpC2CoTKLK009RzOKW3ByISHpQqasIwlUCkIlmJDJceyoKByzbzWEGTq8uRszTTRbu2oOLfVYMEX_CfGk9ykFC4Tk4hBnAuxjrDvmsp-YlbXbR-9kkbpQSIM131rPDKldINbRBuaaCFPu9A7mhgLLCoENm132kthdylKrX-IgQlsCiVrt8y4RsUx3tkIlnYX_B4LQ_1aUr7TXaR3BbUiu0IR6IvaRu805dqbr0Cxv2lkzdZhMIDfKKeSaApnw5-Hw6HN5HjFJIpK7YjMmFQaXjo75hblouUNVc-t_u1GqP1z9FAm7nmjstTbCOultFUBcLOO3B4Oqd_04rSbLtyhQbd90HMY9l8-5s6P3HtlnhMRiLJsvcKIcweipOQ9WccpgczwxHbOHqvY5N_3jPj6Cyk2LLUn0sqoosMyRzckN5bSXz_MHAPElCkinyJCr6jyipoyRChfn33r2IBnZcouuD4DeK2MRi3VXVadv3ad61Slu0nXCKoU1rHGfkyXKKNu6GL6lKf2_7NuSUEJWBeR80NMzaCLtgsGO8733i-anFd-fJ_KhxwYqyQ1nD0U7MfkUIkrkmfVpHUh3rU8ySRNOxWasNyvzneaxViI_Nr64pgyue74EkLISK7oIZcnQ3kztCW-RpZ9SX8aKnfNY1RNPpyTGyaMynel3FBgjRIx9NZeVmolTJ6C3xLc9aOlob1CKBz-GwruSCNFtx9mV2dtfsoTq45ICmsCqmY_DbQrV6-IBrxt3K1qxz0OHYus9q9rmDN7rzYPjnKheKZikuNFhKM_JZoATj-TaIPZwkaifR56t2Ix0xjsOT6oH8_f9RBSWmBxbkB5NNiWGNEjOy6W1RU2arMHue28ImpFlX0wpPYSslU_RizfE8tLYvmjxGa5QOIs8xbylf_cqbTmbp9OW3kn-tRcKo-z1UO3C4NFORbbP34XL9BK3x5J9mNIDgtrLv0oEl6zonTqttRweYkPEStjfIfg_xVRbtl2uFzP6nUXi3vnrq6HerDPC7K1R6U1sd3rLqUZsSfJX9Ku-aeuMvtsGI2aA7sxfzFHuXjfuFwa25-4RB69O2zIAQevG-Gta-bfyhmBp6jNy5Puo05dVLcuZCyl2La3LYmzTNQvD2L9yodP9f7wFHpLnJXGmsWLiqujmlzIE_bxMe3tvKZO9i0sXfwr6bLxWLq2ppSMRqh5StJI4QOfRpXwj2q7forOWkU8prJHy_GGY42FcP9wAlFLhv-LDpvl_hcVNguvdCgp3gsc7LvtAb9Z2u9r-A1O1l9AA75JkrPzl-YMzaLQ_ttp0EMAV8mLF2evoEdksrpS6-jVEHo9fBbqnus_vup6uXsEYmTO4kbOJElyXNz4L5m4gMONfw8ewsqOYmrZUS_Bwzf3-w_h-DkJ)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QbUXFqhgZ1aToKgKogluZIWoXbZvUhWDf97Z5ZLibTlNGgfGj9E3LmdOTM7Q95FhhsjlDTR6Pe7SOTR6CSOCiaXji15NIoyl7MojoxyOqPnwbO5hGfwSpU7LZYrC52sC6fD0x9juPw4fT2dwKur6_dX15Ob6dVln3S9_juRcWl4Dk7mXINdcZiULMP_giSGj1wTEDjtD6FDCvMoyOZRd-y97JSDNduBVBac4ehGGFiIggO_zXhpQUjI1LosBJMZh62wKx8q-PFw4HNwolLLUJ-hRYlPi6YmMLuHTv9W1pajwWC73faZh91XejkoKmUzeDd99eZy9qaH0PdmH2SBzILmfzqhMfF0B6xEZBlLEW_BtqA0sKXmKLOKkG-1sEIuYzBqYbdMc-8nF8ZqkTrbIq_Gifk3FZA-JpG4yQyms3kEP01m01ns_Xya3vxy9eEGPk2uryeXN9M3M7i6xmJdvp5SqfDpLUwuP8Ov08vXMXCkDkPx21JTFghVEK08rziccd6CsVAVLFPyTCxEBnUHwVJtuJaYFpRcr0XVawgy934KsRaWWX_2KDkfajCXc_mdkFnhcg5nmcvUIC2UWidYd8t1P3Ori7aOXWln7CBTTto-CR-Jcr7BEMmGZ1bp4yr8lmeOgCWlwqLtjmsZrC7HVus_xCAUFoWztT8eDOAL34HdlRVTHj9U-OfSGWIHFRKvcI6tYMfB7K2TGYGgDkld8QVlhmtLj0RUqdVG5NRcDY9ELvANKxyzHN6-v0YeLcfq0eMZxZBsjQKvfIOPiHGjRO4DJFWABH0k5CNpct052HwfgsWEFqRbJ5iA6c7lHRYNEPk0AEWUC6GNhRUrFpSlIQsVBL6CZEFeMuyCypctkYbaKbbJsP9ifERLNrV6wdJTB1AVaDRqlfqsZvnCI-nscY-bJnVROyTqp3wpZKcbe4s-lzn9PumGMMxRLmXitRBO02TcVEBDnCPnB93nB7y1lty7qQxaotrBAUaA8CTZYaR5A1IMdLM879Qo4hCp-1XWUqWKC9LEWeAK26mAx7BgheFt7o4ayqZhq1gNHyGV3xzXu0Z37EcLDqKNUM4Uu3AJsO33qaHhDe0Cs1KuyKEKB366W-14r1QGZ-uGg6ZLgMTcvL8-P2mygm1F68A8oiY-JF4X9lvg-jTtilnAQe4XFk1SLj2qB82PfmaCFlbzGhtviDc8ZQVOeJypObMMcKy4zDr0FTfcBC_8diVSYdG04vVB3jgLLoaQc5wAOc0cVSHBUqSIHFnxtqLOrZEVjZS26oqZFTe0hXOa7JTvcTrlnk4Z6JRP0LkoFNJFbBPYczoKh526vfxM77Sb7UiBmkfhxlIbdLswCA6r9qt6t4q7-BdxH2bSPPpKXFlnbGzufVs4O8OtTU2JSxt_1yT4Y6pbOF8czr0tBinQ1z25o-G4Rs47NMzrSbyfmLTJD8PST5KT4Q_D4XBMtR6cwc9ccu-7OvYN_NiDcWlS9-c5vveROQQP-ETyXi2ntmivurrZS7alq9naWv5lLecLRje32rm-7R549U1GbwGjUXM1NSZ7dXDXQHp_uK9HluiTu_NIhX5qmLfW7D-lMoKqgM2qwTds3HrLNjdVSGWitdq2AwfaiAdSqt4sGOmFt5j6FaMisJKkiyA8cDiu2XqUlTf5hjiVTTtcdbGOVi0OKvzW4og8M-IvntgLOs3ZaIT3jrM8MfiizpP68j1K62L8DQ3Wwv9Uez2ZSfjdlDf7LIa7-yPAWq34v_RiM-v_1omt3B-3pea4miQMaSbhVyN9i-HXmT58XEZyk2Unpy_cCYpVaasvz6iHaM6z589PXkKP6Wx1btbJyyH0ejh0bM93Bb7m9gq2Tv3naCHShs8sywo83FQfkHiAQ1d-ie7jWo691ZJjD0X3f_i_vwEK5UGw)) \ No newline at end of file From 8eb5b3cc7ee01ad9ee54fb1bf2d3d8a6ee74f780 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 00:01:20 +0000 Subject: [PATCH 08/29] Update bloom filter policies structure --- benchmarks/bloom_filter/add_bench.cu | 6 +- benchmarks/bloom_filter/contains_bench.cu | 10 +- examples/bloom_filter/host_bulk_example.cu | 2 +- include/cuco/arrow_bf_policy.cuh | 114 ------------------ include/cuco/bloom_filter.cuh | 9 +- ...r_policy.cuh => bloom_filter_policies.cuh} | 27 +++-- include/cuco/bloom_filter_ref.cuh | 4 +- .../detail/bloom_filter/arrow_bf_policy.inl | 52 -------- ...olicy_impl.cuh => arrow_filter_policy.cuh} | 67 ++++++++-- ...r_policy.inl => default_filter_policy.inl} | 21 ++-- ...mpl.cuh => default_filter_policy_impl.cuh} | 5 +- tests/bloom_filter/unique_sequence_test.cu | 21 ++-- 12 files changed, 114 insertions(+), 224 deletions(-) delete mode 100644 include/cuco/arrow_bf_policy.cuh rename include/cuco/{bloom_filter_policy.cuh => bloom_filter_policies.cuh} (81%) delete mode 100644 include/cuco/detail/bloom_filter/arrow_bf_policy.inl rename include/cuco/detail/bloom_filter/{arrow_bf_policy_impl.cuh => arrow_filter_policy.cuh} (55%) rename include/cuco/detail/bloom_filter/{bloom_filter_policy.inl => default_filter_policy.inl} (59%) rename include/cuco/detail/bloom_filter/{bloom_filter_policy_impl.cuh => default_filter_policy_impl.cuh} (95%) diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu index 00e2de775..409e08540 100644 --- a/benchmarks/bloom_filter/add_bench.cu +++ b/benchmarks/bloom_filter/add_bench.cu @@ -41,9 +41,9 @@ template , Dist>) { - using policy_type = cuco::bloom_filter_policy, - Word, - static_cast(WordsPerBlock)>; + using policy_type = cuco::default_filter_policy, + Word, + static_cast(WordsPerBlock)>; using filter_type = cuco::bloom_filter, cuda::thread_scope_device, policy_type>; diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index 8060cbf4a..7af9781f0 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/benchmarks/bloom_filter/contains_bench.cu @@ -43,9 +43,9 @@ void bloom_filter_contains( { // cudaDeviceSetLimit(cudaLimitMaxL2FetchGranularity, 32); // slightly improves peformance if // filter block fits into a 32B sector - using policy_type = cuco::bloom_filter_policy, - Word, - static_cast(WordsPerBlock)>; + using policy_type = cuco::default_filter_policy, + Word, + static_cast(WordsPerBlock)>; using filter_type = cuco::bloom_filter, cuda::thread_scope_device, policy_type>; @@ -90,14 +90,14 @@ void bloom_filter_contains( /** * @brief A benchmark evaluating `cuco::bloom_filter::contains_async` performance with - * `arrow_bf_policy` + * `arrow_filter_policy` */ template void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list) { // cudaDeviceSetLimit(cudaLimitMaxL2FetchGranularity, 32); // slightly improves peformance if // filter block fits into a 32B sector - using policy_type = cuco::arrow_bf_policy; + using policy_type = cuco::arrow_filter_policy; using filter_type = cuco::bloom_filter, cuda::thread_scope_device, policy_type>; diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu index b3881c29e..4364db2b3 100644 --- a/examples/bloom_filter/host_bulk_example.cu +++ b/examples/bloom_filter/host_bulk_example.cu @@ -79,7 +79,7 @@ int main(void) bulk_insert_and_eval_bloom_filter(filter, num_keys); // Arrow bloom filter policy type - using arrow_policy_type = cuco::arrow_bf_policy; + using arrow_policy_type = cuco::arrow_filter_policy; // bloom filter with arrow policy type using arrow_policy_filter_type = cuco:: bloom_filter, cuda::thread_scope_device, arrow_policy_type>; diff --git a/include/cuco/arrow_bf_policy.cuh b/include/cuco/arrow_bf_policy.cuh deleted file mode 100644 index 1af5df917..000000000 --- a/include/cuco/arrow_bf_policy.cuh +++ /dev/null @@ -1,114 +0,0 @@ -/* - * 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 - -#include - -namespace cuco { - -/** - * @brief A policy that defines how a Blocked Bloom Filter generates and stores a key's fingerprint. - * - * @note `Word` type must be an atomically updatable integral type. `WordsPerBlock` must - * be a power-of-two. - * - * @tparam Hash Hash function used to generate a key's fingerprint - * @tparam Word Underlying word/segment type of a filter block - * @tparam WordsPerBlock Number of words/segments in each block - */ -template -class arrow_bf_policy { - using impl_type = cuco::detail::arrow_bf_policy_impl; - - public: - using hasher = typename impl_type::hasher; ///< Type of the hash function - using hash_argument_type = typename impl_type::hash_argument_type; ///< Hash function input type - using hash_result_type = typename impl_type::hash_result_type; ///< hash function output type - using word_type = - typename impl_type::word_type; ///< Underlying word/segment type of a filter block - - static constexpr std::uint32_t words_per_block = - impl_type::words_per_block; ///< Number of words/segments in each filter block - - static constexpr std::uint32_t bits_set_per_block = - impl_type::bits_set_per_block; ///< Number of words/segments in each filter block - - public: - /** - * @brief Constructs the `arrow_bf_policy` object. - * - * @throws If number of filter blocks (`num_blocks`) is smaller than 1 - * or larger than 4194304. If called from host: throws exception; - * If called from device: Traps the kernel. - * - * @param num_blocks Number of bloom filter blocks - * @param hash Hash function used to generate a key's fingerprint - */ - __host__ __device__ constexpr arrow_bf_policy(std::uint32_t num_blocks, hasher hash = {}); - - /** - * @brief Generates the hash value for a given key. - * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * - * @param key The key to hash - * - * @return The hash value of the key - */ - __device__ constexpr hash_result_type hash(hash_argument_type const& key) const; - - /** - * @brief Determines the filter block a key is added into. - * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * - * @tparam Extent Size type that is used to determine the number of blocks in the filter - * - * @param hash Hash value of the key - * @param num_blocks Number of block in the filter - * - * @return The block index for the given key's hash value - */ - template - __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const; - - /** - * @brief Determines the fingerprint pattern for a word/segment within the filter block for a - * given key's hash value. - * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * - * @param hash Hash value of the key - * @param word_index Target word/segment within the filter block - * - * @return The bit pattern for the word/segment in the filter block - */ - __device__ constexpr word_type word_pattern(hash_result_type hash, - std::uint32_t word_index) const; - - private: - impl_type impl_; ///< Policy implementation -}; - -} // namespace cuco - -#include \ No newline at end of file diff --git a/include/cuco/bloom_filter.cuh b/include/cuco/bloom_filter.cuh index 6f26ab947..432934a16 100644 --- a/include/cuco/bloom_filter.cuh +++ b/include/cuco/bloom_filter.cuh @@ -16,8 +16,7 @@ #pragma once -#include -#include +#include #include #include #include @@ -56,13 +55,13 @@ namespace cuco { * @tparam Extent Size type that is used to determine the number of blocks in the filter * @tparam Scope The scope in which operations will be performed by individual threads * @tparam Policy Type that defines how to generate and store key fingerprints (see - * `cuco/bloom_filter_policy.cuh`) + * `cuco/default_filter_policy.cuh`) * @tparam Allocator Type of allocator used for device-accessible storage */ template , cuda::thread_scope Scope = cuda::thread_scope_device, - class Policy = cuco::bloom_filter_policy, std::uint32_t, 8>, + class Policy = cuco::default_filter_policy, std::uint32_t, 8>, class Allocator = cuco::cuda_allocator> class bloom_filter { public: @@ -110,7 +109,7 @@ class bloom_filter { * * @param num_blocks Number of sub-filters or blocks * @param scope The scope in which operations will be performed - * @param policy Fingerprint generation policy (see `cuco/bloom_filter_policy.cuh`) + * @param policy Fingerprint generation policy (see `cuco/default_filter_policy.cuh`) * @param alloc Allocator used for allocating device-accessible storage * @param stream CUDA stream used to initialize the filter */ diff --git a/include/cuco/bloom_filter_policy.cuh b/include/cuco/bloom_filter_policies.cuh similarity index 81% rename from include/cuco/bloom_filter_policy.cuh rename to include/cuco/bloom_filter_policies.cuh index 5b8f87d95..cf9ddb371 100644 --- a/include/cuco/bloom_filter_policy.cuh +++ b/include/cuco/bloom_filter_policies.cuh @@ -16,14 +16,25 @@ #pragma once -#include +#include +#include #include namespace cuco { /** - * @brief A policy that defines how a Blocked Bloom Filter generates and stores a key's fingerprint. + * @brief A policy that defines how Arrow Block-Split Bloom Filter generates and stores a key's + * fingerprint. + * + * @tparam Key The type of the values to generate a fingerprint for. + */ +template +using arrow_filter_policy = detail::arrow_filter_policy; + +/** + * @brief The default policy that defines how a Blocked Bloom Filter generates and stores a key's + * fingerprint. * * @note `Word` type must be an atomically updatable integral type. `WordsPerBlock` must * be a power-of-two. @@ -33,8 +44,8 @@ namespace cuco { * @tparam WordsPerBlock Number of words/segments in each block */ template -class bloom_filter_policy { - using impl_type = cuco::detail::bloom_filter_policy_impl; +class default_filter_policy { + using impl_type = cuco::detail::default_filter_policy_impl; public: using hasher = typename impl_type::hasher; ///< Type of the hash function @@ -48,7 +59,7 @@ class bloom_filter_policy { public: /** - * @brief Constructs the `bloom_filter_policy` object. + * @brief Constructs the `default_filter_policy` object. * * @throws Compile-time error if the specified number of words in a filter block is not a * power-of-two or is larger than 32. If called from host: throws exception; If called from @@ -64,8 +75,8 @@ class bloom_filter_policy { * @param pattern_bits Number of bits in a key's fingerprint * @param hash Hash function used to generate a key's fingerprint */ - __host__ __device__ constexpr bloom_filter_policy(std::uint32_t pattern_bits = words_per_block, - Hash hash = {}); + __host__ __device__ constexpr default_filter_policy(std::uint32_t pattern_bits = words_per_block, + Hash hash = {}); /** * @brief Generates the hash value for a given key. @@ -116,4 +127,4 @@ class bloom_filter_policy { } // namespace cuco -#include \ No newline at end of file +#include \ No newline at end of file diff --git a/include/cuco/bloom_filter_ref.cuh b/include/cuco/bloom_filter_ref.cuh index e4434845f..8b1e633af 100644 --- a/include/cuco/bloom_filter_ref.cuh +++ b/include/cuco/bloom_filter_ref.cuh @@ -33,7 +33,7 @@ namespace cuco { * @tparam Extent Size type that is used to determine the number of blocks in the filter * @tparam Scope The scope in which operations will be performed by individual threads * @tparam Policy Type that defines how to generate and store key fingerprints (see - * `cuco/bloom_filter_policy.cuh`) + * `cuco/bloom_filter_policies.cuh`) */ template class bloom_filter_ref { @@ -60,7 +60,7 @@ class bloom_filter_ref { * @param data Pointer to the storage span of the filter * @param num_blocks Number of sub-filters or blocks * @param scope The scope in which operations will be performed - * @param policy Fingerprint generation policy (see `cuco/bloom_filter_policy.cuh`) + * @param policy Fingerprint generation policy (see `cuco/default_filter_policy.cuh`) */ __host__ __device__ explicit constexpr bloom_filter_ref(word_type* data, Extent num_blocks, diff --git a/include/cuco/detail/bloom_filter/arrow_bf_policy.inl b/include/cuco/detail/bloom_filter/arrow_bf_policy.inl deleted file mode 100644 index cfdf3c84b..000000000 --- a/include/cuco/detail/bloom_filter/arrow_bf_policy.inl +++ /dev/null @@ -1,52 +0,0 @@ -/* - * 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 - -namespace cuco { - -template -__host__ __device__ constexpr arrow_bf_policy::arrow_bf_policy(std::uint32_t num_blocks, - hasher hash) - : impl_{num_blocks, hash} -{ -} - -template -__device__ constexpr typename arrow_bf_policy::hash_result_type arrow_bf_policy::hash( - typename arrow_bf_policy::hash_argument_type const& key) const -{ - return impl_.hash(key); -} - -template -template -__device__ constexpr auto arrow_bf_policy::block_index( - typename arrow_bf_policy::hash_result_type hash, Extent num_blocks) const -{ - return impl_.block_index(hash, num_blocks); -} - -template -__device__ constexpr typename arrow_bf_policy::word_type arrow_bf_policy::word_pattern( - arrow_bf_policy::hash_result_type hash, std::uint32_t word_index) const -{ - return impl_.word_pattern(hash, word_index); -} - -} // namespace cuco \ No newline at end of file diff --git a/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh similarity index 55% rename from include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh rename to include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 1964bc1f3..6a4c8251e 100644 --- a/include/cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -29,39 +29,70 @@ namespace cuco::detail { -template -class arrow_bf_policy_impl { +/** + * @brief A policy that defines how Arrow Block-Split Bloom Filter generates and stores a key's + * fingerprint. + * + * @tparam Key The type of the values to generate a fingerprint for. + */ +template +class arrow_filter_policy { public: - using hasher = cuco::xxhash_64; ///< xxhash_64 hasher for Arrow bloom filter policy - using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy + using hasher = cuco::xxhash_64; ///< xxhash_64 hasher for Arrow bloom filter policy + using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy using hash_argument_type = typename hasher::argument_type; using hash_result_type = decltype(std::declval()(std::declval())); - static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow bf block - static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow bf block + static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block + static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow filter block - __host__ __device__ explicit constexpr arrow_bf_policy_impl(std::uint32_t num_blocks, hasher hash) + __host__ __device__ constexpr arrow_filter_policy(std::uint32_t num_blocks, hasher hash = {}) : hash_{hash} { - constexpr std::uint32_t bytes_per_filter_block = 32; ///< Number of bytes in one Arrow bf block - constexpr std::uint32_t max_arrow_bf_bytes = + constexpr std::uint32_t bytes_per_filter_block = + 32; ///< Number of bytes in one Arrow filter block + constexpr std::uint32_t max_arrow_filter_bytes = 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter NV_DISPATCH_TARGET( NV_IS_HOST, - (CUCO_EXPECTS(num_blocks >= 1 and num_blocks <= (max_arrow_bf_bytes / bytes_per_filter_block), - "`num_blocks` must be in the range of [1, 4194304]");), + (CUCO_EXPECTS( + num_blocks >= 1 and num_blocks <= (max_arrow_filter_bytes / bytes_per_filter_block), + "`num_blocks` must be in the range of [1, 4194304]");), NV_IS_DEVICE, - (if (num_blocks < 1 or num_blocks > (max_arrow_bf_bytes / bytes_per_filter_block)) { + (if (num_blocks < 1 or num_blocks > (max_arrow_filter_bytes / bytes_per_filter_block)) { __trap(); // TODO this kills the kernel and corrupts the CUDA context. Not ideal. })); } + /** + * @brief Generates the hash value for a given key. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @param key The key to hash + * + * @return The hash value of the key + */ __device__ constexpr hash_result_type hash(hash_argument_type const& key) const { return hash_(key); } + /** + * @brief Determines the filter block a key is added into. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @tparam Extent Size type that is used to determine the number of blocks in the filter + * + * @param hash Hash value of the key + * @param num_blocks Number of block in the filter + * + * @return The block index for the given key's hash value + */ template __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const { @@ -69,6 +100,18 @@ class arrow_bf_policy_impl { return static_cast(((hash >> hash_bits) * num_blocks) >> hash_bits); } + /** + * @brief Determines the fingerprint pattern for a word/segment within the filter block for a + * given key's hash value. + * + * @note This function is meant as a customization point and is only used in the internals of the + * `bloom_filter(_ref)` implementation. + * + * @param hash Hash value of the key + * @param word_index Target word/segment within the filter block + * + * @return The bit pattern for the word/segment in the filter block + */ __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const { // Arrow's block-based bloom filter algorithm needs these eight odd SALT values to calculate diff --git a/include/cuco/detail/bloom_filter/bloom_filter_policy.inl b/include/cuco/detail/bloom_filter/default_filter_policy.inl similarity index 59% rename from include/cuco/detail/bloom_filter/bloom_filter_policy.inl rename to include/cuco/detail/bloom_filter/default_filter_policy.inl index f117798f9..eb8dbf703 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_policy.inl +++ b/include/cuco/detail/bloom_filter/default_filter_policy.inl @@ -21,33 +21,34 @@ namespace cuco { template -__host__ __device__ constexpr bloom_filter_policy::bloom_filter_policy( - uint32_t pattern_bits, Hash hash) +__host__ + __device__ constexpr default_filter_policy::default_filter_policy( + uint32_t pattern_bits, Hash hash) : impl_{pattern_bits, hash} { } template -__device__ constexpr typename bloom_filter_policy::hash_result_type -bloom_filter_policy::hash( - typename bloom_filter_policy::hash_argument_type const& key) const +__device__ constexpr typename default_filter_policy::hash_result_type +default_filter_policy::hash( + typename default_filter_policy::hash_argument_type const& key) const { return impl_.hash(key); } template template -__device__ constexpr auto bloom_filter_policy::block_index( - typename bloom_filter_policy::hash_result_type hash, +__device__ constexpr auto default_filter_policy::block_index( + typename default_filter_policy::hash_result_type hash, Extent num_blocks) const { return impl_.block_index(hash, num_blocks); } template -__device__ constexpr typename bloom_filter_policy::word_type -bloom_filter_policy::word_pattern( - bloom_filter_policy::hash_result_type hash, +__device__ constexpr typename default_filter_policy::word_type +default_filter_policy::word_pattern( + default_filter_policy::hash_result_type hash, std::uint32_t word_index) const { return impl_.word_pattern(hash, word_index); diff --git a/include/cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh b/include/cuco/detail/bloom_filter/default_filter_policy_impl.cuh similarity index 95% rename from include/cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh rename to include/cuco/detail/bloom_filter/default_filter_policy_impl.cuh index d7af7667c..ae2331b44 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh +++ b/include/cuco/detail/bloom_filter/default_filter_policy_impl.cuh @@ -29,7 +29,7 @@ namespace cuco::detail { template -class bloom_filter_policy_impl { +class default_filter_policy_impl { public: using hasher = Hash; using word_type = Word; @@ -43,7 +43,8 @@ class bloom_filter_policy_impl { static constexpr std::uint32_t bit_index_width = cuda::std::bit_width(word_bits - 1); public: - __host__ __device__ explicit constexpr bloom_filter_policy_impl(uint32_t pattern_bits, Hash hash) + __host__ __device__ explicit constexpr default_filter_policy_impl(uint32_t pattern_bits, + Hash hash) : pattern_bits_{pattern_bits}, min_bits_per_word_{pattern_bits_ / words_per_block}, remainder_bits_{pattern_bits_ % words_per_block}, diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 428d6e3f0..9a0771514 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -84,13 +84,14 @@ void test_unique_sequence(Filter& filter, size_type num_keys) // TODO test FPR but how? } -TEMPLATE_TEST_CASE_SIG("Unique sequence", - "", - ((class Key, class Policy), Key, Policy), - (int32_t, cuco::bloom_filter_policy, uint32_t, 1>), - (int32_t, cuco::bloom_filter_policy, uint32_t, 8>), - (int32_t, cuco::bloom_filter_policy, uint64_t, 1>), - (int32_t, cuco::bloom_filter_policy, uint64_t, 8>)) +TEMPLATE_TEST_CASE_SIG( + "Unique sequence", + "", + ((class Key, class Policy), Key, Policy), + (int32_t, cuco::default_filter_policy, uint32_t, 1>), + (int32_t, cuco::default_filter_policy, uint32_t, 8>), + (int32_t, cuco::default_filter_policy, uint64_t, 1>), + (int32_t, cuco::default_filter_policy, uint64_t, 8>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; @@ -104,11 +105,11 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence", test_unique_sequence(filter, num_keys); } -TEMPLATE_TEST_CASE_SIG("Unique sequence Arrow BF", +TEMPLATE_TEST_CASE_SIG("Unique sequence Arrow policy", "", ((class Key, class Policy), Key, Policy), - (int32_t, cuco::arrow_bf_policy), - (float, cuco::arrow_bf_policy)) + (int32_t, cuco::arrow_filter_policy), + (float, cuco::arrow_filter_policy)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; From 708aac379682cbea918dccab5c973a9993a8dd41 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Sat, 26 Oct 2024 00:01:32 +0000 Subject: [PATCH 09/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 39990b323..2578c46fc 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QbUXFqhgZ1aToKgKogluZIWoXbZvUhWDf97Z5ZLibTlNGgfGj9E3LmdOTM7Q95FhhsjlDTR6Pe7SOTR6CSOCiaXji15NIoyl7MojoxyOqPnwbO5hGfwSpU7LZYrC52sC6fD0x9juPw4fT2dwKur6_dX15Ob6dVln3S9_juRcWl4Dk7mXINdcZiULMP_giSGj1wTEDjtD6FDCvMoyOZRd-y97JSDNduBVBac4ehGGFiIggO_zXhpQUjI1LosBJMZh62wKx8q-PFw4HNwolLLUJ-hRYlPi6YmMLuHTv9W1pajwWC73faZh91XejkoKmUzeDd99eZy9qaH0PdmH2SBzILmfzqhMfF0B6xEZBlLEW_BtqA0sKXmKLOKkG-1sEIuYzBqYbdMc-8nF8ZqkTrbIq_Gifk3FZA-JpG4yQyms3kEP01m01ns_Xya3vxy9eEGPk2uryeXN9M3M7i6xmJdvp5SqfDpLUwuP8Ov08vXMXCkDkPx21JTFghVEK08rziccd6CsVAVLFPyTCxEBnUHwVJtuJaYFpRcr0XVawgy934KsRaWWX_2KDkfajCXc_mdkFnhcg5nmcvUIC2UWidYd8t1P3Ori7aOXWln7CBTTto-CR-Jcr7BEMmGZ1bp4yr8lmeOgCWlwqLtjmsZrC7HVus_xCAUFoWztT8eDOAL34HdlRVTHj9U-OfSGWIHFRKvcI6tYMfB7K2TGYGgDkld8QVlhmtLj0RUqdVG5NRcDY9ELvANKxyzHN6-v0YeLcfq0eMZxZBsjQKvfIOPiHGjRO4DJFWABH0k5CNpct052HwfgsWEFqRbJ5iA6c7lHRYNEPk0AEWUC6GNhRUrFpSlIQsVBL6CZEFeMuyCypctkYbaKbbJsP9ifERLNrV6wdJTB1AVaDRqlfqsZvnCI-nscY-bJnVROyTqp3wpZKcbe4s-lzn9PumGMMxRLmXitRBO02TcVEBDnCPnB93nB7y1lty7qQxaotrBAUaA8CTZYaR5A1IMdLM879Qo4hCp-1XWUqWKC9LEWeAK26mAx7BgheFt7o4ayqZhq1gNHyGV3xzXu0Z37EcLDqKNUM4Uu3AJsO33qaHhDe0Cs1KuyKEKB366W-14r1QGZ-uGg6ZLgMTcvL8-P2mygm1F68A8oiY-JF4X9lvg-jTtilnAQe4XFk1SLj2qB82PfmaCFlbzGhtviDc8ZQVOeJypObMMcKy4zDr0FTfcBC_8diVSYdG04vVB3jgLLoaQc5wAOc0cVSHBUqSIHFnxtqLOrZEVjZS26oqZFTe0hXOa7JTvcTrlnk4Z6JRP0LkoFNJFbBPYczoKh526vfxM77Sb7UiBmkfhxlIbdLswCA6r9qt6t4q7-BdxH2bSPPpKXFlnbGzufVs4O8OtTU2JSxt_1yT4Y6pbOF8czr0tBinQ1z25o-G4Rs47NMzrSbyfmLTJD8PST5KT4Q_D4XBMtR6cwc9ccu-7OvYN_NiDcWlS9-c5vveROQQP-ETyXi2ntmivurrZS7alq9naWv5lLecLRje32rm-7R549U1GbwGjUXM1NSZ7dXDXQHp_uK9HluiTu_NIhX5qmLfW7D-lMoKqgM2qwTds3HrLNjdVSGWitdq2AwfaiAdSqt4sGOmFt5j6FaMisJKkiyA8cDiu2XqUlTf5hjiVTTtcdbGOVi0OKvzW4og8M-IvntgLOs3ZaIT3jrM8MfiizpP68j1K62L8DQ3Wwv9Uez2ZSfjdlDf7LIa7-yPAWq34v_RiM-v_1omt3B-3pea4miQMaSbhVyN9i-HXmT58XEZyk2Unpy_cCYpVaasvz6iHaM6z589PXkKP6Wx1btbJyyH0ejh0bM93Bb7m9gq2Tv3naCHShs8sywo83FQfkHiAQ1d-ie7jWo691ZJjD0X3f_i_vwEK5UGw)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8nqvxXB7dwRYox3_l36sZ_3furGR-8PW1BzXk4Q-zSX8cqTvMfxC04cPzEhu0vTk9IU7QbEqbPn1GXUQzXn6_PnJS-gwna7OzTp52YdOBweP7fiuwFfdTs7Wc_9Jmot5zWeapjkebsqPSDzAwSu_RXdxJcfeasixh6K7P_zf3_9HQ24=)) \ No newline at end of file From 0872cfb48807f646e24a034cdc8d70bf294e9597 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 00:13:01 +0000 Subject: [PATCH 10/29] doc updates --- .../bloom_filter/arrow_filter_policy.cuh | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 6a4c8251e..ebff70328 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -46,6 +46,16 @@ class arrow_filter_policy { static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow filter block + /** + * @brief Constructs the `arrow_filter_policy` object. + * + * @throws If number of filter blocks (`num_blocks`) is smaller than 1 + * or larger than 4194304. If called from host: throws exception; + * If called from device: Traps the kernel. + * + * @param num_blocks Number of bloom filter blocks + * @param hash Hash function used to generate a key's fingerprint + */ __host__ __device__ constexpr arrow_filter_policy(std::uint32_t num_blocks, hasher hash = {}) : hash_{hash} { @@ -68,9 +78,6 @@ class arrow_filter_policy { /** * @brief Generates the hash value for a given key. * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * * @param key The key to hash * * @return The hash value of the key @@ -83,9 +90,6 @@ class arrow_filter_policy { /** * @brief Determines the filter block a key is added into. * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * * @tparam Extent Size type that is used to determine the number of blocks in the filter * * @param hash Hash value of the key @@ -104,9 +108,6 @@ class arrow_filter_policy { * @brief Determines the fingerprint pattern for a word/segment within the filter block for a * given key's hash value. * - * @note This function is meant as a customization point and is only used in the internals of the - * `bloom_filter(_ref)` implementation. - * * @param hash Hash value of the key * @param word_index Target word/segment within the filter block * From a0c7d32a88f4825fffceeea5d05578c7f7bebfc0 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 00:37:22 +0000 Subject: [PATCH 11/29] Minor updates --- README.md | 2 +- .../detail/bloom_filter/arrow_filter_policy.cuh | 14 +++++++------- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/README.md b/README.md index 39990b323..2578c46fc 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QbUXFqhgZ1aToKgKogluZIWoXbZvUhWDf97Z5ZLibTlNGgfGj9E3LmdOTM7Q95FhhsjlDTR6Pe7SOTR6CSOCiaXji15NIoyl7MojoxyOqPnwbO5hGfwSpU7LZYrC52sC6fD0x9juPw4fT2dwKur6_dX15Ob6dVln3S9_juRcWl4Dk7mXINdcZiULMP_giSGj1wTEDjtD6FDCvMoyOZRd-y97JSDNduBVBac4ehGGFiIggO_zXhpQUjI1LosBJMZh62wKx8q-PFw4HNwolLLUJ-hRYlPi6YmMLuHTv9W1pajwWC73faZh91XejkoKmUzeDd99eZy9qaH0PdmH2SBzILmfzqhMfF0B6xEZBlLEW_BtqA0sKXmKLOKkG-1sEIuYzBqYbdMc-8nF8ZqkTrbIq_Gifk3FZA-JpG4yQyms3kEP01m01ns_Xya3vxy9eEGPk2uryeXN9M3M7i6xmJdvp5SqfDpLUwuP8Ov08vXMXCkDkPx21JTFghVEK08rziccd6CsVAVLFPyTCxEBnUHwVJtuJaYFpRcr0XVawgy934KsRaWWX_2KDkfajCXc_mdkFnhcg5nmcvUIC2UWidYd8t1P3Ori7aOXWln7CBTTto-CR-Jcr7BEMmGZ1bp4yr8lmeOgCWlwqLtjmsZrC7HVus_xCAUFoWztT8eDOAL34HdlRVTHj9U-OfSGWIHFRKvcI6tYMfB7K2TGYGgDkld8QVlhmtLj0RUqdVG5NRcDY9ELvANKxyzHN6-v0YeLcfq0eMZxZBsjQKvfIOPiHGjRO4DJFWABH0k5CNpct052HwfgsWEFqRbJ5iA6c7lHRYNEPk0AEWUC6GNhRUrFpSlIQsVBL6CZEFeMuyCypctkYbaKbbJsP9ifERLNrV6wdJTB1AVaDRqlfqsZvnCI-nscY-bJnVROyTqp3wpZKcbe4s-lzn9PumGMMxRLmXitRBO02TcVEBDnCPnB93nB7y1lty7qQxaotrBAUaA8CTZYaR5A1IMdLM879Qo4hCp-1XWUqWKC9LEWeAK26mAx7BgheFt7o4ayqZhq1gNHyGV3xzXu0Z37EcLDqKNUM4Uu3AJsO33qaHhDe0Cs1KuyKEKB366W-14r1QGZ-uGg6ZLgMTcvL8-P2mygm1F68A8oiY-JF4X9lvg-jTtilnAQe4XFk1SLj2qB82PfmaCFlbzGhtviDc8ZQVOeJypObMMcKy4zDr0FTfcBC_8diVSYdG04vVB3jgLLoaQc5wAOc0cVSHBUqSIHFnxtqLOrZEVjZS26oqZFTe0hXOa7JTvcTrlnk4Z6JRP0LkoFNJFbBPYczoKh526vfxM77Sb7UiBmkfhxlIbdLswCA6r9qt6t4q7-BdxH2bSPPpKXFlnbGzufVs4O8OtTU2JSxt_1yT4Y6pbOF8czr0tBinQ1z25o-G4Rs47NMzrSbyfmLTJD8PST5KT4Q_D4XBMtR6cwc9ccu-7OvYN_NiDcWlS9-c5vveROQQP-ETyXi2ntmivurrZS7alq9naWv5lLecLRje32rm-7R549U1GbwGjUXM1NSZ7dXDXQHp_uK9HluiTu_NIhX5qmLfW7D-lMoKqgM2qwTds3HrLNjdVSGWitdq2AwfaiAdSqt4sGOmFt5j6FaMisJKkiyA8cDiu2XqUlTf5hjiVTTtcdbGOVi0OKvzW4og8M-IvntgLOs3ZaIT3jrM8MfiizpP68j1K62L8DQ3Wwv9Uez2ZSfjdlDf7LIa7-yPAWq34v_RiM-v_1omt3B-3pea4miQMaSbhVyN9i-HXmT58XEZyk2Unpy_cCYpVaasvz6iHaM6z589PXkKP6Wx1btbJyyH0ejh0bM93Bb7m9gq2Tv3naCHShs8sywo83FQfkHiAQ1d-ie7jWo691ZJjD0X3f_i_vwEK5UGw)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8nqvxXB7dwRYox3_l36sZ_3furGR-8PW1BzXk4Q-zSX8cqTvMfxC04cPzEhu0vTk9IU7QbEqbPn1GXUQzXn6_PnJS-gwna7OzTp52YdOBweP7fiuwFfdTs7Wc_9Jmot5zWeapjkebsqPSDzAwSu_RXdxJcfeasixh6K7P_zf3_9HQ24=)) \ No newline at end of file diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index ebff70328..1f81232a6 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -43,9 +43,14 @@ class arrow_filter_policy { using hash_argument_type = typename hasher::argument_type; using hash_result_type = decltype(std::declval()(std::declval())); - static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block + static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow filter block + static constexpr std::uint32_t bytes_per_filter_block = + 32; ///< Number of bytes in one Arrow filter block + static constexpr std::uint32_t max_arrow_filter_bytes = + 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter + /** * @brief Constructs the `arrow_filter_policy` object. * @@ -59,16 +64,11 @@ class arrow_filter_policy { __host__ __device__ constexpr arrow_filter_policy(std::uint32_t num_blocks, hasher hash = {}) : hash_{hash} { - constexpr std::uint32_t bytes_per_filter_block = - 32; ///< Number of bytes in one Arrow filter block - constexpr std::uint32_t max_arrow_filter_bytes = - 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter - NV_DISPATCH_TARGET( NV_IS_HOST, (CUCO_EXPECTS( num_blocks >= 1 and num_blocks <= (max_arrow_filter_bytes / bytes_per_filter_block), - "`num_blocks` must be in the range of [1, 4194304]");), + "The `num_blocks` in Arrow filter policy must be in the range of [1, 4194304]");), NV_IS_DEVICE, (if (num_blocks < 1 or num_blocks > (max_arrow_filter_bytes / bytes_per_filter_block)) { __trap(); // TODO this kills the kernel and corrupts the CUDA context. Not ideal. From 05b69679b24e162577be1c3aadb9e22dd1e8b970 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 01:14:59 +0000 Subject: [PATCH 12/29] Minor --- tests/bloom_filter/unique_sequence_test.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 9a0771514..e13a89697 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -85,7 +85,7 @@ void test_unique_sequence(Filter& filter, size_type num_keys) } TEMPLATE_TEST_CASE_SIG( - "Unique sequence", + "Unique sequence with default policy", "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::default_filter_policy, uint32_t, 1>), @@ -105,10 +105,11 @@ TEMPLATE_TEST_CASE_SIG( test_unique_sequence(filter, num_keys); } -TEMPLATE_TEST_CASE_SIG("Unique sequence Arrow policy", +TEMPLATE_TEST_CASE_SIG("Unique sequence with arrow policy", "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::arrow_filter_policy), + (uint64_t, cuco::arrow_filter_policy), (float, cuco::arrow_filter_policy)) { using filter_type = From e389a7b620db0e0e4049d376ac5473080a63fb70 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 01:38:58 +0000 Subject: [PATCH 13/29] Minor updates --- include/cuco/detail/bloom_filter/arrow_filter_policy.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 1f81232a6..3e442a768 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -100,8 +100,8 @@ class arrow_filter_policy { template __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const { - constexpr auto hash_bits = cuda::std::numeric_limits::digits; - return static_cast(((hash >> hash_bits) * num_blocks) >> hash_bits); + constexpr auto hash_bits = cuda::std::numeric_limits::digits; + return static_cast(((hash >> hash_bits) * num_blocks) >> hash_bits); } /** From 8c536bdb97fe7aef30ce04ac322606b7bc325a2d Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 02:14:57 +0000 Subject: [PATCH 14/29] Minor --- examples/bloom_filter/host_bulk_example.cu | 2 +- .../bloom_filter/arrow_filter_policy.cuh | 34 +++++++++---------- tests/bloom_filter/unique_sequence_test.cu | 4 +-- 3 files changed, 19 insertions(+), 21 deletions(-) diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu index 4364db2b3..a8748db55 100644 --- a/examples/bloom_filter/host_bulk_example.cu +++ b/examples/bloom_filter/host_bulk_example.cu @@ -85,7 +85,7 @@ int main(void) bloom_filter, cuda::thread_scope_device, arrow_policy_type>; // Spawn a bloom filter with arrow policy and 200 sub-filters. - arrow_policy_filter_type filter_arrow_policy{sub_filters, {}, arrow_policy_type{sub_filters}}; + arrow_policy_filter_type filter_arrow_policy{sub_filters}; // bulk insert to the bloom filter and evaluate std::cout << "Bulk insert and evaluate bloom filter with arrow policy: " << std::endl; diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 3e442a768..cbc647016 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -50,30 +51,20 @@ class arrow_filter_policy { 32; ///< Number of bytes in one Arrow filter block static constexpr std::uint32_t max_arrow_filter_bytes = 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter - + static constexpr std::uint32_t max_filter_blocks = + (max_arrow_filter_bytes / + bytes_per_filter_block); ///< Max sub-filter blocks allowed in Arrow bloom filter /** * @brief Constructs the `arrow_filter_policy` object. * - * @throws If number of filter blocks (`num_blocks`) is smaller than 1 - * or larger than 4194304. If called from host: throws exception; - * If called from device: Traps the kernel. + * @note The number of filter blocks with Arrow policy must be in the + * range of [1, 4194304]. If the bloom filter is constructed with a larger + * number of blocks, only the first 4194304 (128MB) blocks will be used. * * @param num_blocks Number of bloom filter blocks * @param hash Hash function used to generate a key's fingerprint */ - __host__ __device__ constexpr arrow_filter_policy(std::uint32_t num_blocks, hasher hash = {}) - : hash_{hash} - { - NV_DISPATCH_TARGET( - NV_IS_HOST, - (CUCO_EXPECTS( - num_blocks >= 1 and num_blocks <= (max_arrow_filter_bytes / bytes_per_filter_block), - "The `num_blocks` in Arrow filter policy must be in the range of [1, 4194304]");), - NV_IS_DEVICE, - (if (num_blocks < 1 or num_blocks > (max_arrow_filter_bytes / bytes_per_filter_block)) { - __trap(); // TODO this kills the kernel and corrupts the CUDA context. Not ideal. - })); - } + __host__ __device__ constexpr arrow_filter_policy(hasher hash = {}) : hash_{hash} {} /** * @brief Generates the hash value for a given key. @@ -90,6 +81,10 @@ class arrow_filter_policy { /** * @brief Determines the filter block a key is added into. * + * @note The number of filter blocks with Arrow policy must be in the + * range of [1, 4194304]. Passing a larger `num_blocks` will still + * upperbound the number of blocks used to the mentioned range. + * * @tparam Extent Size type that is used to determine the number of blocks in the filter * * @param hash Hash value of the key @@ -101,7 +96,10 @@ class arrow_filter_policy { __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const { constexpr auto hash_bits = cuda::std::numeric_limits::digits; - return static_cast(((hash >> hash_bits) * num_blocks) >> hash_bits); + // TODO: assert if num_blocks > max_filter_blocks + auto const max_blocks = cuda::std::min(num_blocks, max_filter_blocks); + // Make sure we are only contained withing the `max_filter_blocks` blocks + return static_cast(((hash >> hash_bits) * max_blocks) >> hash_bits) % max_blocks; } /** diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index e13a89697..7069f4c6a 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -57,7 +57,7 @@ void test_unique_sequence(Filter& filter, size_type num_keys) REQUIRE(cuco::test::all_of(contained.begin(), contained.end(), thrust::identity{})); } - SECTION("After clearing the flter no keys should be contained.") + SECTION("After clearing the filter no keys should be contained.") { filter.clear(); filter.contains(keys.begin(), keys.end(), contained.begin()); @@ -116,7 +116,7 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence with arrow policy", cuco::bloom_filter, cuda::thread_scope_device, Policy>; constexpr size_type num_keys{400}; - auto filter = filter_type{1000, {}, {1000}}; + auto filter = filter_type{1000}; test_unique_sequence(filter, num_keys); } From f61e79cb6b903458a03f29b9910dbb98fbed7093 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Sat, 26 Oct 2024 02:15:08 +0000 Subject: [PATCH 15/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 2578c46fc..fd4bc68ec 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8nqvxXB7dwRYox3_l36sZ_3furGR-8PW1BzXk4Q-zSX8cqTvMfxC04cPzEhu0vTk9IU7QbEqbPn1GXUQzXn6_PnJS-gwna7OzTp52YdOBweP7fiuwFfdTs7Wc_9Jmot5zWeapjkebsqPSDzAwSu_RXdxJcfeasixh6K7P_zf3_9HQ24=)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8v-91-oZ_bdOa-T1sO00x9UjoU8zB78K6VsLv7704eMxkps0PTl94U5QrApbfllGHURznj5_fvISOkynq3OzTl72odPBoWI7vuL4GtvJ2XruPzdzMa_5TNM0x8NN-YGIBzhU5bfoLq7k2DcNOfZHdPeH__sbf0s1CQ==)) \ No newline at end of file From e587a21ff9d3e3fe7fdda881c25e26fb0bf1fec7 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 26 Oct 2024 02:16:46 +0000 Subject: [PATCH 16/29] Update example link --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 2578c46fc..fd4bc68ec 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8nqvxXB7dwRYox3_l36sZ_3furGR-8PW1BzXk4Q-zSX8cqTvMfxC04cPzEhu0vTk9IU7QbEqbPn1GXUQzXn6_PnJS-gwna7OzTp52YdOBweP7fiuwFfdTs7Wc_9Jmot5zWeapjkebsqPSDzAwSu_RXdxJcfeasixh6K7P_zf3_9HQ24=)) \ No newline at end of file +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8v-91-oZ_bdOa-T1sO00x9UjoU8zB78K6VsLv7704eMxkps0PTl94U5QrApbfllGHURznj5_fvISOkynq3OzTl72odPBoWI7vuL4GtvJ2XruPzdzMa_5TNM0x8NN-YGIBzhU5bfoLq7k2DcNOfZHdPeH__sbf0s1CQ==)) \ No newline at end of file From 8d54b164c00218aa1445d0413d50bf62b24ef27b Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 10:12:06 -0700 Subject: [PATCH 17/29] Apply suggestions from code review Co-authored-by: Yunsong Wang --- include/cuco/bloom_filter.cuh | 4 ++-- include/cuco/bloom_filter_ref.cuh | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/include/cuco/bloom_filter.cuh b/include/cuco/bloom_filter.cuh index 432934a16..b05e9469e 100644 --- a/include/cuco/bloom_filter.cuh +++ b/include/cuco/bloom_filter.cuh @@ -55,7 +55,7 @@ namespace cuco { * @tparam Extent Size type that is used to determine the number of blocks in the filter * @tparam Scope The scope in which operations will be performed by individual threads * @tparam Policy Type that defines how to generate and store key fingerprints (see - * `cuco/default_filter_policy.cuh`) + * `cuco/bloom_filter_policies.cuh`) * @tparam Allocator Type of allocator used for device-accessible storage */ template Date: Mon, 28 Oct 2024 10:14:09 -0700 Subject: [PATCH 18/29] Apply suggestion from code review Co-authored-by: Yunsong Wang --- include/cuco/detail/bloom_filter/arrow_filter_policy.cuh | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index cbc647016..d307d534f 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -16,14 +16,11 @@ #pragma once -#include #include -#include #include +#include #include -#include -#include #include #include From db6b8f0a73002ca3ef9d3e70d3dfb9c6a89b48e2 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 18:08:35 +0000 Subject: [PATCH 19/29] Doxygen fix. Benchmarks --- benchmarks/bloom_filter/add_bench.cu | 54 +++++++++++++++++++ benchmarks/bloom_filter/contains_bench.cu | 11 ++-- .../bloom_filter/arrow_filter_policy.cuh | 2 +- 3 files changed, 60 insertions(+), 7 deletions(-) diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu index 409e08540..8e4d0b762 100644 --- a/benchmarks/bloom_filter/add_bench.cu +++ b/benchmarks/bloom_filter/add_bench.cu @@ -83,6 +83,51 @@ void bloom_filter_add(nvbench::state& state, }); } +/** + * @brief A benchmark evaluating `cuco::bloom_filter::add_async` performance with + * `arrow_filter_policy` + */ +template +void arrow_bloom_filter_add(nvbench::state& state, nvbench::type_list) +{ + using policy_type = cuco::arrow_filter_policy; + using filter_type = + cuco::bloom_filter, cuda::thread_scope_device, policy_type>; + + auto const num_keys = state.get_int64("NumInputs"); + auto const filter_size_mb = state.get_int64("FilterSizeMB"); + + std::size_t const num_sub_filters = + (filter_size_mb * 1024 * 1024) / + (sizeof(typename filter_type::word_type) * filter_type::words_per_block); + + if (num_sub_filters > policy_type::max_filter_blocks) { + state.skip("bloom filter with arrow policy should have <= 4194304 blocks"); // skip invalid + // configurations + } + + thrust::device_vector keys(num_keys); + + key_generator gen; + gen.generate(dist_from_state(state), keys.begin(), keys.end()); + + state.add_element_count(num_keys); + + filter_type filter{num_sub_filters}; + + state.collect_dram_throughput(); + state.collect_l1_hit_rates(); + state.collect_l2_hit_rates(); + state.collect_loads_efficiency(); + state.collect_stores_efficiency(); + + add_fpr_summary(state, filter); + + state.exec([&](nvbench::launch& launch) { + filter.add_async(keys.begin(), keys.end(), {launch.get_stream()}); + }); +} + NVBENCH_BENCH_TYPES(bloom_filter_add, NVBENCH_TYPE_AXES(nvbench::type_list, nvbench::type_list, @@ -118,3 +163,12 @@ NVBENCH_BENCH_TYPES(bloom_filter_add, .set_max_noise(defaults::MAX_NOISE) .add_int64_axis("NumInputs", {defaults::BF_N}) .add_int64_axis("FilterSizeMB", {defaults::BF_SIZE_MB}); + +NVBENCH_BENCH_TYPES(arrow_bloom_filter_add, + NVBENCH_TYPE_AXES(nvbench::type_list, + nvbench::type_list)) + .set_name("arrow_bloom_filter_add_unique_size") + .set_type_axes_names({"Key", "Distribution"}) + .set_max_noise(defaults::MAX_NOISE) + .add_int64_axis("NumInputs", {defaults::BF_N}) + .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); \ No newline at end of file diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index 7af9781f0..bf81f5d83 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/benchmarks/bloom_filter/contains_bench.cu @@ -108,10 +108,9 @@ void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list(num_sub_filters)}; - } catch (std::exception const& e) { - state.skip(e.what()); // skip invalid configurations + if (num_sub_filters > policy_type::max_filter_blocks) { + state.skip("bloom filter with arrow policy should have <= 4194304 blocks"); // skip invalid + // configurations } thrust::device_vector keys(num_keys); @@ -122,7 +121,7 @@ void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list(num_sub_filters)}}; + filter_type filter{num_sub_filters}; state.collect_dram_throughput(); state.collect_l1_hit_rates(); @@ -181,5 +180,5 @@ NVBENCH_BENCH_TYPES(arrow_bloom_filter_contains, .set_name("arrow_bloom_filter_contains_unique_size") .set_type_axes_names({"Key", "Distribution"}) .set_max_noise(defaults::MAX_NOISE) - .add_int64_axis("NumInputs", {defaults::BF_N / 2}) + .add_int64_axis("NumInputs", {defaults::BF_N}) .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); \ No newline at end of file diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index d307d534f..ca0ef5a73 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -18,6 +18,7 @@ #include +#include #include #include #include @@ -58,7 +59,6 @@ class arrow_filter_policy { * range of [1, 4194304]. If the bloom filter is constructed with a larger * number of blocks, only the first 4194304 (128MB) blocks will be used. * - * @param num_blocks Number of bloom filter blocks * @param hash Hash function used to generate a key's fingerprint */ __host__ __device__ constexpr arrow_filter_policy(hasher hash = {}) : hash_{hash} {} From 59f97911b31a1193c444789a44fd02b2730bbd0c Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 18:36:23 +0000 Subject: [PATCH 20/29] Separate the two examples. --- examples/CMakeLists.txt | 1 + .../host_bulk_arrow_policy_example.cu | 80 +++++++++++++++++++ examples/bloom_filter/host_bulk_example.cu | 49 ++++-------- 3 files changed, 95 insertions(+), 35 deletions(-) create mode 100644 examples/bloom_filter/host_bulk_arrow_policy_example.cu diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 8e7be947a..d29052028 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -47,3 +47,4 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/ 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") +ConfigureExample(BLOOM_FILTER_ARROW_POLICY_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/host_bulk_arrow_policy_example.cu") diff --git a/examples/bloom_filter/host_bulk_arrow_policy_example.cu b/examples/bloom_filter/host_bulk_arrow_policy_example.cu new file mode 100644 index 000000000..b59b730fd --- /dev/null +++ b/examples/bloom_filter/host_bulk_arrow_policy_example.cu @@ -0,0 +1,80 @@ +/* + * 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 + +#include +#include +#include +#include + +#include + +int main(void) +{ + int constexpr num_keys = 10'000; ///< Generate 10'000 keys + int constexpr num_tp = num_keys * 0.5; ///< Insert the first half keys into the filter. + int constexpr num_tn = num_keys - num_tp; + int constexpr sub_filters = 200; ///< 200 sub-filters per bloom filter + + // key type for bloom filter + using key_type = int; + + // We will use the Arrow filter policy for bloom filter fingerprint generation + using policy_type = cuco::arrow_filter_policy; + // Bloom filter type with Arrow filter policy + using filter_type = + cuco::bloom_filter, cuda::thread_scope_device, policy_type>; + + // Spawn a bloom filter with arrow policy and 200 sub-filters. + filter_type filter{sub_filters}; + + std::cout << "Bulk insert into bloom filter with Arrow fingerprint generation policy: " + << std::endl; + + thrust::device_vector keys(num_keys); + thrust::sequence(keys.begin(), keys.end(), 1); + + auto tp_begin = keys.begin(); + auto tp_end = tp_begin + num_tp; + auto tn_begin = tp_end; + auto tn_end = keys.end(); + + // Insert the first half of the keys. + filter.add(tp_begin, tp_end); + + thrust::device_vector tp_result(num_tp, false); + thrust::device_vector tn_result(num_keys - num_tp, false); + + // Query the filter for the previously inserted keys. + // This should result in a true-positive rate of TPR=1. + filter.contains(tp_begin, tp_end, tp_result.begin()); + + // Query the filter for the keys that are not present in the filter. + // Since bloom filters are probalistic data structures, the filter + // exhibits a false-positive rate FPR>0 depending on the number of bits in + // the filter and the number of hashes used per key. + filter.contains(tn_begin, tn_end, tn_result.begin()); + + float tp_rate = + float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); + float fp_rate = + float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); + + std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; + + return 0; +} diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu index a8748db55..f02f6e657 100644 --- a/examples/bloom_filter/host_bulk_example.cu +++ b/examples/bloom_filter/host_bulk_example.cu @@ -23,16 +23,21 @@ #include -// key type for bloom filter -using key_type = int; - -// Function to bulk insert to the provided bloom filter and evaluate FPR. -template -void bulk_insert_and_eval_bloom_filter(FilterType& filter, int num_keys) +int main(void) { - // Insert the first half keys into the filter. - int const num_tp = num_keys * 0.5; - int const num_tn = num_keys - num_tp; + int constexpr num_keys = 10'000; ///< Generate 10'000 keys + int constexpr num_tp = num_keys * 0.5; ///< Insert the first half keys into the filter. + int constexpr num_tn = num_keys - num_tp; + int constexpr sub_filters = 200; ///< 200 sub-filters per bloom filter + + // key type for bloom filter + using key_type = int; + + // Spawn a bloom filter with default policy and 200 sub-filters. + cuco::bloom_filter filter{sub_filters}; + + std::cout << "Bulk insert into bloom filter with default fingerprint generation policy: " + << std::endl; thrust::device_vector keys(num_keys); thrust::sequence(keys.begin(), keys.end(), 1); @@ -64,32 +69,6 @@ void bulk_insert_and_eval_bloom_filter(FilterType& filter, int num_keys) float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; -} - -int main(void) -{ - int constexpr num_keys = 10'000; ///< Generate 10'000 keys - int constexpr sub_filters = 200; ///< 200 sub-filters per bloom filter - - // Spawn a bloom filter with default policy and 200 sub-filters. - cuco::bloom_filter filter{sub_filters}; - - // bulk insert to the bloom filter and evaluate - std::cout << "Bulk insert and evaluate bloom filter with default policy: " << std::endl; - bulk_insert_and_eval_bloom_filter(filter, num_keys); - - // Arrow bloom filter policy type - using arrow_policy_type = cuco::arrow_filter_policy; - // bloom filter with arrow policy type - using arrow_policy_filter_type = cuco:: - bloom_filter, cuda::thread_scope_device, arrow_policy_type>; - - // Spawn a bloom filter with arrow policy and 200 sub-filters. - arrow_policy_filter_type filter_arrow_policy{sub_filters}; - - // bulk insert to the bloom filter and evaluate - std::cout << "Bulk insert and evaluate bloom filter with arrow policy: " << std::endl; - bulk_insert_and_eval_bloom_filter(filter_arrow_policy, num_keys); return 0; } \ No newline at end of file From 26ef2a8c3c766f78601e490783cadd2814b15610 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 11:40:19 -0700 Subject: [PATCH 21/29] Update README with both example links --- README.md | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index fd4bc68ec..689b96d70 100644 --- a/README.md +++ b/README.md @@ -254,4 +254,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzFV9tu20YQ_ZUBC7RSQl1sNAgg2QaUWyM0kFPLSRBUBbEiV9Ii1C67F8mq4X_vzHIpkbacGu1D44eIO7czZ2ZnyNvIcGOEkiYa_H4biSwanMRRzuTSsSWPBlHqMhbFkVFOp_TcezaT8Axeq2KnxXJloZW24bR_-nMMk8_jN-MRvL68-nh5NboeX066pOv1P4iUS8MzcDLjGuyKw6hgKf4XJDF85pqAwGm3Dy1SmEVBNovaQ-9lpxys2Q6ksuAMRzfCwELkHPhNygsLQkKq1kUumEw5bIVd-VDBj4cDX4MTNbcM9RlaFPi0qGsCs3vo9G9lbTHo9bbbbZd52F2ll728VDa9D-PXbyfTtx2Evjf7JHNkFjT_0wmNic93wApElrI54s3ZFpQGttQcZVYR8q0WVshlDEYt7JZp7v1kwlgt5s42yKtwYv51BaSPSSRuNIXxdBbBq9F0PI29ny_j6_eXn67hy-jqajS5Hr-dwuUVFmvyZkylwqd3MJp8hV_HkzcxcKQOQ_GbQlMWCFUQrTwrOZxy3oCxUCUsU_BULEQKVQfBUm24lpgWFFyvRdlrCDLzfnKxFpZZf_YgOR-qN5Mz-YOQae4yDmepS1Vvniu1TrDulutu6lYXTR270s7YXqqctF0SPhBlfIMhkg1PrdLHVfgNTx0BSwqFRdsd1zJYXY6t1r2PQSgsCmdrf9zrwTe-A7srSqY8fijxz6QzxA4qJF7hHFvBDoPZOydTAkEdMnf5N5QZri09ElGFVhuRUXPVPBK5wDcsd8xyePfxCnm0HKtHj2cUQ7I1CrzyNT4ixo0SmQ-QlAES9JGQj6TOdetg82MIFhNakG6dYAKmPZO3WDRA5OMAFFEuhDYWVixfUJaGLFQQ-AqSBXlJsQtKX7ZAGiqn2Cb97ovhES1Z1-oES08dQFmgwaBR6rOK5QuPpLXHPaybVEVtkag750shW-3YW3S5zOj3STuEYY5yKRKvhXDqJsO6AhriHDk_6D4_4K205N5NadAQVQ4OMAKER8kOI80bkGKgm2VZq0IRh0jt77I2Vyq_IE2cBS63rRJ4DAuWG97k7qihrBs2ilXzEVL5zXG9q3XHfrTgINoI5Uy-C5cA236fGhpe0y4wK-XyDMpw4Ke71Y53CmVwtm44aLoESMz1x6vzkzor2Fa0DswDauJD4lVhnwLXp2lXzAIOcr-waJJy6VHda370MxW0sOrX2HhDvOFzluOEx5maMcsAx4pLrUNfcc1N8MJvVmIuLJqWvN7LG2fBRR8yjhMgo5mjSiRYijkiR1a8rahyq2VFI6WpumJmxQ1t4YwmO-V7nE65p1MGOuUjdC5yhXQR2wT2nI7CYatqLz_TW81mO1Kg-lG4sdQG7Tb0gsOy_creLeMu_kXc-5nUj74TV1YZG5t53xbOznBrU1Pi0sbfFQn-mOoWzheHc2-LQXL0dUfuaDiukfMWDfNqEu8nJm3yw7D0k-Sk_1O_3x9SrXtn8AuX3Psuj30DP_Rg3Dyp-vMc3_vIHIIHfCJ5p5JTWzRXXdXsBdvS1WxsLf-ylvEFo5tb7lzfdve8-iajt4DBoL6aapO9PLitIb073NcjS_TR3XmkQq9q5o01-0-pDKAsYL1q8ISNW23Z-qYKqYy0Vttm4EAb8UBK5ZsFI73wFlO9YpQElpLSNigceBxWjD3IzJs9IVZw3AhZXq6jlYuDCr-xOCbPjPiLJ_aCTjM2GODd4yxLDL6s86S6gA9Suxg-ocka-B9rsUczCb_r8v-91-oZ_bdOa-T1sO00x9UjoU8zB78K6VsLv7704eMxkps0PTl94U5QrApbfllGHURznj5_fvISOkynq3OzTl72odPBoWI7vuL4GtvJ2XruPzdzMa_5TNM0x8NN-YGIBzhU5bfoLq7k2DcNOfZHdPeH__sbf0s1CQ==)) \ No newline at end of file +- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) + +#### Examples: +- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV22PGjcQ_iuj7YdCbnk7NYq03J3EXZIWNeJSuCSKSoWMdwAri721vXD0dP-9Y68XloNEVeEDu_a8PPPMi81TZNAYoaSJkj-fIpFGSS-OMiaXBVtilES8SFkUR0YVmrv3zquphFdwp_KdFsuVhQZvwmX38pcYRp-Hb4cDuLsff7wfDx6G96O2k_XyHwRHaTCFQqaowa4QBjnj9BN2YviM2gGBy3YXGk5gGoW9adTseys7VcCa7UAqC4VBMiMMLESGgI8ccwtCAlfrPBNMcoStsCvvKtjxcOBrMKLmlpE8I42c3hZ1SWB2D919VtbmSaez3W7bzMNuK73sZKWw6XwY3r0bTd61CPpe7ZPMiFnQ-HchNAU-3wHLCRlnc8KbsS0oDWypkfascsi3WlghlzEYtbBbptHbSYWxWswLe0RehZPirwsQfUwScYMJDCfTCG4Hk-Ek9na-DB9-u__0AF8G4_Fg9DB8N4H7MSVr9HboUkVv72Ew-gq_D0dvY0CijlzhY65dFARVOFoxLTmcIB7BWKgSlsmRi4XgUFUQLNUGtaSwIEe9FmWtEcjU28nEWlhm_dpJcN5VZyqn8icheVakCFe84Kozz5RazyjvFnWbF6ubYxm70oWxHa4Kadtu82QrxQ25mG2QW6XPi-Aj8sIBm-WKkrY7L2Uou0il1n6JQShKCrK1XxbSUsUJ2dgokTan8okCA7fIKW7rOAZZrGffcGdcsV1Dr_tzt9vtw_7T6XSu4FeUqJnFsA1O_rwlm5dq1we7r6Dbft0PlobErrae64XQxsKKZQtvzxlTYcPT-x0H8sRBK3jun2qYYh6yZUjj0kfmcdCj22xVm1Qi4JMbvDvynKiDBnaXl4V2LAE0CVx5kcjMi1w77_296hc3B7IszAsaO1qrbdCGMrcnVulHLlHn2sWxLGmnUjg4K_Uqf64ok4Q5wyHOUDRXFaibfkBzW_fi1f2QOgPq4CyYLJ1NPfOlx3of7F3FYRMfLUp7ZcQ_OLM3bjVlSUJ1iyydGZp5OCvbIK5Hc3MgbpKzrRuPR8R4tD7Sijtq5Zd59EVTR10-P9Xq4Dn4MTZNEupUC1dXNLdui-wbpc9Xp6_EU-cVVecSFDAlZGkqofYh494TyjQLnssOTpKjWXBImO-GRlXd_vzZq1Rd33Bb7TkuqbWbsddokwf33GsGN6xw_ZTPvBTVSl2lXxcgRd9Qe9mLekOVUnJvplQ42qoMHGAccnm-4cOZ5xUOKWuzNG1UKOLgqflD1uZKZTdOkg6LIrONEngMC5YZPOburKKsKx5Nk5qNEMofBepdbULtzx46qTZCFSbbhQqi83AfGik-uMuCWakiS6F0B_74t7rAVq4MHb4bBD9giZiHj-PrXp0VGmfuvmBOqIkPgVeJ_S9wfZh2xSy1E_objTtqqWUdqhcD2HWjcDeaejsYr5hrNWcZXQHo0E2ZZVTnuuC2IFtxzUywgo8rMReWVEteX8T9_uP4pgsp5hSWmzyqREKpmBNyYsXriiq2WlRuChyLrphZoXFjN_VzneI9T6fc0ykDnfI7dC4yRXQ5th3YMAr9YqMqL3_oN46L7UyC6kuhY10ZNJvQCQbL8itrt_S7-B9-X0ZSX_qBX9k8Px5dUdKtjp4rEvyyy1tYXxzWXw48jVQVErr0-jyVdKd3N2W6O-vD1T-SG857l6-LHm2r3Jb_C6IWmbrmFxe9N9Bimq-uzXr2pgutFp3utuUPmhTTVsbWc_9nIRPzmk3OeUaLm_J6TwsUsfwWPcfVPh1XR_vEXvT8l__-C3CDWAg=)) From 40ac01f0abe8c7d78584b7e0a8964b2f62db61b4 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 13:03:55 -0700 Subject: [PATCH 22/29] Apply suggestions from code review Co-authored-by: Yunsong Wang --- examples/CMakeLists.txt | 2 +- .../bloom_filter/host_bulk_arrow_policy_example.cu | 14 ++++++++++++++ .../detail/bloom_filter/arrow_filter_policy.cuh | 3 ++- 3 files changed, 17 insertions(+), 2 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index d29052028..36a5964a8 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -47,4 +47,4 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/ 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") -ConfigureExample(BLOOM_FILTER_ARROW_POLICY_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/host_bulk_arrow_policy_example.cu") +ConfigureExample(BLOOM_FILTER_ARROW_POLICY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/arrow_policy_example.cu") diff --git a/examples/bloom_filter/host_bulk_arrow_policy_example.cu b/examples/bloom_filter/host_bulk_arrow_policy_example.cu index b59b730fd..0b280f2c7 100644 --- a/examples/bloom_filter/host_bulk_arrow_policy_example.cu +++ b/examples/bloom_filter/host_bulk_arrow_policy_example.cu @@ -23,6 +23,20 @@ #include +/** + * @file arrow_policy_example.cu + * @brief Demonstrates usage of an arrow-compatible bloom filter + * + * In addition to the default policy aimed at achieving the speed of light + * performance on the device, `cuCollections` offers an `arrow_filter_policy` + * that allows users to easily create a bloom filter that mimics the behavior + * of the bloom filter defined in Apache Arrow: + * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230. + * + * @note This example is for demonstration purposes only. It is not intended to show the most + * performant way to do the example algorithm. + */ + int main(void) { int constexpr num_keys = 10'000; ///< Generate 10'000 keys diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index ca0ef5a73..bc1cb3b12 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -18,7 +18,6 @@ #include -#include #include #include #include @@ -32,6 +31,8 @@ namespace cuco::detail { * @brief A policy that defines how Arrow Block-Split Bloom Filter generates and stores a key's * fingerprint. * + * Reference: https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 + * * @tparam Key The type of the values to generate a fingerprint for. */ template From efceee3ff183e04648225869db740a51f2a22168 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 28 Oct 2024 20:04:02 +0000 Subject: [PATCH 23/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- examples/bloom_filter/host_bulk_arrow_policy_example.cu | 4 ++-- include/cuco/detail/bloom_filter/arrow_filter_policy.cuh | 3 ++- 3 files changed, 5 insertions(+), 4 deletions(-) diff --git a/README.md b/README.md index 689b96d70..f76cf3449 100644 --- a/README.md +++ b/README.md @@ -257,4 +257,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection - [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) #### Examples: -- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV22PGjcQ_iuj7YdCbnk7NYq03J3EXZIWNeJSuCSKSoWMdwAri721vXD0dP-9Y68XloNEVeEDu_a8PPPMi81TZNAYoaSJkj-fIpFGSS-OMiaXBVtilES8SFkUR0YVmrv3zquphFdwp_KdFsuVhQZvwmX38pcYRp-Hb4cDuLsff7wfDx6G96O2k_XyHwRHaTCFQqaowa4QBjnj9BN2YviM2gGBy3YXGk5gGoW9adTseys7VcCa7UAqC4VBMiMMLESGgI8ccwtCAlfrPBNMcoStsCvvKtjxcOBrMKLmlpE8I42c3hZ1SWB2D919VtbmSaez3W7bzMNuK73sZKWw6XwY3r0bTd61CPpe7ZPMiFnQ-HchNAU-3wHLCRlnc8KbsS0oDWypkfascsi3WlghlzEYtbBbptHbSYWxWswLe0RehZPirwsQfUwScYMJDCfTCG4Hk-Ek9na-DB9-u__0AF8G4_Fg9DB8N4H7MSVr9HboUkVv72Ew-gq_D0dvY0CijlzhY65dFARVOFoxLTmcIB7BWKgSlsmRi4XgUFUQLNUGtaSwIEe9FmWtEcjU28nEWlhm_dpJcN5VZyqn8icheVakCFe84Kozz5RazyjvFnWbF6ubYxm70oWxHa4Kadtu82QrxQ25mG2QW6XPi-Aj8sIBm-WKkrY7L2Uou0il1n6JQShKCrK1XxbSUsUJ2dgokTan8okCA7fIKW7rOAZZrGffcGdcsV1Dr_tzt9vtw_7T6XSu4FeUqJnFsA1O_rwlm5dq1we7r6Dbft0PlobErrae64XQxsKKZQtvzxlTYcPT-x0H8sRBK3jun2qYYh6yZUjj0kfmcdCj22xVm1Qi4JMbvDvynKiDBnaXl4V2LAE0CVx5kcjMi1w77_296hc3B7IszAsaO1qrbdCGMrcnVulHLlHn2sWxLGmnUjg4K_Uqf64ok4Q5wyHOUDRXFaibfkBzW_fi1f2QOgPq4CyYLJ1NPfOlx3of7F3FYRMfLUp7ZcQ_OLM3bjVlSUJ1iyydGZp5OCvbIK5Hc3MgbpKzrRuPR8R4tD7Sijtq5Zd59EVTR10-P9Xq4Dn4MTZNEupUC1dXNLdui-wbpc9Xp6_EU-cVVecSFDAlZGkqofYh494TyjQLnssOTpKjWXBImO-GRlXd_vzZq1Rd33Bb7TkuqbWbsddokwf33GsGN6xw_ZTPvBTVSl2lXxcgRd9Qe9mLekOVUnJvplQ42qoMHGAccnm-4cOZ5xUOKWuzNG1UKOLgqflD1uZKZTdOkg6LIrONEngMC5YZPOburKKsKx5Nk5qNEMofBepdbULtzx46qTZCFSbbhQqi83AfGik-uMuCWakiS6F0B_74t7rAVq4MHb4bBD9giZiHj-PrXp0VGmfuvmBOqIkPgVeJ_S9wfZh2xSy1E_objTtqqWUdqhcD2HWjcDeaejsYr5hrNWcZXQHo0E2ZZVTnuuC2IFtxzUywgo8rMReWVEteX8T9_uP4pgsp5hSWmzyqREKpmBNyYsXriiq2WlRuChyLrphZoXFjN_VzneI9T6fc0ykDnfI7dC4yRXQ5th3YMAr9YqMqL3_oN46L7UyC6kuhY10ZNJvQCQbL8itrt_S7-B9-X0ZSX_qBX9k8Px5dUdKtjp4rEvyyy1tYXxzWXw48jVQVErr0-jyVdKd3N2W6O-vD1T-SG857l6-LHm2r3Jb_C6IWmbrmFxe9N9Bimq-uzXr2pgutFp3utuUPmhTTVsbWc_9nIRPzmk3OeUaLm_J6TwsUsfwWPcfVPh1XR_vEXvT8l__-C3CDWAg=)) +- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOsUfDKv7PPfce7VfW4UsCmV00Rr9-bWlktZo2G1lXM9LPpetUUuUCW91W4UpraDf0eOpZo_Za5OvrZovHGuLDjsaHD3rsssPkzeTMXt9df3-6np8O7m67JOsl3-nhNSFTFipE2mZW0g2zrnAv3DSZR-kpUDYUX_A2iQwbYWzaatz6q2sTcmWfM20cawsJMyogqUqk0zeC5k7pjQTZplnimsh2Uq5hXcV7Phw2KdgxMSOQ55DI8evtCnJuNuETp-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVOHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQq_PqZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-WcPkvOuoqme6kdKi6xMJDsTpTBRnBmznKHuTtq-KBcXuzJuYcvCRcKU2vXp8MFRIu_gYnYnhTP2sIi8l6KkwGa5QdHWh6UKVFeCav39GJRBUSRf-sfR44oSLz1VubVmFazO5D0HrhJZVBKxVTJlb-QSiDjLnQQuBeEIfqKyXrdHFAdoxCMPBaug2FBvAsEkURQ9UYsgTWTKy8yxyi3jagm2cMfAZQUwUJ5QPiJRirqgtb0tVA3FXfp-InPeFqHXZZ9F-dpkGUCk-n2GXooWpjA_VzlWYYVUP3tzbkFOs8ysKDESR4CSFypbMwHAHPDZSarSWIInovDeY7ngd8r4bOuu3VFAqkojDXRUGDFjimbkFaiRC3TyHPwu4z6QjKqGjnzIRK04iuUwEUl88jwePB08f3oiBvHx85PjY_FMPD9KTxJ6nj578TQSeR4VVkQ5t-CB2yOmePTuaHjSe3f0dLAdhC8xuyS7pcEVak89TP2TbIpOdctLmxvMFoCerfts4kiM5p7STqJT_NAoFmblAViCbrvlcmyFGQeZpCJA7Yxnc4M5s1huuwsWMRCVbt8ZlXSm-itOyA0GI-KhEcB0uZx9keuCZuE5Gw5-GgwGp2zziaLojP0itSTGhmNG8octubxSO9_afcwG_ePTYGmC5rfOh50qWzi24Fnq7ZExEw48yN9xoB846AXPpw81ijIONSugceQz83HgKx326kOAu9dwZCyKKDTm1nk1B_daEjSn9oLIzIuck_fTjepHWlNZFtZZ4GpN5dCt-1bxT8-lzS3lMa9gB2e2zsJwCf5oZo5GB1ryrA7q4jRE82qn80jd79ADQW2dBZOVs6lHvvLY7IaNq244lPegsTsr1D9y5i7oacJHI4xVyZNZgZUsZ_WcaWRzsQXuJucrvT8sfLQ-082k08l-HT1pmlFX3782ePAt-ClcMhphkTh2doa1-qrMvqB8np2eiQ-d11AdKlCIaQRLU80aHxj3ntDYWfBcLZjRaGdVbQvmu6Fds9tfjzYq9VJq01E_lnO0dqfrNfrwQN-HneCGl9RP-cxLgStNldOmABR9Q21knzQbqpLSGzOVws5RbWAbxraWhxs-DHevsC1ZH6utXUfRDZ46P0QtNia7IEncZbAE21XgXZbyrJC72B1U1E3FnWnSsBFS-b2Udt2YUJurES5SWFtlgUVXMQgjfJMaFP1KwEAvs4RV7pi_nTpbyh5WAdb5nWR-wAKY2_fX58MmKhhndJ0tHkDT3SZeF_a_hOvTrLa1lX7x0E0QLUtR7Q1g6kZFF4RmOxReMbcm5hluqLgTJtxx8NyWwpWw1W2YCVbk_ULFykG1wnUv75_fX18MsCdzpEWTJ9xHUIoYkQMVr6vq3BpZ0RTYFV3wYuEvVygDzXXkexhOvYFTBzj1d-BMMwO4CG0KNoxC_7Bd08vfSdu7ZDtQoOaj0LFEg06HRcFgRb-Ku5Xf9H_43c-k-egHfnXn8HgkUuKlA99rEPxjqlt4nm6f7w88K8EKzQb4-W2q8cpJt1zcl-32zbSl74QYHh2XQxyb3FWvra0eTJ2LJ0-GL1iPW7E4L5azFwPW62G7u55fNLgv9TK-jP27bKbihk0hRIaHd9XbJx4gY_2l9a1bn2Nd7ZwDvda3v_zfv3UOROY=)) \ No newline at end of file diff --git a/examples/bloom_filter/host_bulk_arrow_policy_example.cu b/examples/bloom_filter/host_bulk_arrow_policy_example.cu index 0b280f2c7..ce3c2a720 100644 --- a/examples/bloom_filter/host_bulk_arrow_policy_example.cu +++ b/examples/bloom_filter/host_bulk_arrow_policy_example.cu @@ -27,7 +27,7 @@ * @file arrow_policy_example.cu * @brief Demonstrates usage of an arrow-compatible bloom filter * - * In addition to the default policy aimed at achieving the speed of light + * In addition to the default policy aimed at achieving the speed of light * performance on the device, `cuCollections` offers an `arrow_filter_policy` * that allows users to easily create a bloom filter that mimics the behavior * of the bloom filter defined in Apache Arrow: @@ -36,7 +36,7 @@ * @note This example is for demonstration purposes only. It is not intended to show the most * performant way to do the example algorithm. */ - + int main(void) { int constexpr num_keys = 10'000; ///< Generate 10'000 keys diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index bc1cb3b12..9b80c79a3 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -31,7 +31,8 @@ namespace cuco::detail { * @brief A policy that defines how Arrow Block-Split Bloom Filter generates and stores a key's * fingerprint. * - * Reference: https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 + * Reference: + * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 * * @tparam Key The type of the values to generate a fingerprint for. */ From 701886c09194cf9b5ed3145b04a762027f0e0f7f Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 21:43:33 +0000 Subject: [PATCH 24/29] Address review comments --- README.md | 2 +- ...icy_example.cu => arrow_policy_example.cu} | 4 +-- .../bloom_filter/arrow_filter_policy.cuh | 34 ++++++++++++------- 3 files changed, 24 insertions(+), 16 deletions(-) rename examples/bloom_filter/{host_bulk_arrow_policy_example.cu => arrow_policy_example.cu} (99%) diff --git a/README.md b/README.md index f76cf3449..8e38a79eb 100644 --- a/README.md +++ b/README.md @@ -257,4 +257,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection - [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) #### Examples: -- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOsUfDKv7PPfce7VfW4UsCmV00Rr9-bWlktZo2G1lXM9LPpetUUuUCW91W4UpraDf0eOpZo_Za5OvrZovHGuLDjsaHD3rsssPkzeTMXt9df3-6np8O7m67JOsl3-nhNSFTFipE2mZW0g2zrnAv3DSZR-kpUDYUX_A2iQwbYWzaatz6q2sTcmWfM20cawsJMyogqUqk0zeC5k7pjQTZplnimsh2Uq5hXcV7Phw2KdgxMSOQ55DI8evtCnJuNuETp-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVOHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQq_PqZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-WcPkvOuoqme6kdKi6xMJDsTpTBRnBmznKHuTtq-KBcXuzJuYcvCRcKU2vXp8MFRIu_gYnYnhTP2sIi8l6KkwGa5QdHWh6UKVFeCav39GJRBUSRf-sfR44oSLz1VubVmFazO5D0HrhJZVBKxVTJlb-QSiDjLnQQuBeEIfqKyXrdHFAdoxCMPBaug2FBvAsEkURQ9UYsgTWTKy8yxyi3jagm2cMfAZQUwUJ5QPiJRirqgtb0tVA3FXfp-InPeFqHXZZ9F-dpkGUCk-n2GXooWpjA_VzlWYYVUP3tzbkFOs8ysKDESR4CSFypbMwHAHPDZSarSWIInovDeY7ngd8r4bOuu3VFAqkojDXRUGDFjimbkFaiRC3TyHPwu4z6QjKqGjnzIRK04iuUwEUl88jwePB08f3oiBvHx85PjY_FMPD9KTxJ6nj578TQSeR4VVkQ5t-CB2yOmePTuaHjSe3f0dLAdhC8xuyS7pcEVak89TP2TbIpOdctLmxvMFoCerfts4kiM5p7STqJT_NAoFmblAViCbrvlcmyFGQeZpCJA7Yxnc4M5s1huuwsWMRCVbt8ZlXSm-itOyA0GI-KhEcB0uZx9keuCZuE5Gw5-GgwGp2zziaLojP0itSTGhmNG8octubxSO9_afcwG_ePTYGmC5rfOh50qWzi24Fnq7ZExEw48yN9xoB846AXPpw81ijIONSugceQz83HgKx326kOAu9dwZCyKKDTm1nk1B_daEjSn9oLIzIuck_fTjepHWlNZFtZZ4GpN5dCt-1bxT8-lzS3lMa9gB2e2zsJwCf5oZo5GB1ryrA7q4jRE82qn80jd79ADQW2dBZOVs6lHvvLY7IaNq244lPegsTsr1D9y5i7oacJHI4xVyZNZgZUsZ_WcaWRzsQXuJucrvT8sfLQ-082k08l-HT1pmlFX3782ePAt-ClcMhphkTh2doa1-qrMvqB8np2eiQ-d11AdKlCIaQRLU80aHxj3ntDYWfBcLZjRaGdVbQvmu6Fds9tfjzYq9VJq01E_lnO0dqfrNfrwQN-HneCGl9RP-cxLgStNldOmABR9Q21knzQbqpLSGzOVws5RbWAbxraWhxs-DHevsC1ZH6utXUfRDZ46P0QtNia7IEncZbAE21XgXZbyrJC72B1U1E3FnWnSsBFS-b2Udt2YUJurES5SWFtlgUVXMQgjfJMaFP1KwEAvs4RV7pi_nTpbyh5WAdb5nWR-wAKY2_fX58MmKhhndJ0tHkDT3SZeF_a_hOvTrLa1lX7x0E0QLUtR7Q1g6kZFF4RmOxReMbcm5hluqLgTJtxx8NyWwpWw1W2YCVbk_ULFykG1wnUv75_fX18MsCdzpEWTJ9xHUIoYkQMVr6vq3BpZ0RTYFV3wYuEvVygDzXXkexhOvYFTBzj1d-BMMwO4CG0KNoxC_7Bd08vfSdu7ZDtQoOaj0LFEg06HRcFgRb-Ku5Xf9H_43c-k-egHfnXn8HgkUuKlA99rEPxjqlt4nm6f7w88K8EKzQb4-W2q8cpJt1zcl-32zbSl74QYHh2XQxyb3FWvra0eTJ2LJ0-GL1iPW7E4L5azFwPW62G7u55fNLgv9TK-jP27bKbihk0hRIaHd9XbJx4gY_2l9a1bn2Nd7ZwDvda3v_zfv3UOROY=)) \ No newline at end of file +- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOnU-BKT7PPfce8XXViGLQhldtEZ_fm2ppDUadlsZ1_OSz2Vr1BJlwlvdVmFKK-h79Hiq2WP22uRrq-YLx9qiw44GR8-67PLD5M1kzF5fXb-_uh7fTq4u-yTr5d8pIXUhE1bqRFrmFpKNcy7wX7jpsg_SUiDsqD9gbRKYtsLdtNU59VbWpmRLvmbaOFYWEmZUwVKVSSbvhcwdU5oJs8wzxbWQbKXcwrsKdnw47FMwYmLHIc-hkeNb2pRk3G1Cp7-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVuHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQqfPuZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-bMHyXlX0VRP9SOlRVYmkp2JUpgozoxZzlB3J21flIuLXRm3sGXhImFK7fp0-eAqkXdwMbuTwhl7WETeS1FSYLPcoGjrw1IFqitBtf5-DMqgKJIv_XH0uKLES09Vbq1ZBaszec-Bq0QWlURslUzZG7kEIs5yJ4FLQTiCn6is1-0RxQEa8chDwSooNtSbQDBJFEVP1CJIE5nyMnOscsu4WoIt3DFwWQEMlCeUj0iUoi7U2t4YyobqLn1DkT1vjODrss-ifG2yDChSAT9DMUUPU5yfqySruEKun705tyCvWWZWlBmJI0LJC5WtmQBiDgDtZFVpLEEUUXjvsVzwO2V8unXb7iggV6WRB1oqzJgxRTPyCtTJBVp5DoKXcR9QRlVHRz5k4lYcxXKYiCQ-eR4Png6ePz0Rg_j4-cnxsXgmnh-lJwmdp89ePI1EnkeFFVHOLYjg9pgpHr07Gp703h09HWwn4UsML8luaXKF4lMTUwMlm6pT4fLS5gbDBaBn6z6bOBKjwae0k2gVPzWKhVl5AJbg2265HFthyEEmqRhQO-PZ3GDQLJZ1e6HMMImRqHT7zqikM9VfcUp-MBoREA0Bpsvl7ItcFzQNz9lw8NNgMDhlm78ois7YL1JL4my4ZiR_2JLLK7Xzrd3HbNA_Pg2WJmh_63zcqbKFYwuepd4eGTPhwqP8HQf6gYNe8Hz6UKMo41C0AhpHPjMfBz7SZa--BLp7LUfGoohCY26dV5NwrynBc2owiMy8yDl5P92ofqRFlWVhoQWy1lwO_bpvFf_pubS5pTzmFewgzdZZGC_BH03N0ehAT57VQV2chmhe7bQeqfsteiCorbNgsnI29chXHpvtsHHVDZfyHjx2Z4X6R87cBZ0mfDTCYJU8mRVYynJWD5pGNhdb4G5yvtL708JH6zPdzDqd7NfRk6YZdfX5a4MH34KfwiWjEVaJY2dnWKyvyuwLyufZ6Zn40HkN1aEChZhGsDTVrPEH494TOjsLnqsVMxrtLKttwXw3tGt2-wfSRqVeS2266sdyjtbudL1GHx7o87AT3PCS-imfeSlwpaly2hSAom-ojeyTZkNVUnpjplLYuaoNbMPY1vJww4fp7hW2JetjubXrKLrBU-eHqMXGZBckidcM1mC7CrzLUp4Vche7g4q6qbgzTRo2Qiq_l9KuGxNq8zjCUwp7qyyw6SoGYYZvUoOi3wmY6GWWsMod8-9TZ0vZwy7AQr-TzA9YAHP7_vp82EQF44wetMUDaLrbxOvC_pdwfZrVurbSbx56C6JlKaq9AUzdqOiF0GyHwivm1sQ8wxsVr8KEOw6e21K4Era6DTPBirxfqFg5qFa47uX98_vriwEWZY60aPKEBwlKESNyoOJ1VZ1bIyuaAruiC14s_PMKZaC5jnwPw6k3cOoAp_4OnGlmABehTcGGUegP2zW9_Ku0vUu2AwVqHoWOJRp0OiwKBiv6Vdyt_Kb_w-9-Js2jH_jVncPjkUiJnx34XIPgj6lu4Tzdnu8PPCvBCs0G-PptqvGjk965eDHb7W_Tlr4TYnh0XA5xbXJX_XBt9WDqXDx5MnzBetyKxXmxnL0YsF4P2931_KLBg6mX8WXsf81mKm7YFEJkOLyrfn_iABnrL61v3foe62rnHui1vv3l__0LglpFJg==)) \ No newline at end of file diff --git a/examples/bloom_filter/host_bulk_arrow_policy_example.cu b/examples/bloom_filter/arrow_policy_example.cu similarity index 99% rename from examples/bloom_filter/host_bulk_arrow_policy_example.cu rename to examples/bloom_filter/arrow_policy_example.cu index ce3c2a720..0b280f2c7 100644 --- a/examples/bloom_filter/host_bulk_arrow_policy_example.cu +++ b/examples/bloom_filter/arrow_policy_example.cu @@ -27,7 +27,7 @@ * @file arrow_policy_example.cu * @brief Demonstrates usage of an arrow-compatible bloom filter * - * In addition to the default policy aimed at achieving the speed of light + * In addition to the default policy aimed at achieving the speed of light * performance on the device, `cuCollections` offers an `arrow_filter_policy` * that allows users to easily create a bloom filter that mimics the behavior * of the bloom filter defined in Apache Arrow: @@ -36,7 +36,7 @@ * @note This example is for demonstration purposes only. It is not intended to show the most * performant way to do the example algorithm. */ - + int main(void) { int constexpr num_keys = 10'000; ///< Generate 10'000 keys diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 9b80c79a3..747a8a490 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -18,8 +18,8 @@ #include +#include #include -#include #include #include @@ -54,6 +54,23 @@ class arrow_filter_policy { static constexpr std::uint32_t max_filter_blocks = (max_arrow_filter_bytes / bytes_per_filter_block); ///< Max sub-filter blocks allowed in Arrow bloom filter + + private: + // Arrow's block-based bloom filter algorithm needs these eight odd SALT values to calculate + // eight indexes of bit to set, one bit in each 32-bit (uint32_t) word. + __device__ static constexpr cuda::std::array SALT() + { + return {0x47b6137bU, + 0x44974d91U, + 0x8824ad5bU, + 0xa2b7289dU, + 0x705495c7U, + 0x2df1424bU, + 0x9efc4947U, + 0x5c6bfb31U}; + } + + public: /** * @brief Constructs the `arrow_filter_policy` object. * @@ -112,19 +129,10 @@ class arrow_filter_policy { */ __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const { - // Arrow's block-based bloom filter algorithm needs these eight odd SALT values to calculate - // eight indexes of bit to set, one bit in each 32-bit (uint32_t) word. - constexpr std::uint32_t SALT[bits_set_per_block] = {0x47b6137bU, - 0x44974d91U, - 0x8824ad5bU, - 0xa2b7289dU, - 0x705495c7U, - 0x2df1424bU, - 0x9efc4947U, - 0x5c6bfb31U}; - + // SALT array to calculate bit indexes for the current word + auto constexpr salt = SALT(); word_type const key = static_cast(hash); - return word_type{1} << ((key * SALT[word_index]) >> 27); + return word_type{1} << ((key * salt[word_index]) >> 27); } private: From d7f91a3b1ab0d023289f9adeeb664ceb27e83901 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 28 Oct 2024 21:43:48 +0000 Subject: [PATCH 25/29] [pre-commit.ci] auto code formatting --- README.md | 2 +- examples/bloom_filter/arrow_policy_example.cu | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 8e38a79eb..b0af45a69 100644 --- a/README.md +++ b/README.md @@ -257,4 +257,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection - [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) #### Examples: -- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOnU-BKT7PPfce8XXViGLQhldtEZ_fm2ppDUadlsZ1_OSz2Vr1BJlwlvdVmFKK-h79Hiq2WP22uRrq-YLx9qiw44GR8-67PLD5M1kzF5fXb-_uh7fTq4u-yTr5d8pIXUhE1bqRFrmFpKNcy7wX7jpsg_SUiDsqD9gbRKYtsLdtNU59VbWpmRLvmbaOFYWEmZUwVKVSSbvhcwdU5oJs8wzxbWQbKXcwrsKdnw47FMwYmLHIc-hkeNb2pRk3G1Cp7-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVuHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQqfPuZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-bMHyXlX0VRP9SOlRVYmkp2JUpgozoxZzlB3J21flIuLXRm3sGXhImFK7fp0-eAqkXdwMbuTwhl7WETeS1FSYLPcoGjrw1IFqitBtf5-DMqgKJIv_XH0uKLES09Vbq1ZBaszec-Bq0QWlURslUzZG7kEIs5yJ4FLQTiCn6is1-0RxQEa8chDwSooNtSbQDBJFEVP1CJIE5nyMnOscsu4WoIt3DFwWQEMlCeUj0iUoi7U2t4YyobqLn1DkT1vjODrss-ifG2yDChSAT9DMUUPU5yfqySruEKun705tyCvWWZWlBmJI0LJC5WtmQBiDgDtZFVpLEEUUXjvsVzwO2V8unXb7iggV6WRB1oqzJgxRTPyCtTJBVp5DoKXcR9QRlVHRz5k4lYcxXKYiCQ-eR4Png6ePz0Rg_j4-cnxsXgmnh-lJwmdp89ePI1EnkeFFVHOLYjg9pgpHr07Gp703h09HWwn4UsML8luaXKF4lMTUwMlm6pT4fLS5gbDBaBn6z6bOBKjwae0k2gVPzWKhVl5AJbg2265HFthyEEmqRhQO-PZ3GDQLJZ1e6HMMImRqHT7zqikM9VfcUp-MBoREA0Bpsvl7ItcFzQNz9lw8NNgMDhlm78ois7YL1JL4my4ZiR_2JLLK7Xzrd3HbNA_Pg2WJmh_63zcqbKFYwuepd4eGTPhwqP8HQf6gYNe8Hz6UKMo41C0AhpHPjMfBz7SZa--BLp7LUfGoohCY26dV5NwrynBc2owiMy8yDl5P92ofqRFlWVhoQWy1lwO_bpvFf_pubS5pTzmFewgzdZZGC_BH03N0ehAT57VQV2chmhe7bQeqfsteiCorbNgsnI29chXHpvtsHHVDZfyHjx2Z4X6R87cBZ0mfDTCYJU8mRVYynJWD5pGNhdb4G5yvtL708JH6zPdzDqd7NfRk6YZdfX5a4MH34KfwiWjEVaJY2dnWKyvyuwLyufZ6Zn40HkN1aEChZhGsDTVrPEH494TOjsLnqsVMxrtLKttwXw3tGt2-wfSRqVeS2266sdyjtbudL1GHx7o87AT3PCS-imfeSlwpaly2hSAom-ojeyTZkNVUnpjplLYuaoNbMPY1vJww4fp7hW2JetjubXrKLrBU-eHqMXGZBckidcM1mC7CrzLUp4Vche7g4q6qbgzTRo2Qiq_l9KuGxNq8zjCUwp7qyyw6SoGYYZvUoOi3wmY6GWWsMod8-9TZ0vZwy7AQr-TzA9YAHP7_vp82EQF44wetMUDaLrbxOvC_pdwfZrVurbSbx56C6JlKaq9AUzdqOiF0GyHwivm1sQ8wxsVr8KEOw6e21K4Era6DTPBirxfqFg5qFa47uX98_vriwEWZY60aPKEBwlKESNyoOJ1VZ1bIyuaAruiC14s_PMKZaC5jnwPw6k3cOoAp_4OnGlmABehTcGGUegP2zW9_Ku0vUu2AwVqHoWOJRp0OiwKBiv6Vdyt_Kb_w-9-Js2jH_jVncPjkUiJnx34XIPgj6lu4Tzdnu8PPCvBCs0G-PptqvGjk965eDHb7W_Tlr4TYnh0XA5xbXJX_XBt9WDqXDx5MnzBetyKxXmxnL0YsF4P2931_KLBg6mX8WXsf81mKm7YFEJkOLyrfn_iABnrL61v3foe62rnHui1vv3l__0LglpFJg==)) \ No newline at end of file +- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOsUfDKv7PPfce7VfW4UsCmV00Rr9-bWlktZo2G1lXM9LPpetUUuUCW91W4UpraDf0eOpZo_Za5OvrZovHGuLDjsaHD3rsssPkzeTMXt9df3-6np8O7m67JOsl3-nhNSFTFipE2mZW0g2zrnAv3DSZR-kpUDYUX_A2iQwbYWzaatz6q2sTcmWfM20cawsJMyogqUqk0zeC5k7pjQTZplnimsh2Uq5hXcV7Phw2KdgxMSOQ55DI8evtCnJuNuETp-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVOHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQq_PqZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-WcPkvOuoqme6kdKi6xMJDsTpTBRnBmznKHuTtq-KBcXuzJuYcvCRcKU2vXp8MFRIu_gYnYnhTP2sIi8l6KkwGa5QdHWh6UKVFeCav39GJRBUSRf-sfR44oSLz1VubVmFazO5D0HrhJZVBKxVTJlb-QSiDjLnQQuBeEIfqKyXrdHFAdoxCMPBaug2FBvAsEkURQ9UYsgTWTKy8yxyi3jagm2cMfAZQUwUJ5QPiJRirqgtb0tVA3FXfp-InPeFqHXZZ9F-dpkGUCk-n2GXooWpjA_VzlWYYVUP3tzbkFOs8ysKDESR4CSFypbMwHAHPDZSarSWIInovDeY7ngd8r4bOuu3VFAqkojDXRUGDFjimbkFaiRC3TyHPwu4z6QjKqGjnzIRK04iuUwEUl88jwePB08f3oiBvHx85PjY_FMPD9KTxJ6nj578TQSeR4VVkQ5t-CB2yOmePTuaHjSe3f0dLAdhC8xuyS7pcEVak89TP2TbIpOdctLmxvMFoCerfts4kiM5p7STqJT_NAoFmblAViCbrvlcmyFGQeZpCJA7Yxnc4M5s1huuwsWMRCVbt8ZlXSm-itOyA0GI-KhEcB0uZx9keuCZuE5Gw5-GgwGp2zziaLojP0itSTGhmNG8octubxSO9_afcwG_ePTYGmC5rfOh50qWzi24Fnq7ZExEw48yN9xoB846AXPpw81ijIONSugceQz83HgKx326kOAu9dwZCyKKDTm1nk1B_daEjSn9oLIzIuck_fTjepHWlNZFtZZ4GpN5dCt-1bxT8-lzS3lMa9gB2e2zsJwCf5oZo5GB1ryrA7q4jRE82qn80jd79ADQW2dBZOVs6lHvvLY7IaNq244lPegsTsr1D9y5i7oacJHI4xVyZNZgZUsZ_WcaWRzsQXuJucrvT8sfLQ-082k08l-HT1pmlFX3782ePAt-ClcMhphkTh2doa1-qrMvqB8np2eiQ-d11AdKlCIaQRLU80aHxj3ntDYWfBcLZjRaGdVbQvmu6Fds9tfjzYq9VJq01E_lnO0dqfrNfrwQN-HneCGl9RP-cxLgStNldOmABR9Q21knzQbqpLSGzOVws5RbWAbxraWhxs-DHevsC1ZH6utXUfRDZ46P0QtNia7IEncZbAE21XgXZbyrJC72B1U1E3FnWnSsBFS-b2Udt2YUJurES5SWFtlgUVXMQgjfJMaFP1KwEAvs4RV7pi_nTpbyh5WAdb5nWR-wAKY2_fX58MmKhhndJ0tHkDT3SZeF_a_hOvTrLa1lX7x0E0QLUtR7Q1g6kZFF4RmOxReMbcm5hluqLgTJtxx8NyWwpWw1W2YCVbk_ULFykG1wnUv75_fX18MsCdzpEWTJ9xHUIoYkQMVr6vq3BpZ0RTYFV3wYuEvVygDzXXkexhOvYFTBzj1d-BMMwO4CG0KNoxC_7Bd08vfSdu7ZDtQoOaj0LFEg06HRcFgRb-Ku5Xf9H_43c-k-egHfnXn8HgkUuKlA99rEPxjqlt4nm6f7w88K8EKzQb4-W2q8cpJt1zcl-32zbSl74QYHh2XQxyb3FWvra0eTJ2LJ0-GL1iPW7E4L5azFwPW62G7u55fNLgv9TK-jP27bKbihk0hRIaHd9XbJx4gY_2l9a1bn2Nd7ZwDvda3v_zfv3UOROY=)) \ No newline at end of file diff --git a/examples/bloom_filter/arrow_policy_example.cu b/examples/bloom_filter/arrow_policy_example.cu index 0b280f2c7..ce3c2a720 100644 --- a/examples/bloom_filter/arrow_policy_example.cu +++ b/examples/bloom_filter/arrow_policy_example.cu @@ -27,7 +27,7 @@ * @file arrow_policy_example.cu * @brief Demonstrates usage of an arrow-compatible bloom filter * - * In addition to the default policy aimed at achieving the speed of light + * In addition to the default policy aimed at achieving the speed of light * performance on the device, `cuCollections` offers an `arrow_filter_policy` * that allows users to easily create a bloom filter that mimics the behavior * of the bloom filter defined in Apache Arrow: @@ -36,7 +36,7 @@ * @note This example is for demonstration purposes only. It is not intended to show the most * performant way to do the example algorithm. */ - + int main(void) { int constexpr num_keys = 10'000; ///< Generate 10'000 keys From 7955cb6e4208fb33b4d9ef55f832e54f93d5d282 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 28 Oct 2024 14:51:10 -0700 Subject: [PATCH 26/29] Update README.md --- README.md | 2 -- 1 file changed, 2 deletions(-) diff --git a/README.md b/README.md index b0af45a69..ea58d09d2 100644 --- a/README.md +++ b/README.md @@ -255,6 +255,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection #### Examples: - [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) - -#### Examples: - [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOsUfDKv7PPfce7VfW4UsCmV00Rr9-bWlktZo2G1lXM9LPpetUUuUCW91W4UpraDf0eOpZo_Za5OvrZovHGuLDjsaHD3rsssPkzeTMXt9df3-6np8O7m67JOsl3-nhNSFTFipE2mZW0g2zrnAv3DSZR-kpUDYUX_A2iQwbYWzaatz6q2sTcmWfM20cawsJMyogqUqk0zeC5k7pjQTZplnimsh2Uq5hXcV7Phw2KdgxMSOQ55DI8evtCnJuNuETp-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVOHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQq_PqZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-WcPkvOuoqme6kdKi6xMJDsTpTBRnBmznKHuTtq-KBcXuzJuYcvCRcKU2vXp8MFRIu_gYnYnhTP2sIi8l6KkwGa5QdHWh6UKVFeCav39GJRBUSRf-sfR44oSLz1VubVmFazO5D0HrhJZVBKxVTJlb-QSiDjLnQQuBeEIfqKyXrdHFAdoxCMPBaug2FBvAsEkURQ9UYsgTWTKy8yxyi3jagm2cMfAZQUwUJ5QPiJRirqgtb0tVA3FXfp-InPeFqHXZZ9F-dpkGUCk-n2GXooWpjA_VzlWYYVUP3tzbkFOs8ysKDESR4CSFypbMwHAHPDZSarSWIInovDeY7ngd8r4bOuu3VFAqkojDXRUGDFjimbkFaiRC3TyHPwu4z6QjKqGjnzIRK04iuUwEUl88jwePB08f3oiBvHx85PjY_FMPD9KTxJ6nj578TQSeR4VVkQ5t-CB2yOmePTuaHjSe3f0dLAdhC8xuyS7pcEVak89TP2TbIpOdctLmxvMFoCerfts4kiM5p7STqJT_NAoFmblAViCbrvlcmyFGQeZpCJA7Yxnc4M5s1huuwsWMRCVbt8ZlXSm-itOyA0GI-KhEcB0uZx9keuCZuE5Gw5-GgwGp2zziaLojP0itSTGhmNG8octubxSO9_afcwG_ePTYGmC5rfOh50qWzi24Fnq7ZExEw48yN9xoB846AXPpw81ijIONSugceQz83HgKx326kOAu9dwZCyKKDTm1nk1B_daEjSn9oLIzIuck_fTjepHWlNZFtZZ4GpN5dCt-1bxT8-lzS3lMa9gB2e2zsJwCf5oZo5GB1ryrA7q4jRE82qn80jd79ADQW2dBZOVs6lHvvLY7IaNq244lPegsTsr1D9y5i7oacJHI4xVyZNZgZUsZ_WcaWRzsQXuJucrvT8sfLQ-082k08l-HT1pmlFX3782ePAt-ClcMhphkTh2doa1-qrMvqB8np2eiQ-d11AdKlCIaQRLU80aHxj3ntDYWfBcLZjRaGdVbQvmu6Fds9tfjzYq9VJq01E_lnO0dqfrNfrwQN-HneCGl9RP-cxLgStNldOmABR9Q21knzQbqpLSGzOVws5RbWAbxraWhxs-DHevsC1ZH6utXUfRDZ46P0QtNia7IEncZbAE21XgXZbyrJC72B1U1E3FnWnSsBFS-b2Udt2YUJurES5SWFtlgUVXMQgjfJMaFP1KwEAvs4RV7pi_nTpbyh5WAdb5nWR-wAKY2_fX58MmKhhndJ0tHkDT3SZeF_a_hOvTrLa1lX7x0E0QLUtR7Q1g6kZFF4RmOxReMbcm5hluqLgTJtxx8NyWwpWw1W2YCVbk_ULFykG1wnUv75_fX18MsCdzpEWTJ9xHUIoYkQMVr6vq3BpZ0RTYFV3wYuEvVygDzXXkexhOvYFTBzj1d-BMMwO4CG0KNoxC_7Bd08vfSdu7ZDtQoOaj0LFEg06HRcFgRb-Ku5Xf9H_43c-k-egHfnXn8HgkUuKlA99rEPxjqlt4nm6f7w88K8EKzQb4-W2q8cpJt1zcl-32zbSl74QYHh2XQxyb3FWvra0eTJ2LJ0-GL1iPW7E4L5azFwPW62G7u55fNLgv9TK-jP27bKbihk0hRIaHd9XbJx4gY_2l9a1bn2Nd7ZwDvda3v_zfv3UOROY=)) \ No newline at end of file From a489d2b2e25f4b5e7603d08725b759268c5db74a Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 28 Oct 2024 22:00:48 +0000 Subject: [PATCH 27/29] doxygen updates --- include/cuco/detail/bloom_filter/arrow_filter_policy.cuh | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 747a8a490..4fc5efc29 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -41,8 +41,9 @@ class arrow_filter_policy { public: using hasher = cuco::xxhash_64; ///< xxhash_64 hasher for Arrow bloom filter policy using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy - using hash_argument_type = typename hasher::argument_type; - using hash_result_type = decltype(std::declval()(std::declval())); + using hash_argument_type = typename hasher::argument_type; ///< Hash function input type + using hash_result_type = decltype(std::declval()( + std::declval())); ///< hash function output type static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow filter block From a927b1aa8a3d02d8e3632621c8f8eeb83eb0a131 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 29 Oct 2024 20:25:03 +0000 Subject: [PATCH 28/29] Remove Arrow policy example and add a @code docstring to demo the policy --- README.md | 3 +- examples/CMakeLists.txt | 1 - examples/bloom_filter/arrow_policy_example.cu | 94 ------------------- .../bloom_filter/arrow_filter_policy.cuh | 41 ++++++++ 4 files changed, 42 insertions(+), 97 deletions(-) delete mode 100644 examples/bloom_filter/arrow_policy_example.cu diff --git a/README.md b/README.md index ea58d09d2..bbc6fc18b 100644 --- a/README.md +++ b/README.md @@ -254,5 +254,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) -- [Host-bulk APIs (Arrow fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/arrow_policy_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydV2tzGjcU_Ssa8qGQAAtOnIzxY0IebZlm7NR2k8mUDtFqtaDJIm21WmOayX_vuVotLJhkOsUfDKv7PPfce7VfW4UsCmV00Rr9-bWlktZo2G1lXM9LPpetUUuUCW91W4UpraDf0eOpZo_Za5OvrZovHGuLDjsaHD3rsssPkzeTMXt9df3-6np8O7m67JOsl3-nhNSFTFipE2mZW0g2zrnAv3DSZR-kpUDYUX_A2iQwbYWzaatz6q2sTcmWfM20cawsJMyogqUqk0zeC5k7pjQTZplnimsh2Uq5hXcV7Phw2KdgxMSOQ55DI8evtCnJuNuETp-Fc_koilarVZ_7sPvGzqOsEi6id5PXby9v3vYQ-kbtD50BWWbl36WySDxeM54jMsFjxJvxFTOW8bmVOHOGIl9Z5ZSed1lhUrfiVno7iSqcVXHpdsCr40T-TQHAxzWAG9-wyc20xV6NbyY3XW_n4-T216s_btnH8fX1-PJ28vaGXV2jWJdvJlQq_PqZjS8_sd8ml2-6TAI6uJL3uaUsEKoiWGVSYXgj5U4YqanCKnIpVKoEqxnE5uZOWo20WC7tUlVcQ5CJt5OppXLc-WcPkvOuoqme6kdKi6xMJDsTpTBRnBmznKHuTtq-KBcXuzJuYcvCRcKU2vXp8MFRIu_gYnYnhTP2sIi8l6KkwGa5QdHWh6UKVFeCav39GJRBUSRf-sfR44oSLz1VubVmFazO5D0HrhJZVBKxVTJlb-QSiDjLnQQuBeEIfqKyXrdHFAdoxCMPBaug2FBvAsEkURQ9UYsgTWTKy8yxyi3jagm2cMfAZQUwUJ5QPiJRirqgtb0tVA3FXfp-InPeFqHXZZ9F-dpkGUCk-n2GXooWpjA_VzlWYYVUP3tzbkFOs8ysKDESR4CSFypbMwHAHPDZSarSWIInovDeY7ngd8r4bOuu3VFAqkojDXRUGDFjimbkFaiRC3TyHPwu4z6QjKqGjnzIRK04iuUwEUl88jwePB08f3oiBvHx85PjY_FMPD9KTxJ6nj578TQSeR4VVkQ5t-CB2yOmePTuaHjSe3f0dLAdhC8xuyS7pcEVak89TP2TbIpOdctLmxvMFoCerfts4kiM5p7STqJT_NAoFmblAViCbrvlcmyFGQeZpCJA7Yxnc4M5s1huuwsWMRCVbt8ZlXSm-itOyA0GI-KhEcB0uZx9keuCZuE5Gw5-GgwGp2zziaLojP0itSTGhmNG8octubxSO9_afcwG_ePTYGmC5rfOh50qWzi24Fnq7ZExEw48yN9xoB846AXPpw81ijIONSugceQz83HgKx326kOAu9dwZCyKKDTm1nk1B_daEjSn9oLIzIuck_fTjepHWlNZFtZZ4GpN5dCt-1bxT8-lzS3lMa9gB2e2zsJwCf5oZo5GB1ryrA7q4jRE82qn80jd79ADQW2dBZOVs6lHvvLY7IaNq244lPegsTsr1D9y5i7oacJHI4xVyZNZgZUsZ_WcaWRzsQXuJucrvT8sfLQ-082k08l-HT1pmlFX3782ePAt-ClcMhphkTh2doa1-qrMvqB8np2eiQ-d11AdKlCIaQRLU80aHxj3ntDYWfBcLZjRaGdVbQvmu6Fds9tfjzYq9VJq01E_lnO0dqfrNfrwQN-HneCGl9RP-cxLgStNldOmABR9Q21knzQbqpLSGzOVws5RbWAbxraWhxs-DHevsC1ZH6utXUfRDZ46P0QtNia7IEncZbAE21XgXZbyrJC72B1U1E3FnWnSsBFS-b2Udt2YUJurES5SWFtlgUVXMQgjfJMaFP1KwEAvs4RV7pi_nTpbyh5WAdb5nWR-wAKY2_fX58MmKhhndJ0tHkDT3SZeF_a_hOvTrLa1lX7x0E0QLUtR7Q1g6kZFF4RmOxReMbcm5hluqLgTJtxx8NyWwpWw1W2YCVbk_ULFykG1wnUv75_fX18MsCdzpEWTJ9xHUIoYkQMVr6vq3BpZ0RTYFV3wYuEvVygDzXXkexhOvYFTBzj1d-BMMwO4CG0KNoxC_7Bd08vfSdu7ZDtQoOaj0LFEg06HRcFgRb-Ku5Xf9H_43c-k-egHfnXn8HgkUuKlA99rEPxjqlt4nm6f7w88K8EKzQb4-W2q8cpJt1zcl-32zbSl74QYHh2XQxyb3FWvra0eTJ2LJ0-GL1iPW7E4L5azFwPW62G7u55fNLgv9TK-jP27bKbihk0hRIaHd9XbJx4gY_2l9a1bn2Nd7ZwDvda3v_zfv3UOROY=)) \ No newline at end of file +- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) \ No newline at end of file diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 36a5964a8..8e7be947a 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -47,4 +47,3 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/ 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") -ConfigureExample(BLOOM_FILTER_ARROW_POLICY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/arrow_policy_example.cu") diff --git a/examples/bloom_filter/arrow_policy_example.cu b/examples/bloom_filter/arrow_policy_example.cu deleted file mode 100644 index ce3c2a720..000000000 --- a/examples/bloom_filter/arrow_policy_example.cu +++ /dev/null @@ -1,94 +0,0 @@ -/* - * 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 - -#include -#include -#include -#include - -#include - -/** - * @file arrow_policy_example.cu - * @brief Demonstrates usage of an arrow-compatible bloom filter - * - * In addition to the default policy aimed at achieving the speed of light - * performance on the device, `cuCollections` offers an `arrow_filter_policy` - * that allows users to easily create a bloom filter that mimics the behavior - * of the bloom filter defined in Apache Arrow: - * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230. - * - * @note This example is for demonstration purposes only. It is not intended to show the most - * performant way to do the example algorithm. - */ - -int main(void) -{ - int constexpr num_keys = 10'000; ///< Generate 10'000 keys - int constexpr num_tp = num_keys * 0.5; ///< Insert the first half keys into the filter. - int constexpr num_tn = num_keys - num_tp; - int constexpr sub_filters = 200; ///< 200 sub-filters per bloom filter - - // key type for bloom filter - using key_type = int; - - // We will use the Arrow filter policy for bloom filter fingerprint generation - using policy_type = cuco::arrow_filter_policy; - // Bloom filter type with Arrow filter policy - using filter_type = - cuco::bloom_filter, cuda::thread_scope_device, policy_type>; - - // Spawn a bloom filter with arrow policy and 200 sub-filters. - filter_type filter{sub_filters}; - - std::cout << "Bulk insert into bloom filter with Arrow fingerprint generation policy: " - << std::endl; - - thrust::device_vector keys(num_keys); - thrust::sequence(keys.begin(), keys.end(), 1); - - auto tp_begin = keys.begin(); - auto tp_end = tp_begin + num_tp; - auto tn_begin = tp_end; - auto tn_end = keys.end(); - - // Insert the first half of the keys. - filter.add(tp_begin, tp_end); - - thrust::device_vector tp_result(num_tp, false); - thrust::device_vector tn_result(num_keys - num_tp, false); - - // Query the filter for the previously inserted keys. - // This should result in a true-positive rate of TPR=1. - filter.contains(tp_begin, tp_end, tp_result.begin()); - - // Query the filter for the keys that are not present in the filter. - // Since bloom filters are probalistic data structures, the filter - // exhibits a false-positive rate FPR>0 depending on the number of bits in - // the filter and the number of hashes used per key. - filter.contains(tn_begin, tn_end, tn_result.begin()); - - float tp_rate = - float(thrust::count(thrust::device, tp_result.begin(), tp_result.end(), true)) / float(num_tp); - float fp_rate = - float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn); - - std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl; - - return 0; -} diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index 4fc5efc29..bb8ff307a 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -34,6 +34,47 @@ namespace cuco::detail { * Reference: * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 * + * Example: + * @code{.cpp} + * template + * void bulk_insert_and_eval_arrow_policy_bloom_filter(device_vector const& positive_keys, + * device_vector const& negative_keys) + * { + * // Arrow filter policy type + * using policy_type = cuco::arrow_filter_policy; + * + * // Create a bloom filter with arrow_filter_policy + * cuco::bloom_filter, + * cuda::thread_scope_device, policy_type> filter{NUM_FILTER_BLOCKS}; + * + * // Add positive keys to the bloom filter + * filter.add(positive_keys.begin(), positive_keys.end()); + * + * // Number of true positives and true negatives + * auto const num_tp = positive_keys.size(); + * auto const num_tn = negative_keys.size(); + * + * // Vectors to store true positive and true negative filter query results. + * thrust::device_vector true_positive_result(num_tp, false); + * thrust::device_vector true_negative_result(num_tn, false); + * + * // Query the bloom filter for the inserted positive keys. + * filter.contains(positive_keys.begin(), positive_keys.end(), true_positive_result.begin()); + * + * // Query the bloom filter for the non-inserted true negative_keys. + * filter.contains(negative_keys.begin(), negative_keys.end(), true_negative_result.begin()); + * + * // We should see a true-positive rate of 1. + * float true_positive_rate = float(thrust::count(thrust::device, + * true_positive_result.begin(), true_positive_result.end(), true)) / float(num_tp); + * + * // Since bloom filters are probalistic data structures, we may see a false-positive rate > 0 + * // depending on the number of bits in the filter and the number of hashes used per key. + * float false_positive_rate = float(thrust::count(thrust::device, + * true_negative_result.begin(), true_negative_result.end(), true)) / float(num_tn); + * } + * @endcode + * * @tparam Key The type of the values to generate a fingerprint for. */ template From fd9861fdedddf9509d7ac97ea7ec208d31152258 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Wed, 30 Oct 2024 00:19:01 +0000 Subject: [PATCH 29/29] docstring updates --- .../bloom_filter/arrow_filter_policy.cuh | 24 ++++++++++--------- 1 file changed, 13 insertions(+), 11 deletions(-) diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index bb8ff307a..23b95793e 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -36,40 +36,42 @@ namespace cuco::detail { * * Example: * @code{.cpp} - * template + * template * void bulk_insert_and_eval_arrow_policy_bloom_filter(device_vector const& positive_keys, * device_vector const& negative_keys) * { - * // Arrow filter policy type * using policy_type = cuco::arrow_filter_policy; * - * // Create a bloom filter with arrow_filter_policy + * // Warn or throw if the number of filter blocks is greater than maximum used by Arrow policy. + * static_assert(NUM_FILTER_BLOCKS <= policy_type::max_filter_blocks, "NUM_FILTER_BLOCKS must be + * in range: [1, 4194304]"); + * + * // Create a bloom filter with Arrow policy * cuco::bloom_filter, * cuda::thread_scope_device, policy_type> filter{NUM_FILTER_BLOCKS}; * * // Add positive keys to the bloom filter * filter.add(positive_keys.begin(), positive_keys.end()); * - * // Number of true positives and true negatives * auto const num_tp = positive_keys.size(); * auto const num_tn = negative_keys.size(); * - * // Vectors to store true positive and true negative filter query results. + * // Vectors to store query results. * thrust::device_vector true_positive_result(num_tp, false); * thrust::device_vector true_negative_result(num_tn, false); * - * // Query the bloom filter for the inserted positive keys. + * // Query the bloom filter for the inserted keys. * filter.contains(positive_keys.begin(), positive_keys.end(), true_positive_result.begin()); * - * // Query the bloom filter for the non-inserted true negative_keys. - * filter.contains(negative_keys.begin(), negative_keys.end(), true_negative_result.begin()); - * * // We should see a true-positive rate of 1. * float true_positive_rate = float(thrust::count(thrust::device, * true_positive_result.begin(), true_positive_result.end(), true)) / float(num_tp); * - * // Since bloom filters are probalistic data structures, we may see a false-positive rate > 0 - * // depending on the number of bits in the filter and the number of hashes used per key. + * // Query the bloom filter for the non-inserted keys. + * filter.contains(negative_keys.begin(), negative_keys.end(), true_negative_result.begin()); + * + * // We may see a false-positive rate > 0 depending on the number of bits in the + * // filter and the number of hashes used per key. * float false_positive_rate = float(thrust::count(thrust::device, * true_negative_result.begin(), true_negative_result.end(), true)) / float(num_tn); * }