Add Arrow bloom filter policy

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 <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_bf_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/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")

examples/bloom_filter/arrow_bf_host_bulk_example.cu

@@ -0,0 +1,75 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cuco/bloom_filter.cuh>
+
+#include <thrust/count.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/sequence.h>
+
+#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;
+
+  using policy_type = cuco::arrow_bf_policy<int>;
+  using filter_type =
+    cuco::bloom_filter<int, cuco::extent<size_t>, 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<int> 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<bool> tp_result(num_tp, false);
+  thrust::device_vector<bool> 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;
+}

include/cuco/arrow_bf_policy.cuh

@@ -0,0 +1,114 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <cuco/detail/bloom_filter/arrow_bf_policy_impl.cuh>
+
+#include <cstdint>
+
+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 T>
+class arrow_bf_policy {
+  using impl_type = cuco::detail::arrow_bf_policy_impl<T>;
+
+ 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 <class Extent>
+  __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 <cuco/detail/bloom_filter/arrow_bf_policy.inl>

include/cuco/bloom_filter.cuh

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cuco/arrow_bf_policy.cuh>
 #include <cuco/bloom_filter_policy.cuh>
 #include <cuco/bloom_filter_ref.cuh>
 #include <cuco/detail/storage/storage_base.cuh>

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 <cstdint>
+
+namespace cuco {
+
+template <class T>
+__host__ __device__ constexpr arrow_bf_policy<T>::arrow_bf_policy(std::uint32_t num_blocks,
+                                                                  hasher hash)
+  : impl_{num_blocks, hash}
+{
+}
+
+template <class T>
+__device__ constexpr typename arrow_bf_policy<T>::hash_result_type arrow_bf_policy<T>::hash(
+  typename arrow_bf_policy<T>::hash_argument_type const& key) const
+{
+  return impl_.hash(key);
+}
+
+template <class T>
+template <class Extent>
+__device__ constexpr auto arrow_bf_policy<T>::block_index(
+  typename arrow_bf_policy<T>::hash_result_type hash, Extent num_blocks) const
+{
+  return impl_.block_index(hash, num_blocks);
+}
+
+template <class T>
+__device__ constexpr typename arrow_bf_policy<T>::word_type arrow_bf_policy<T>::word_pattern(
+  arrow_bf_policy<T>::hash_result_type hash, std::uint32_t word_index) const
+{
+  return impl_.word_pattern(hash, word_index);
+}
+
+}  // namespace cuco