-
Notifications
You must be signed in to change notification settings - Fork 211
CAGRA Bloom Filter #2236
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
CAGRA Bloom Filter #2236
Changes from all commits
21e5f20
16946a4
a3d7b96
80027b5
56d959a
93db300
4e6aa9c
cf40ea1
c575717
2bebdb9
4b33230
863e9fb
0569eb7
241ca5d
dfe8f9f
5106f94
d59ca43
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,89 @@ | ||
| /* | ||
| * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. | ||
| * SPDX-License-Identifier: Apache-2.0 | ||
| */ | ||
|
|
||
| #pragma once | ||
|
|
||
| #include <cuvs/core/export.hpp> | ||
| #include <raft/core/device_mdarray.hpp> | ||
| #include <raft/core/resources.hpp> | ||
|
|
||
| #include <cstddef> | ||
| #include <cstdint> | ||
| #include <memory> | ||
|
|
||
| namespace CUVS_EXPORT cuvs { | ||
| namespace core { | ||
|
|
||
| /** | ||
| * @brief cuVS-owned Bloom filter wrapper with opaque implementation. | ||
| * | ||
| * This class intentionally hides cuCollections types from the cuVS public API. | ||
| * The wrapper supports the expected bulk host APIs used by ANN workflows. | ||
| */ | ||
| class CUVS_EXPORT bloom_filter { | ||
| private: | ||
| struct impl; | ||
|
|
||
| public: | ||
| using key_type = std::uint32_t; | ||
|
|
||
| /** | ||
| * @brief Construct a Bloom filter with user-facing quality knobs. | ||
| * | ||
| * @p dataset_rows is the number of rows in the indexed dataset. The filter uses it with | ||
| * @p filtering_rate to estimate the number of inserted valid ids and compute a target filter | ||
| * size that satisfies the requested false-positive rate. | ||
| * | ||
| * The primary tuning knobs are: | ||
| * - @p filtering_rate: expected fraction of dataset rows that will be inserted as valid ids. | ||
| * - @p target_false_positive_rate: desired Bloom filter false-positive probability. | ||
| * | ||
| * Sizing math used internally: | ||
| * - `expected_insertions = ceil(dataset_rows * filtering_rate)` | ||
| * - `required_bits = -expected_insertions * ln(target_false_positive_rate) / (ln(2)^2)` | ||
| * - `required_blocks = ceil(required_bits / 256)` (default cuco policy uses 256-bit blocks) | ||
| * | ||
| * Practical knob behavior: | ||
| * - Lower @p target_false_positive_rate -> larger filter, fewer false positives, typically higher | ||
| * filtered-search recall. | ||
| * - Higher @p filtering_rate -> larger filter for the same target false-positive rate. | ||
| */ | ||
| bloom_filter(raft::resources const& res, | ||
| std::size_t dataset_rows, | ||
| float filtering_rate = 1.0f, | ||
| float target_false_positive_rate = 0.01f); | ||
| ~bloom_filter(); | ||
|
|
||
| bloom_filter(bloom_filter const&) = delete; | ||
| bloom_filter& operator=(bloom_filter const&) = delete; | ||
| bloom_filter(bloom_filter&&) noexcept; | ||
| bloom_filter& operator=(bloom_filter&&) noexcept; | ||
|
|
||
| void clear(raft::resources const& res); | ||
| void clear_async(raft::resources const& res); | ||
|
|
||
| void add(raft::resources const& res, raft::device_vector_view<const key_type, int64_t> keys); | ||
| void add_async(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys); | ||
|
|
||
| void contains(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys, | ||
| raft::device_vector_view<std::uint8_t, int64_t> output) const; | ||
| void contains_async(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys, | ||
| raft::device_vector_view<std::uint8_t, int64_t> output) const; | ||
|
|
||
| [[nodiscard]] std::size_t num_blocks() const noexcept; | ||
| [[nodiscard]] float estimate_filtering_rate(raft::resources const& res, | ||
| std::size_t dataset_rows) const; | ||
|
|
||
| private: | ||
| friend impl const& get_bloom_filter_impl(bloom_filter const& filter) noexcept; | ||
|
|
||
| std::unique_ptr<impl> impl_; | ||
| }; | ||
|
|
||
| } // namespace core | ||
| } // namespace CUVS_EXPORT cuvs |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,184 @@ | ||
| /* | ||
| * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. | ||
| * SPDX-License-Identifier: Apache-2.0 | ||
| */ | ||
|
|
||
| #include "../neighbors/detail/sample_filter_data.cuh" | ||
| #include <cuvs/core/bloom_filter.hpp> | ||
|
|
||
| #include <cuco/bloom_filter.cuh> | ||
|
|
||
| #include <raft/core/error.hpp> | ||
| #include <raft/core/resource/cuda_stream.hpp> | ||
| #include <raft/core/resource/thrust_policy.hpp> | ||
|
|
||
| #include <rmm/device_uvector.hpp> | ||
| #include <thrust/count.h> | ||
| #include <thrust/iterator/counting_iterator.h> | ||
|
|
||
| #include <algorithm> | ||
| #include <cmath> | ||
| #include <utility> | ||
|
|
||
| namespace cuvs::core { | ||
|
|
||
| std::size_t compute_num_blocks_from_rates(std::size_t dataset_rows, | ||
| float filtering_rate, | ||
| float target_false_positive_rate) | ||
| { | ||
| RAFT_EXPECTS(dataset_rows > 0, | ||
| "dataset_rows must be greater than zero when deriving bloom size."); | ||
| RAFT_EXPECTS(filtering_rate > 0.0f && filtering_rate <= 1.0f, | ||
| "filtering_rate must be in (0, 1]."); | ||
| RAFT_EXPECTS(target_false_positive_rate > 0.0f && target_false_positive_rate < 1.0f, | ||
| "target_false_positive_rate must be in (0, 1)."); | ||
|
|
||
| // Bloom sizing: m = -n * ln(p) / (ln(2)^2), then blocks = ceil(m / 256 bits-per-block). | ||
| constexpr double kBitsPerBlock = 256.0; | ||
| constexpr double kLn2 = 0.6931471805599453; | ||
| constexpr double kLn2Sq = kLn2 * kLn2; | ||
|
|
||
| auto expected_insertions = std::max<std::size_t>( | ||
| 1, static_cast<std::size_t>(std::ceil(static_cast<double>(dataset_rows) * filtering_rate))); | ||
| auto required_bits = -static_cast<double>(expected_insertions) * | ||
| std::log(static_cast<double>(target_false_positive_rate)) / kLn2Sq; | ||
| return std::max<std::size_t>(1, | ||
| static_cast<std::size_t>(std::ceil(required_bits / kBitsPerBlock))); | ||
| } | ||
|
|
||
| struct bloom_filter::impl { | ||
| using key_type = bloom_filter::key_type; | ||
| using cuco_filter_type = cuco::bloom_filter<key_type>; | ||
| using sample_filter_payload = cuvs::neighbors::detail::bloom_filter_data_t<key_type>; | ||
|
|
||
| cuco_filter_type filter; | ||
| std::size_t dataset_rows; | ||
| float filtering_rate; | ||
| float target_false_positive_rate; | ||
|
|
||
| impl(raft::resources const& res, | ||
| std::size_t num_blocks, | ||
| std::size_t dataset_rows_, | ||
| float filtering_rate_, | ||
| float target_false_positive_rate_) | ||
| : filter(num_blocks, {}, {}, {}, raft::resource::get_cuda_stream(res)), | ||
| dataset_rows(dataset_rows_), | ||
| filtering_rate(filtering_rate_), | ||
| target_false_positive_rate(target_false_positive_rate_) | ||
| { | ||
| } | ||
|
|
||
| void validate_keys(raft::device_vector_view<const key_type, int64_t> keys) const | ||
| { | ||
| if (keys.extent(0) == 0) { return; } | ||
| auto num_insertions = static_cast<std::size_t>(keys.extent(0)); | ||
| RAFT_EXPECTS(num_insertions <= dataset_rows, | ||
| "Number of keys to add must not exceed dataset_rows."); | ||
| } | ||
| }; | ||
|
|
||
| bloom_filter::bloom_filter(raft::resources const& res, | ||
| std::size_t dataset_rows, | ||
| float filtering_rate, | ||
| float target_false_positive_rate) | ||
| : impl_(std::make_unique<impl>( | ||
| res, | ||
| compute_num_blocks_from_rates(dataset_rows, filtering_rate, target_false_positive_rate), | ||
| dataset_rows, | ||
| filtering_rate, | ||
| target_false_positive_rate)) | ||
| { | ||
| } | ||
|
|
||
| bloom_filter::~bloom_filter() = default; | ||
| bloom_filter::bloom_filter(bloom_filter&&) noexcept = default; | ||
| bloom_filter& bloom_filter::operator=(bloom_filter&&) noexcept = default; | ||
|
|
||
| void bloom_filter::clear(raft::resources const& res) | ||
| { | ||
| impl_->filter.clear(raft::resource::get_cuda_stream(res)); | ||
| } | ||
|
|
||
| void bloom_filter::clear_async(raft::resources const& res) | ||
| { | ||
| impl_->filter.clear_async(raft::resource::get_cuda_stream(res)); | ||
| } | ||
|
|
||
| void bloom_filter::add(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys) | ||
| { | ||
| auto stream = raft::resource::get_cuda_stream(res); | ||
| impl_->validate_keys(keys); | ||
| impl_->filter.add(keys.data_handle(), keys.data_handle() + keys.extent(0), stream); | ||
| } | ||
|
|
||
| void bloom_filter::add_async(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys) | ||
| { | ||
| auto stream = raft::resource::get_cuda_stream(res); | ||
| impl_->validate_keys(keys); | ||
| impl_->filter.add_async(keys.data_handle(), keys.data_handle() + keys.extent(0), stream); | ||
| } | ||
|
|
||
| void bloom_filter::contains(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys, | ||
| raft::device_vector_view<std::uint8_t, int64_t> output) const | ||
| { | ||
| RAFT_EXPECTS(output.extent(0) >= keys.extent(0), | ||
| "Bloom filter contains output size must be at least keys.size()."); | ||
| impl_->filter.contains(keys.data_handle(), | ||
|
divyegala marked this conversation as resolved.
|
||
| keys.data_handle() + keys.extent(0), | ||
| output.data_handle(), | ||
| raft::resource::get_cuda_stream(res)); | ||
| } | ||
|
|
||
| void bloom_filter::contains_async(raft::resources const& res, | ||
| raft::device_vector_view<const key_type, int64_t> keys, | ||
| raft::device_vector_view<std::uint8_t, int64_t> output) const | ||
| { | ||
| RAFT_EXPECTS(output.extent(0) >= keys.extent(0), | ||
| "Bloom filter contains output size must be at least keys.size()."); | ||
| impl_->filter.contains_async(keys.data_handle(), | ||
| keys.data_handle() + keys.extent(0), | ||
| output.data_handle(), | ||
| raft::resource::get_cuda_stream(res)); | ||
| } | ||
|
|
||
| std::size_t bloom_filter::num_blocks() const noexcept { return impl_->filter.block_extent(); } | ||
|
|
||
| float bloom_filter::estimate_filtering_rate(raft::resources const& res, | ||
| std::size_t dataset_rows) const | ||
| { | ||
| if (dataset_rows == 0) { return 0.0f; } | ||
| auto stream = raft::resource::get_cuda_stream(res); | ||
| auto policy = raft::resource::get_thrust_policy(res); | ||
|
|
||
| rmm::device_uvector<std::uint8_t> hits(dataset_rows, stream); | ||
|
|
||
| auto first_id = thrust::counting_iterator<key_type>(0); | ||
| impl_->filter.contains_async(first_id, first_id + dataset_rows, hits.data(), stream); | ||
|
Contributor
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. The default CAGRA path calls this estimate on every search, so each search allocates an N-byte buffer, probes every index row, and synchronizes before ANN search begins, that could be costly, no? Could we cache an estimate invalidated by add/clear, retain insertion metadata, or use bounded sampling instead?
Contributor
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Also fixed by bounding to size ahead of time in the constructor. Very, very good catch. |
||
|
|
||
| auto positives = thrust::count_if( | ||
| policy, hits.begin(), hits.end(), [] __device__(std::uint8_t v) { return v != 0; }); | ||
| raft::resource::sync_stream(res); | ||
| auto filtering_rate = static_cast<float>(dataset_rows - static_cast<std::size_t>(positives)) / | ||
| static_cast<float>(dataset_rows); | ||
| return std::clamp(filtering_rate, 0.0f, 0.999f); | ||
| } | ||
|
|
||
| auto get_bloom_filter_impl(bloom_filter const& filter) noexcept -> bloom_filter::impl const& | ||
| { | ||
| return *filter.impl_; | ||
| } | ||
|
|
||
| } // namespace cuvs::core | ||
|
|
||
| namespace cuvs::neighbors::detail { | ||
|
|
||
| bloom_filter_data_t<std::uint32_t> bloom_filter_factory::make( | ||
| cuvs::core::bloom_filter const& filter) | ||
| { | ||
| return bloom_filter_data_t<std::uint32_t>{get_bloom_filter_impl(filter).filter.ref()}; | ||
| } | ||
|
|
||
| } // namespace cuvs::neighbors::detail | ||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The referenced core filter must outlive the adapter and all searches, and must not be moved or mutated concurrently, right? The stored pointer otherwise becomes invalid or refers to a moved-from object without a diagnostic, so we should document this if I'm right.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yes, I'll add the note but it is a problem with all our filters as we separate the implementation from the interface.