Add Arrow bloom filter policy

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))
+- [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=))

benchmarks/bloom_filter/add_bench.cu

@@ -41,9 +41,9 @@ template <typename Key, typename Hash, typename Word, nvbench::int32_t WordsPerB
 void bloom_filter_add(nvbench::state& state,
                       nvbench::type_list<Key, Hash, Word, nvbench::enum_type<WordsPerBlock>, Dist>)
 {
-  using policy_type = cuco::bloom_filter_policy<rebind_hasher_t<Hash, Key>,
-                                                Word,
-                                                static_cast<std::uint32_t>(WordsPerBlock)>;
+  using policy_type = cuco::default_filter_policy<rebind_hasher_t<Hash, Key>,
+                                                  Word,
+                                                  static_cast<std::uint32_t>(WordsPerBlock)>;
   using filter_type =
     cuco::bloom_filter<Key, cuco::extent<size_t>, cuda::thread_scope_device, policy_type>;
 

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<rebind_hasher_t<Hash, Key>,
-                                                Word,
-                                                static_cast<std::uint32_t>(WordsPerBlock)>;
+  using policy_type = cuco::default_filter_policy<rebind_hasher_t<Hash, Key>,
+                                                  Word,
+                                                  static_cast<std::uint32_t>(WordsPerBlock)>;
   using filter_type =
     cuco::bloom_filter<Key, cuco::extent<size_t>, cuda::thread_scope_device, policy_type>;
 
@@ -88,6 +88,57 @@ void bloom_filter_contains(
   });
 }
 
+/**
+ * @brief A benchmark evaluating `cuco::bloom_filter::contains_async` performance with
+ * `arrow_filter_policy`
+ */
+template <typename Key, typename Dist>
+void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list<Key, Dist>)
+{
+  // cudaDeviceSetLimit(cudaLimitMaxL2FetchGranularity, 32); // slightly improves peformance if
+  // filter block fits into a 32B sector
+  using policy_type = cuco::arrow_filter_policy<Key>;
+  using filter_type =
+    cuco::bloom_filter<Key, cuco::extent<size_t>, 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<uint32_t>(num_sub_filters)};
+  } catch (std::exception const& e) {
+    state.skip(e.what());  // skip invalid configurations
+  }
+
+  thrust::device_vector<Key> keys(num_keys);
+  thrust::device_vector<bool> result(num_keys, false);
+
+  key_generator gen;
+  gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
+
+  state.add_element_count(num_keys);
+
+  filter_type filter{num_sub_filters, {}, {static_cast<uint32_t>(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<defaults::BF_KEY>,
                                       nvbench::type_list<defaults::BF_HASH>,
@@ -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});
+  .add_int64_axis("FilterSizeMB", {defaults::BF_SIZE_MB});
+
+NVBENCH_BENCH_TYPES(arrow_bloom_filter_contains,
+                    NVBENCH_TYPE_AXES(nvbench::type_list<defaults::BF_KEY>,
+                                      nvbench::type_list<distribution::unique>))
+  .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);

examples/bloom_filter/host_bulk_example.cu

@@ -23,17 +23,18 @@
 
 #include <iostream>
 
-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;
+// key type for bloom filter
+using key_type = int;
 
-  // Spawn a filter with 200 sub-filters.
-  cuco::bloom_filter<int> filter{200};
+// Function to bulk insert to the provided bloom filter and evaluate FPR.
+template <typename FilterType>
+void bulk_insert_and_eval_bloom_filter(FilterType& filter, int num_keys)
+{
+  // Insert the first half keys into the filter.
+  int const num_tp = num_keys * 0.5;
+  int const num_tn = num_keys - num_tp;
 
-  thrust::device_vector<int> keys(num_keys);
+  thrust::device_vector<key_type> keys(num_keys);
   thrust::sequence(keys.begin(), keys.end(), 1);
 
   auto tp_begin = keys.begin();
@@ -63,6 +64,32 @@ int main(void)
     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<key_type> 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<key_type>;
+  // bloom filter with arrow policy type
+  using arrow_policy_filter_type = cuco::
+    bloom_filter<key_type, cuco::extent<size_t>, 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}};
+
+  // 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;
 }

include/cuco/bloom_filter.cuh

@@ -16,7 +16,7 @@
 
 #pragma once
 
-#include <cuco/bloom_filter_policy.cuh>
+#include <cuco/bloom_filter_policies.cuh>
 #include <cuco/bloom_filter_ref.cuh>
 #include <cuco/detail/storage/storage_base.cuh>
 #include <cuco/extent.cuh>
@@ -55,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 <class Key,
           class Extent             = cuco::extent<std::size_t>,
           cuda::thread_scope Scope = cuda::thread_scope_device,
-          class Policy    = cuco::bloom_filter_policy<cuco::xxhash_64<Key>, std::uint32_t, 8>,
+          class Policy    = cuco::default_filter_policy<cuco::xxhash_64<Key>, std::uint32_t, 8>,
           class Allocator = cuco::cuda_allocator<cuda::std::byte>>
 class bloom_filter {
  public:
@@ -109,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
    */

include/cuco/bloom_filter_policy.cuh → include/cuco/bloom_filter_policies.cuh

@@ -16,14 +16,25 @@
 
 #pragma once
 
-#include <cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh>
+#include <cuco/detail/bloom_filter/arrow_filter_policy.cuh>
+#include <cuco/detail/bloom_filter/default_filter_policy_impl.cuh>
 
 #include <cstdint>
 
 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 <class Key>
+using arrow_filter_policy = detail::arrow_filter_policy<Key>;
+
+/**
+ * @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 Hash, class Word, std::uint32_t WordsPerBlock>
-class bloom_filter_policy {
-  using impl_type = cuco::detail::bloom_filter_policy_impl<Hash, Word, WordsPerBlock>;
+class default_filter_policy {
+  using impl_type = cuco::detail::default_filter_policy_impl<Hash, Word, WordsPerBlock>;
 
  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 <cuco/detail/bloom_filter/bloom_filter_policy.inl>
+#include <cuco/detail/bloom_filter/default_filter_policy.inl>

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 Key, class Extent, cuda::thread_scope Scope, class Policy>
 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,