diff --git a/README.md b/README.md
index c11eb5c7d..26d71870e 100644
--- a/README.md
+++ b/README.md
@@ -242,4 +242,11 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 - [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/hyperloglog/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/z/G4qdcTezE))
 - [Device-ref APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/hyperloglog/device_ref_example.cu) (see [live example in godbolt](https://godbolt.org/z/n88713o4n))
 
+### `bloom_filter`
+
+`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/z/EY7T5v5aE))
+
 
diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt
index 56205692e..27b858898 100644
--- a/benchmarks/CMakeLists.txt
+++ b/benchmarks/CMakeLists.txt
@@ -98,3 +98,9 @@ ConfigureBench(HASH_FUNCTION_BENCH
 # - hyperloglog benchmarks -----------------------------------------------------------
 ConfigureBench(HYPERLOGLOG_BENCH
   hyperloglog/hyperloglog_bench.cu)
+
+###################################################################################################
+# - bloom_filter benchmarks -----------------------------------------------------------------------
+ConfigureBench(BLOOM_FILTER_BENCH
+  bloom_filter/add_bench.cu
+  bloom_filter/contains_bench.cu)
diff --git a/benchmarks/benchmark_defaults.hpp b/benchmarks/benchmark_defaults.hpp
index c01e21ca8..1d689d55a 100644
--- a/benchmarks/benchmark_defaults.hpp
+++ b/benchmarks/benchmark_defaults.hpp
@@ -16,6 +16,8 @@
 
 #pragma once
 
+#include <cuco/hash_functions.cuh>
+
 #include <nvbench/nvbench.cuh>
 
 #include <cstdint>
@@ -25,6 +27,12 @@ namespace cuco::benchmark::defaults {
 
 using KEY_TYPE_RANGE   = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
 using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
+using HASH_RANGE       = nvbench::type_list<cuco::identity_hash<char>,
+                                            cuco::xxhash_32<char>,
+                                            cuco::xxhash_64<char>,
+                                            cuco::murmurhash3_32<char>>;  //,
+// cuco::murmurhash3_x86_128<char>,
+// cuco::murmurhash3_x64_128<char>>; // TODO handle tuple-like hash value
 
 auto constexpr N             = 100'000'000;
 auto constexpr OCCUPANCY     = 0.5;
diff --git a/benchmarks/benchmark_utils.hpp b/benchmarks/benchmark_utils.hpp
index 392cafe06..680345ebe 100644
--- a/benchmarks/benchmark_utils.hpp
+++ b/benchmarks/benchmark_utils.hpp
@@ -39,6 +39,17 @@ auto dist_from_state(nvbench::state const& state)
   }
 }
 
+template <typename T, typename NewType>
+struct rebind_hasher;
+
+template <template <typename> class Template, typename OldType, typename NewType>
+struct rebind_hasher<Template<OldType>, NewType> {
+  using type = Template<NewType>;
+};
+
+template <typename T, typename NewType>
+using rebind_hasher_t = typename rebind_hasher<T, NewType>::type;
+
 }  // namespace cuco::benchmark
 
 NVBENCH_DECLARE_TYPE_STRINGS(cuco::utility::distribution::unique, "UNIQUE", "distribution::unique");
diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu
new file mode 100644
index 000000000..b60768d0b
--- /dev/null
+++ b/benchmarks/bloom_filter/add_bench.cu
@@ -0,0 +1,117 @@
+/*
+ * 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 "defaults.hpp"
+#include "utils.hpp"
+
+#include <benchmark_defaults.hpp>
+#include <benchmark_utils.hpp>
+
+#include <cuco/bloom_filter.cuh>
+#include <cuco/utility/key_generator.cuh>
+
+#include <nvbench/nvbench.cuh>
+
+#include <cuda/std/limits>
+#include <thrust/device_vector.h>
+
+#include <cstdint>
+#include <exception>
+
+using namespace cuco::benchmark;  // defaults, dist_from_state, rebind_hasher_t, add_fpr_summary
+using namespace cuco::utility;    // key_generator, distribution
+
+/**
+ * @brief A benchmark evaluating `cuco::bloom_filter::add_async` performance
+ */
+template <typename Key, typename Hash, typename Word, nvbench::int32_t WordsPerBlock, typename Dist>
+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 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");
+  auto const pattern_bits   = state.get_int64("PatternBits");
+
+  try {
+    auto const policy = policy_type{static_cast<uint32_t>(pattern_bits)};
+  } catch (std::exception const& e) {
+    state.skip(e.what());  // skip invalid configurations
+  }
+
+  std::size_t const num_sub_filters =
+    (filter_size_mb * 1024 * 1024) /
+    (sizeof(typename filter_type::word_type) * filter_type::words_per_block);
+
+  thrust::device_vector<Key> keys(num_keys);
+
+  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>(pattern_bits)}};
+
+  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<defaults::BF_KEY>,
+                                      nvbench::type_list<defaults::BF_HASH>,
+                                      nvbench::type_list<defaults::BF_WORD>,
+                                      nvbench::enum_type_list<defaults::BF_WORDS_PER_BLOCK>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_add_unique_size")
+  .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_RANGE_CACHE)
+  .add_int64_axis("PatternBits", {defaults::BF_PATTERN_BITS});
+
+NVBENCH_BENCH_TYPES(bloom_filter_add,
+                    NVBENCH_TYPE_AXES(nvbench::type_list<defaults::BF_KEY>,
+                                      defaults::HASH_RANGE,
+                                      nvbench::type_list<defaults::BF_WORD>,
+                                      nvbench::enum_type_list<defaults::BF_WORDS_PER_BLOCK>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_add_unique_hash")
+  .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("PatternBits", {defaults::BF_PATTERN_BITS});
+
+NVBENCH_BENCH_TYPES(bloom_filter_add,
+                    NVBENCH_TYPE_AXES(nvbench::type_list<defaults::BF_KEY>,
+                                      nvbench::type_list<defaults::BF_HASH>,
+                                      nvbench::type_list<nvbench::uint32_t, nvbench::uint64_t>,
+                                      nvbench::enum_type_list<1, 2, 4, 8>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_add_unique_block_dim")
+  .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("PatternBits", {defaults::BF_PATTERN_BITS});
diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu
new file mode 100644
index 000000000..8b9abcb68
--- /dev/null
+++ b/benchmarks/bloom_filter/contains_bench.cu
@@ -0,0 +1,122 @@
+/*
+ * 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 "defaults.hpp"
+#include "utils.hpp"
+
+#include <benchmark_defaults.hpp>
+#include <benchmark_utils.hpp>
+
+#include <cuco/bloom_filter.cuh>
+#include <cuco/utility/key_generator.cuh>
+
+#include <nvbench/nvbench.cuh>
+
+#include <cuda/std/limits>
+#include <thrust/device_vector.h>
+
+#include <exception>
+
+using namespace cuco::benchmark;  // defaults, dist_from_state, rebind_hasher_t, add_fpr_summary
+using namespace cuco::utility;    // key_generator, distribution
+
+/**
+ * @brief A benchmark evaluating `cuco::bloom_filter::contains_async` performance
+ */
+template <typename Key, typename Hash, typename Word, nvbench::int32_t WordsPerBlock, typename Dist>
+void bloom_filter_contains(
+  nvbench::state& state,
+  nvbench::type_list<Key, Hash, Word, nvbench::enum_type<WordsPerBlock>, Dist>)
+{
+  // 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 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");
+  auto const pattern_bits   = state.get_int64("PatternBits");
+
+  try {
+    auto const policy = policy_type{static_cast<uint32_t>(pattern_bits)};
+  } catch (std::exception const& e) {
+    state.skip(e.what());  // skip invalid configurations
+  }
+
+  std::size_t const num_sub_filters =
+    (filter_size_mb * 1024 * 1024) /
+    (sizeof(typename filter_type::word_type) * filter_type::words_per_block);
+
+  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>(pattern_bits)}};
+
+  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>,
+                                      nvbench::type_list<defaults::BF_WORD>,
+                                      nvbench::enum_type_list<defaults::BF_WORDS_PER_BLOCK>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_contains_unique_size")
+  .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_RANGE_CACHE)
+  .add_int64_axis("PatternBits", {defaults::BF_PATTERN_BITS});
+
+NVBENCH_BENCH_TYPES(bloom_filter_contains,
+                    NVBENCH_TYPE_AXES(nvbench::type_list<defaults::BF_KEY>,
+                                      defaults::HASH_RANGE,
+                                      nvbench::type_list<defaults::BF_WORD>,
+                                      nvbench::enum_type_list<defaults::BF_WORDS_PER_BLOCK>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_contains_unique_hash")
+  .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("PatternBits", {defaults::BF_PATTERN_BITS});
+
+NVBENCH_BENCH_TYPES(bloom_filter_contains,
+                    NVBENCH_TYPE_AXES(nvbench::type_list<defaults::BF_KEY>,
+                                      nvbench::type_list<defaults::BF_HASH>,
+                                      nvbench::type_list<nvbench::uint32_t, nvbench::uint64_t>,
+                                      nvbench::enum_type_list<1, 2, 4, 8>,
+                                      nvbench::type_list<distribution::unique>))
+  .set_name("bloom_filter_contains_unique_block_dim")
+  .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("PatternBits", {defaults::BF_PATTERN_BITS});
\ No newline at end of file
diff --git a/benchmarks/bloom_filter/defaults.hpp b/benchmarks/bloom_filter/defaults.hpp
new file mode 100644
index 000000000..40c0883e2
--- /dev/null
+++ b/benchmarks/bloom_filter/defaults.hpp
@@ -0,0 +1,42 @@
+/*
+ * 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/hash_functions.cuh>
+
+#include <nvbench/nvbench.cuh>
+
+#include <cuda/std/array>
+
+#include <vector>
+
+namespace cuco::benchmark::defaults {
+
+using BF_KEY  = nvbench::int64_t;
+using BF_HASH = cuco::xxhash_64<char>;
+using BF_WORD = nvbench::uint32_t;
+
+static constexpr auto BF_N               = 400'000'000;
+static constexpr auto BF_SIZE_MB         = 2'000;
+static constexpr auto BF_WORDS_PER_BLOCK = 8;
+static constexpr auto BF_PATTERN_BITS    = BF_WORDS_PER_BLOCK;
+
+auto const BF_SIZE_MB_RANGE_CACHE =
+  std::vector<nvbench::int64_t>{1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048};
+auto const BF_PATTERN_BITS_RANGE = std::vector<nvbench::int64_t>{1, 2, 4, 6, 8, 16};
+
+}  // namespace cuco::benchmark::defaults
diff --git a/benchmarks/bloom_filter/utils.hpp b/benchmarks/bloom_filter/utils.hpp
new file mode 100644
index 000000000..cec7e06c3
--- /dev/null
+++ b/benchmarks/bloom_filter/utils.hpp
@@ -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.
+ */
+
+#pragma once
+
+#include <cuco/hash_functions.cuh>
+
+#include <nvbench/nvbench.cuh>
+
+#include <thrust/count.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/sequence.h>
+
+#include <cstdint>
+
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::XXHash_64<char>, "xxhash_64", "cuco::xxhash_64");
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::XXHash_32<char>, "xxhash_32", "cuco::xxhash_32");
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_32<char>,
+                             "murmurhash3_32",
+                             "cuco::murmurhash3_32");
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_x86_128<char>,
+                             "murmurhash3_x86_128",
+                             "cuco::murmurhash3_x86_128");
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_x64_128<char>,
+                             "murmurhash3_x64_128",
+                             "cuco::murmurhash3_x64_128");
+NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::identity_hash<char>,
+                             "identity_hash",
+                             "cuco::identity_hash");
+
+namespace cuco::benchmark {
+
+template <typename FilterType>
+void add_fpr_summary(nvbench::state& state, FilterType& filter)
+{
+  filter.clear();
+
+  auto const num_keys = state.get_int64("NumInputs");
+
+  thrust::device_vector<typename FilterType::key_type> keys(num_keys * 2);
+  thrust::sequence(thrust::device, keys.begin(), keys.end(), 1);
+  thrust::device_vector<bool> result(num_keys, false);
+
+  auto tp_begin = keys.begin();
+  auto tp_end   = tp_begin + num_keys;
+  auto tn_begin = tp_end;
+  auto tn_end   = keys.end();
+  filter.add(tp_begin, tp_end);
+  filter.contains(tn_begin, tn_end, result.begin());
+
+  float fp = thrust::count(thrust::device, result.begin(), result.end(), true);
+
+  auto& summ = state.add_summary("FalsePositiveRate");
+  summ.set_string("hint", "FPR");
+  summ.set_string("short_name", "FPR");
+  summ.set_string("description", "False-positive rate of the bloom filter.");
+  summ.set_float64("value", fp / num_keys);
+
+  filter.clear();
+}
+
+}  // namespace cuco::benchmark
\ No newline at end of file
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
index c8975f604..4058600fa 100644
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@@ -45,3 +45,4 @@ ConfigureExample(STATIC_MAP_COUNT_BY_KEY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/st
 ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multimap/host_bulk_example.cu")
 ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/host_bulk_example.cu")
 ConfigureExample(HYPERLOGLOG_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/device_ref_example.cu")
+ConfigureExample(BLOOM_FILTER_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/host_bulk_example.cu")
diff --git a/examples/bloom_filter/host_bulk_example.cu b/examples/bloom_filter/host_bulk_example.cu
new file mode 100644
index 000000000..14a2a5cfa
--- /dev/null
+++ b/examples/bloom_filter/host_bulk_example.cu
@@ -0,0 +1,68 @@
+/*
+ * 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;
+
+  // Spawn a filter with 200 sub-filters.
+  cuco::bloom_filter<int> filter{200};
+
+  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;
+}
\ No newline at end of file
diff --git a/include/cuco/bloom_filter.cuh b/include/cuco/bloom_filter.cuh
new file mode 100644
index 000000000..16642ab57
--- /dev/null
+++ b/include/cuco/bloom_filter.cuh
@@ -0,0 +1,358 @@
+/*
+ * 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/bloom_filter_policy.cuh>
+#include <cuco/bloom_filter_ref.cuh>
+#include <cuco/detail/storage/storage_base.cuh>
+#include <cuco/extent.cuh>
+#include <cuco/hash_functions.cuh>
+#include <cuco/utility/allocator.hpp>
+#include <cuco/utility/cuda_thread_scope.cuh>
+
+#include <cuda/atomic>
+#include <cuda/std/array>
+#include <cuda/std/cstddef>
+#include <cuda/stream_ref>
+
+#include <cstddef>
+#include <cstdint>
+#include <memory>
+
+namespace cuco {
+
+/**
+ * @brief A GPU-accelerated Blocked Bloom Filter.
+ *
+ * The `bloom_filter` supports two types of operations:
+ * - Host-side "bulk" operations
+ * - Device-side "singular" operations
+ *
+ * The host-side bulk operations include `add`, `contains`, etc. These APIs should be used when
+ * there are a large number of keys to add or lookup. For example, given a range of keys
+ * specified by device-accessible iterators, the bulk `add` function will add all keys into
+ * the filter.
+ *
+ * The singular device-side operations allow individual threads (or Cooperative Groups) to perform
+ * independent add or lookup operations from device code. These operations are accessed through
+ * non-owning, trivially copyable reference types (or "ref").
+ *
+ * @tparam Key Key type
+ * @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`)
+ * @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 Allocator = cuco::cuda_allocator<cuda::std::byte>>
+class bloom_filter {
+ public:
+  /**
+   * @brief Non-owning filter ref type
+   *
+   * @tparam NewScope Thread scope of the to be updated ref type
+   */
+  template <cuda::thread_scope NewScope = Scope>
+  using ref_type = bloom_filter_ref<Key, Extent, NewScope, Policy>;
+
+  static constexpr auto thread_scope = ref_type<>::thread_scope;  ///< CUDA thread scope
+  static constexpr auto words_per_block =
+    ref_type<>::words_per_block;  ///< Number of machine words/segments in each filter block
+
+  using key_type    = typename ref_type<>::key_type;     ///< Key Type
+  using extent_type = typename ref_type<>::extent_type;  ///< Extent type
+  using size_type   = typename extent_type::value_type;  ///< Underlying type of the extent type
+  using word_type =
+    typename ref_type<>::word_type;  ///< Underlying word/segment type of a filter block
+  using allocator_type =
+    typename std::allocator_traits<Allocator>::template rebind_alloc<word_type>;  ///< Allocator
+                                                                                  ///< type
+
+  bloom_filter(bloom_filter const&) = delete;  ///< Copy constructor is not available
+  bloom_filter& operator=(bloom_filter const&) =
+    delete;  ///< Copy-assignment constructor is not available
+
+  bloom_filter(bloom_filter&&) = default;  ///< Move constructor
+
+  /**
+   * @brief Move-assignment operator.
+   *
+   * @return Reference of the current `bloom_filter` object
+   */
+  bloom_filter& operator=(bloom_filter&&) = default;
+
+  ~bloom_filter() = default;  ///< Destructor
+
+  /**
+   * @brief Constructs a statically-sized Bloom filter.
+   *
+   * @note The total number of bits in the filter is determined by `words_per_block * num_blocks *
+   * sizeof(word_type) * CHAR_BIT`.
+   *
+   * @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 alloc Allocator used for allocating device-accessible storage
+   * @param stream CUDA stream used to initialize the filter
+   */
+  __host__ explicit constexpr bloom_filter(Extent num_blocks,
+                                           cuda_thread_scope<Scope> scope = {},
+                                           Policy const& policy           = {},
+                                           Allocator const& alloc         = {},
+                                           cuda::stream_ref stream        = {});
+
+  /**
+   * @brief Erases all information from the filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `clear_async`.
+   *
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  __host__ constexpr void clear(cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously erases all information from the filter.
+   *
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  __host__ constexpr void clear_async(cuda::stream_ref stream = {});
+
+  /**
+   * @brief Adds all keys in the range `[first, last)` to the filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `add_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt>
+  __host__ constexpr void add(InputIt first, InputIt last, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously adds all keys in the range `[first, last)` to the filter.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt>
+  __host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Adds keys in the range `[first, last)` if `pred` of the corresponding `stencil` returns
+   * `true`.
+   *
+   * @note The key `*(first + i)` is added if `pred( *(stencil + i) )` returns `true`.
+   * @note This function synchronizes the given stream and returns the number of successful
+   * insertions. For asynchronous execution use `add_if_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if(
+    InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously adds keys in the range `[first, last)` if `pred` of the corresponding
+   * `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is added if `pred( *(stencil + i) )` returns `true`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if_async(InputIt first,
+                                       InputIt last,
+                                       StencilIt stencil,
+                                       Predicate pred,
+                                       cuda::stream_ref stream = {}) noexcept;
+
+  /**
+   * @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
+   * filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `contains_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains(InputIt first,
+                                   InputIt last,
+                                   OutputIt output_begin,
+                                   cuda::stream_ref stream = {}) const;
+
+  /**
+   * @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
+   * present in the filter.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains_async(InputIt first,
+                                         InputIt last,
+                                         OutputIt output_begin,
+                                         cuda::stream_ref stream = {}) const noexcept;
+
+  /**
+   * @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
+   * filter if `pred` of the corresponding `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is queried if `pred( *(stencil + i) )` returns `true`.
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `contains_if_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if(InputIt first,
+                                      InputIt last,
+                                      StencilIt stencil,
+                                      Predicate pred,
+                                      OutputIt output_begin,
+                                      cuda::stream_ref stream = {}) const;
+
+  /**
+   * @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
+   * present in the filter if `pred` of the corresponding `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is queried if `pred( *(stencil + i) )` returns `true`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if_async(InputIt first,
+                                            InputIt last,
+                                            StencilIt stencil,
+                                            Predicate pred,
+                                            OutputIt output_begin,
+                                            cuda::stream_ref stream = {}) const noexcept;
+
+  /**
+   * @brief Gets a pointer to the underlying filter storage.
+   *
+   * @return Pointer to the underlying filter storage
+   */
+  [[nodiscard]] __host__ constexpr word_type* data() noexcept;
+
+  /**
+   * @brief Gets a pointer to the underlying filter storage.
+   *
+   * @return Pointer to the underlying filter storage
+   */
+  [[nodiscard]] __host__ constexpr word_type const* data() const noexcept;
+
+  /**
+   * @brief Gets the number of sub-filter blocks.
+   *
+   * @return Number of sub-filter blocks
+   */
+  [[nodiscard]] __host__ constexpr extent_type block_extent() const noexcept;
+
+  /**
+   * @brief Gets the allocator.
+   *
+   * @return The allocator
+   */
+  [[nodiscard]] __host__ constexpr allocator_type allocator() const noexcept;
+
+  /**
+   * @brief Get device ref.
+   *
+   * @return Device ref of the current `bloom_filter` object
+   */
+  [[nodiscard]] __host__ constexpr ref_type<> ref() const noexcept;
+
+ private:
+  allocator_type allocator_;  ///< Allocator used to allocate device-accessible storage
+  std::unique_ptr<word_type, detail::custom_deleter<std::size_t, allocator_type>>
+    data_;          ///< Storage of the current `bloom_filter` object
+  ref_type<> ref_;  ///< Device ref of the current `bloom_filter` object
+};
+}  // namespace cuco
+
+#include <cuco/detail/bloom_filter/bloom_filter.inl>
\ No newline at end of file
diff --git a/include/cuco/bloom_filter_policy.cuh b/include/cuco/bloom_filter_policy.cuh
new file mode 100644
index 000000000..5b8f87d95
--- /dev/null
+++ b/include/cuco/bloom_filter_policy.cuh
@@ -0,0 +1,119 @@
+/*
+ * 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/bloom_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.
+ *
+ * @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 Hash, class Word, std::uint32_t WordsPerBlock>
+class bloom_filter_policy {
+  using impl_type = cuco::detail::bloom_filter_policy_impl<Hash, Word, WordsPerBlock>;
+
+ 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
+
+ public:
+  /**
+   * @brief Constructs the `bloom_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
+   * 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 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                  = {});
+
+  /**
+   * @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/bloom_filter_policy.inl>
\ No newline at end of file
diff --git a/include/cuco/bloom_filter_ref.cuh b/include/cuco/bloom_filter_ref.cuh
new file mode 100644
index 000000000..e4434845f
--- /dev/null
+++ b/include/cuco/bloom_filter_ref.cuh
@@ -0,0 +1,366 @@
+/*
+ * 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/bloom_filter_impl.cuh>
+#include <cuco/utility/cuda_thread_scope.cuh>
+
+#include <cuda/atomic>
+#include <cuda/stream_ref>
+
+namespace cuco {
+
+/**
+ * @brief Non-owning "ref" type of `bloom_filter`.
+ *
+ * @note Ref types are trivially-copyable and are intended to be passed by value.
+ *
+ * @tparam Key Key type
+ * @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`)
+ */
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+class bloom_filter_ref {
+  using impl_type = detail::bloom_filter_impl<Key, Extent, Scope, Policy>;  ///< Implementation type
+
+ public:
+  static constexpr auto thread_scope = impl_type::thread_scope;  ///< CUDA thread scope
+  static constexpr auto words_per_block =
+    impl_type::words_per_block;  ///< Number of words/segments in each filter block
+
+  using key_type    = typename impl_type::key_type;      ///< Key Type
+  using extent_type = typename impl_type::extent_type;   ///< Extent type
+  using size_type   = typename extent_type::value_type;  ///< Underlying type of the extent type
+  using word_type =
+    typename impl_type::word_type;  ///< Underlying word/segment type of a filter block
+
+  /**
+   * @brief Constructs the ref object from existing storage.
+   *
+   * @note The storage span starting at `data` must have an extent of at least `num_blocks`
+   * elements.
+   * @note `data` must be aligned to at least `sizeof(word_type) * words_per_block`.
+   *
+   * @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`)
+   */
+  __host__ __device__ explicit constexpr bloom_filter_ref(word_type* data,
+                                                          Extent num_blocks,
+                                                          cuda_thread_scope<Scope> scope,
+                                                          Policy const& policy);
+
+  /**
+   * @brief Device function that cooperatively erases all information from the filter.
+   *
+   * @tparam CG Cooperative Group type
+   *
+   * @param group The Cooperative Group this operation is executed with
+   */
+  template <class CG>
+  __device__ constexpr void clear(CG const& group);
+
+  /**
+   * @brief Erases all information from the filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `clear_async`.
+   *
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  __host__ constexpr void clear(cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously erases all information from the filter.
+   *
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  __host__ constexpr void clear_async(cuda::stream_ref stream = {});
+
+  /**
+   * @brief Device function that adds a key to the filter.
+   *
+   * @tparam ProbeKey Input type that is implicitly convertible to `key_type`
+   *
+   * @param key The key to be added
+   */
+  template <class ProbeKey>
+  __device__ void add(ProbeKey const& key);
+
+  /**
+   * @brief Device function that cooperatively adds a key to the filter.
+   *
+   * @note Best performance is achieved if the size of the CG is equal to `words_per_block`.
+   *
+   * @tparam CG Cooperative Group type
+   * @tparam ProbeKey Input key type
+   *
+   * @param group The Cooperative Group this operation is executed with
+   * @param key The key to be added
+   */
+  template <class CG, class ProbeKey>
+  __device__ void add(CG const& group, ProbeKey const& key);
+
+  /**
+   * @brief Adds all keys in the range `[first, last)` to the filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `add_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt>
+  __host__ constexpr void add(InputIt first, InputIt last, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously adds all keys in the range `[first, last)` to the filter.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt>
+  __host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Adds keys in the range `[first, last)` if `pred` of the corresponding `stencil` returns
+   * `true`.
+   *
+   * @note The key `*(first + i)` is added if `pred( *(stencil + i) )` returns `true`.
+   * @note This function synchronizes the given stream and returns the number of successful
+   * insertions. For asynchronous execution use `add_if_async`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if(
+    InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream = {});
+
+  /**
+   * @brief Asynchronously adds keys in the range `[first, last)` if `pred` of the corresponding
+   * `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is added if `pred( *(stencil + i) )` returns `true`.
+   *
+   * @tparam InputIt Device-accessible random access input key iterator
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if_async(InputIt first,
+                                       InputIt last,
+                                       StencilIt stencil,
+                                       Predicate pred,
+                                       cuda::stream_ref stream = {}) noexcept;
+
+  /**
+   * @brief Device function that tests if a key's fingerprint is present in the filter.
+   *
+   * @tparam ProbeKey Probe key type
+   *
+   * @param key The key to be tested
+   *
+   * @return `true` iff the key's fingerprint was present in the filter
+   */
+  template <class ProbeKey>
+  [[nodiscard]] __device__ bool contains(ProbeKey const& key) const;
+
+  /**
+   * @brief Device function that tests if a key's fingerprint is present in the filter.
+   *
+   * @note Best performance is achieved if the size of the CG is equal to `(words_per_block *
+   * sizeof(word_type)) / 32`.
+   *
+   * @tparam CG Cooperative Group type
+   * @tparam ProbeKey Input type that is implicitly convertible to `key_type`
+   *
+   * @param group The Cooperative Group this operation is executed with
+   * @param key The key to be tested
+   *
+   * @return `true` iff the key's fingerprint was present in the filter
+   */
+  template <class CG, class ProbeKey>
+  [[nodiscard]] __device__ bool contains(CG const& group, ProbeKey const& key) const;
+
+  // TODO
+  // template <class CG, class InputIt, class OutputIt>
+  // __device__ void contains(CG const& group, InputIt first, InputIt last, OutputIt output_begin)
+  // const;
+
+  /**
+   * @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
+   * filter.
+   *
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `contains_async`.
+   *
+   * @tparam InputIt Device-accessible random access input iterator where
+   * <tt>std::is_convertible<std::iterator_traits<InputIt>::value_type,
+   * bloom_filter<K>::key_type></tt> is `true`
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains(InputIt first,
+                                   InputIt last,
+                                   OutputIt output_begin,
+                                   cuda::stream_ref stream = {}) const;
+
+  /**
+   * @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
+   * present in the filter.
+   *
+   * @tparam InputIt Device-accessible random access input iterator where
+   * <tt>std::is_convertible<std::iterator_traits<InputIt>::value_type,
+   * bloom_filter<K>::key_type></tt> is `true`
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains_async(InputIt first,
+                                         InputIt last,
+                                         OutputIt output_begin,
+                                         cuda::stream_ref stream = {}) const noexcept;
+
+  /**
+   * @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
+   * filter if `pred` of the corresponding `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is queried if `pred( *(stencil + i) )` returns `true`.
+   * @note This function synchronizes the given stream. For asynchronous execution use
+   * `contains_if_async`.
+   *
+   * @tparam InputIt Device-accessible random access input iterator where
+   * <tt>std::is_convertible<std::iterator_traits<InputIt>::value_type,
+   * bloom_filter<K>::key_type></tt> is `true`
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if(InputIt first,
+                                      InputIt last,
+                                      StencilIt stencil,
+                                      Predicate pred,
+                                      OutputIt output_begin,
+                                      cuda::stream_ref stream = {}) const;
+
+  /**
+   * @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
+   * present in the filter if `pred` of the corresponding `stencil` returns `true`.
+   *
+   * @note The key `*(first + i)` is queried if `pred( *(stencil + i) )` returns `true`.
+   *
+   * @tparam InputIt Device-accessible random access input iterator where
+   * <tt>std::is_convertible<std::iterator_traits<InputIt>::value_type,
+   * bloom_filter<K>::key_type></tt> is `true`
+   * @tparam StencilIt Device-accessible random-access iterator whose `value_type` is
+   * convertible to Predicate's argument type
+   * @tparam Predicate Unary predicate callable whose return type must be convertible to `bool` and
+   * argument type is convertible from <tt>std::iterator_traits<StencilIt>::value_type</tt>
+   * @tparam OutputIt Device-accessible output iterator assignable from `bool`
+   *
+   * @param first Beginning of the sequence of keys
+   * @param last End of the sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil +
+   * std::distance(first, last))`
+   * @param output_begin Beginning of the sequence of booleans for the presence of each key
+   * @param stream CUDA stream used for device memory operations and kernel launches
+   */
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if_async(InputIt first,
+                                            InputIt last,
+                                            StencilIt stencil,
+                                            Predicate pred,
+                                            OutputIt output_begin,
+                                            cuda::stream_ref stream = {}) const noexcept;
+
+  /**
+   * @brief Gets a pointer to the underlying filter storage.
+   *
+   * @return Pointer to the underlying filter storage
+   */
+  [[nodiscard]] __host__ __device__ constexpr word_type* data() noexcept;
+
+  /**
+   * @brief Gets a pointer to the underlying filter storage.
+   *
+   * @return Pointer to the underlying filter storage
+   */
+  [[nodiscard]] __host__ __device__ constexpr word_type const* data() const noexcept;
+
+  /**
+   * @brief Gets the number of sub-filter blocks.
+   *
+   * @return Number of sub-filter blocks
+   */
+  [[nodiscard]] __host__ __device__ constexpr extent_type block_extent() const noexcept;
+
+ private:
+  impl_type impl_;  ///< Object containing the Blocked Bloom Filter implementation
+};
+}  // namespace cuco
+
+#include <cuco/detail/bloom_filter/bloom_filter_ref.inl>
\ No newline at end of file
diff --git a/include/cuco/detail/bloom_filter/bloom_filter.inl b/include/cuco/detail/bloom_filter/bloom_filter.inl
new file mode 100644
index 000000000..3624d320a
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/bloom_filter.inl
@@ -0,0 +1,173 @@
+/*
+ * 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/storage/storage_base.cuh>
+#include <cuco/utility/cuda_thread_scope.cuh>
+
+#include <cuda/atomic>
+#include <cuda/stream_ref>
+
+#include <cstdint>
+
+namespace cuco {
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+__host__ constexpr bloom_filter<Key, Extent, Scope, Policy, Allocator>::bloom_filter(
+  Extent num_blocks,
+  cuda_thread_scope<Scope>,
+  Policy const& policy,
+  Allocator const& alloc,
+  cuda::stream_ref stream)
+  : allocator_{alloc},
+    data_{allocator_.allocate(num_blocks * words_per_block),
+          detail::custom_deleter<std::size_t, allocator_type>{num_blocks * words_per_block,
+                                                              allocator_}},
+    ref_{data_.get(), num_blocks, {}, policy}
+{
+  this->clear_async(stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::clear(
+  cuda::stream_ref stream)
+{
+  ref_.clear(stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::clear_async(
+  cuda::stream_ref stream)
+{
+  ref_.clear_async(stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::add(
+  InputIt first, InputIt last, cuda::stream_ref stream)
+{
+  ref_.add(first, last, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::add_async(
+  InputIt first, InputIt last, cuda::stream_ref stream)
+{
+  ref_.add_async(first, last, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class StencilIt, class Predicate>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::add_if(
+  InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream)
+{
+  ref_.add_if(first, last, stencil, pred, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class StencilIt, class Predicate>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::add_if_async(
+  InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream) noexcept
+{
+  ref_.add_if_async(first, last, stencil, pred, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class OutputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::contains(
+  InputIt first, InputIt last, OutputIt output_begin, cuda::stream_ref stream) const
+{
+  ref_.contains(first, last, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class OutputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::contains_async(
+  InputIt first, InputIt last, OutputIt output_begin, cuda::stream_ref stream) const noexcept
+{
+  ref_.contains_async(first, last, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::contains_if(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate pred,
+  OutputIt output_begin,
+  cuda::stream_ref stream) const
+{
+  ref_.contains_if(first, last, stencil, pred, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+__host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::contains_if_async(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate pred,
+  OutputIt output_begin,
+  cuda::stream_ref stream) const noexcept
+{
+  ref_.contains_if_async(first, last, stencil, pred, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ constexpr
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::word_type*
+  bloom_filter<Key, Extent, Scope, Policy, Allocator>::data() noexcept
+{
+  return ref_.data();
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ constexpr
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::word_type const*
+  bloom_filter<Key, Extent, Scope, Policy, Allocator>::data() const noexcept
+{
+  return ref_.data();
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ constexpr
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::extent_type
+  bloom_filter<Key, Extent, Scope, Policy, Allocator>::block_extent() const noexcept
+{
+  return ref_.block_extent();
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ constexpr
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::allocator_type
+  bloom_filter<Key, Extent, Scope, Policy, Allocator>::allocator() const noexcept
+{
+  return allocator_;
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ constexpr
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::ref_type<>
+  bloom_filter<Key, Extent, Scope, Policy, Allocator>::ref() const noexcept
+{
+  return ref_;
+}
+
+}  // namespace cuco
\ No newline at end of file
diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh
new file mode 100644
index 000000000..a67db3475
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh
@@ -0,0 +1,368 @@
+/*
+ * 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/kernels.cuh>
+#include <cuco/detail/error.hpp>
+#include <cuco/detail/utility/cuda.cuh>
+#include <cuco/detail/utility/cuda.hpp>
+#include <cuco/detail/utils.hpp>
+#include <cuco/utility/cuda_thread_scope.cuh>
+
+#include <cub/device/device_for.cuh>
+#include <cuda/atomic>
+#include <cuda/std/__algorithm/max.h>
+#include <cuda/std/__algorithm/min.h>  // TODO #include <cuda/std/algorithm> once available
+#include <cuda/std/array>
+#include <cuda/std/bit>
+#include <cuda/std/tuple>
+#include <cuda/std/type_traits>
+#include <cuda/stream_ref>
+#include <thrust/functional.h>
+#include <thrust/iterator/constant_iterator.h>
+
+#include <cstdint>
+#include <nv/target>
+
+namespace cuco::detail {
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+class bloom_filter_impl {
+ public:
+  using key_type    = Key;
+  using extent_type = Extent;
+  using size_type   = typename extent_type::value_type;
+  using policy_type = Policy;
+  using word_type   = typename policy_type::word_type;
+
+  static constexpr auto thread_scope    = Scope;
+  static constexpr auto words_per_block = policy_type::words_per_block;
+
+  static_assert(cuda::std::has_single_bit(words_per_block) and words_per_block <= 32,
+                "Number of words per block must be a power-of-two and less than or equal to 32");
+
+  static_assert(
+    cuda::std::is_constructible_v<cuda::atomic_ref<word_type, Scope>, word_type&> &&
+      cuda::std::is_invocable_r_v<word_type,
+                                  decltype(&cuda::atomic_ref<word_type, Scope>::fetch_or),
+                                  cuda::atomic_ref<word_type, Scope>*,
+                                  word_type,
+                                  cuda::std::memory_order>,
+    "Invalid word type");
+
+  __host__ __device__ explicit constexpr bloom_filter_impl(word_type* filter,
+                                                           Extent num_blocks,
+                                                           cuda_thread_scope<Scope>,
+                                                           Policy policy)
+    : words_{filter}, num_blocks_{num_blocks}, policy_{policy}
+  {
+    auto const alignment =
+      1ull << cuda::std::countr_zero(reinterpret_cast<cuda::std::uintptr_t>(filter));
+
+    NV_DISPATCH_TARGET(
+      NV_IS_HOST,
+      (CUCO_EXPECTS(alignment >= required_alignment(), "Invalid memory alignment");
+       CUCO_EXPECTS(num_blocks_ > 0, "Number of blocks cannot be zero");),
+      NV_IS_DEVICE,
+      (if (alignment < required_alignment() or num_blocks_ == 0) {
+        __trap();  // TODO this kills the kernel and corrupts the CUDA context. Not ideal.
+      }))
+  }
+
+  template <class CG>
+  __device__ constexpr void clear(CG const& group)
+  {
+    for (int i = group.thread_rank(); num_blocks_ * words_per_block; i += group.size()) {
+      words_[i] = 0;
+    }
+  }
+
+  __host__ constexpr void clear(cuda::stream_ref stream)
+  {
+    this->clear_async(stream);
+    stream.wait();
+  }
+
+  __host__ constexpr void clear_async(cuda::stream_ref stream)
+  {
+    CUCO_CUDA_TRY(cub::DeviceFor::ForEachN(
+      words_,
+      num_blocks_ * words_per_block,
+      [] __device__(word_type & word) { word = 0; },
+      stream.get()));
+  }
+
+  template <class ProbeKey>
+  __device__ void add(ProbeKey const& key)
+  {
+    auto const hash_value = policy_.hash(key);
+    auto const idx        = policy_.block_index(hash_value, num_blocks_);
+
+#pragma unroll words_per_block
+    for (uint32_t i = 0; i < words_per_block; ++i) {
+      auto const word = policy_.word_pattern(hash_value, i);
+      if (word != 0) {
+        auto atom_word =
+          cuda::atomic_ref<word_type, thread_scope>{*(words_ + (idx * words_per_block + i))};
+        atom_word.fetch_or(word, cuda::memory_order_relaxed);
+      }
+    }
+  }
+
+  template <class CG, class ProbeKey>
+  __device__ void add(CG const& group, ProbeKey const& key)
+  {
+    constexpr auto num_threads         = tile_size_v<CG>;
+    constexpr auto optimal_num_threads = add_optimal_cg_size();
+    constexpr auto words_per_thread    = words_per_block / optimal_num_threads;
+
+    // If single thread is optimal, use scalar add
+    if constexpr (num_threads == 1 or optimal_num_threads == 1) {
+      this->add(key);
+    } else {
+      auto const rank = group.thread_rank();
+
+      auto const hash_value = policy_.hash(key);
+      auto const idx        = policy_.block_index(hash_value, num_blocks_);
+
+#pragma unroll
+      for (uint32_t i = rank; i < optimal_num_threads; i += num_threads) {
+        auto const word = policy_.word_pattern(hash_value, rank);
+
+        auto atom_word =
+          cuda::atomic_ref<word_type, thread_scope>{*(words_ + (idx * words_per_block + rank))};
+        atom_word.fetch_or(word, cuda::memory_order_relaxed);
+      }
+    }
+  }
+
+  template <class InputIt>
+  __host__ constexpr void add(InputIt first, InputIt last, cuda::stream_ref stream)
+  {
+    this->add_async(first, last, stream);
+    stream.wait();
+  }
+
+  template <class InputIt>
+  __host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream)
+  {
+    auto const num_keys = cuco::detail::distance(first, last);
+    if (num_keys == 0) { return; }
+
+    if constexpr (words_per_block == 1) {
+      CUCO_CUDA_TRY(cub::DeviceFor::ForEachCopyN(
+        first,
+        num_keys,
+        [*this] __device__(key_type const key) mutable { this->add(key); },
+        stream.get()));
+    } else {
+      auto const always_true = thrust::constant_iterator<bool>{true};
+      this->add_if_async(first, last, always_true, thrust::identity{}, stream);
+    }
+  }
+
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if(
+    InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream)
+  {
+    this->add_if_async(first, last, stencil, pred, stream);
+    stream.wait();
+  }
+
+  template <class InputIt, class StencilIt, class Predicate>
+  __host__ constexpr void add_if_async(InputIt first,
+                                       InputIt last,
+                                       StencilIt stencil,
+                                       Predicate pred,
+                                       cuda::stream_ref stream) noexcept
+  {
+    auto const num_keys = cuco::detail::distance(first, last);
+    if (num_keys == 0) { return; }
+
+    auto constexpr cg_size    = add_optimal_cg_size();
+    auto constexpr block_size = cuco::detail::default_block_size();
+    auto const grid_size =
+      cuco::detail::grid_size(num_keys, cg_size, cuco::detail::default_stride(), block_size);
+
+    detail::add_if_n<cg_size, block_size>
+      <<<grid_size, block_size, 0, stream.get()>>>(first, num_keys, stencil, pred, *this);
+  }
+
+  template <class ProbeKey>
+  [[nodiscard]] __device__ bool contains(ProbeKey const& key) const
+  {
+    auto const hash_value = policy_.hash(key);
+
+    auto const stored_pattern = this->vec_load_words<words_per_block>(
+      policy_.block_index(hash_value, num_blocks_) * words_per_block);
+
+#pragma unroll words_per_block
+    for (uint32_t i = 0; i < words_per_block; ++i) {
+      auto const expected_pattern = policy_.word_pattern(hash_value, i);
+      if ((stored_pattern[i] & expected_pattern) != expected_pattern) { return false; }
+    }
+
+    return true;
+  }
+
+  template <class CG, class ProbeKey>
+  [[nodiscard]] __device__ bool contains(CG const& group, ProbeKey const& key) const
+  {
+    constexpr auto num_threads         = tile_size_v<CG>;
+    constexpr auto optimal_num_threads = contains_optimal_cg_size();
+    constexpr auto words_per_thread    = words_per_block / optimal_num_threads;
+
+    // If single thread is optimal, use scalar contains
+    if constexpr (num_threads == 1 or optimal_num_threads == 1) {
+      return this->contains(key);
+    } else {
+      auto const rank       = group.thread_rank();
+      auto const hash_value = policy_.hash(key);
+      bool success          = true;
+
+#pragma unroll
+      for (uint32_t i = rank; i < optimal_num_threads; i += num_threads) {
+        auto const thread_offset  = i * words_per_thread;
+        auto const stored_pattern = this->vec_load_words<words_per_thread>(
+          policy_.block_index(hash_value, num_blocks_) * words_per_block + thread_offset);
+#pragma unroll words_per_thread
+        for (uint32_t j = 0; j < words_per_thread; ++j) {
+          auto const expected_pattern = policy_.word_pattern(hash_value, thread_offset + j);
+          if ((stored_pattern[j] & expected_pattern) != expected_pattern) { success = false; }
+        }
+      }
+
+      return group.all(success);
+    }
+  }
+
+  // TODO
+  // template <class CG, class InputIt, class OutputIt>
+  // __device__ void contains(CG const& group, InputIt first, InputIt last, OutputIt output_begin)
+  // const;
+
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains(InputIt first,
+                                   InputIt last,
+                                   OutputIt output_begin,
+                                   cuda::stream_ref stream) const
+  {
+    this->contains_async(first, last, output_begin, stream);
+    stream.wait();
+  }
+
+  template <class InputIt, class OutputIt>
+  __host__ constexpr void contains_async(InputIt first,
+                                         InputIt last,
+                                         OutputIt output_begin,
+                                         cuda::stream_ref stream) const noexcept
+  {
+    auto const always_true = thrust::constant_iterator<bool>{true};
+    this->contains_if_async(first, last, always_true, thrust::identity{}, output_begin, stream);
+  }
+
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if(InputIt first,
+                                      InputIt last,
+                                      StencilIt stencil,
+                                      Predicate pred,
+                                      OutputIt output_begin,
+                                      cuda::stream_ref stream) const
+  {
+    this->contains_if_async(first, last, stencil, pred, output_begin, stream);
+    stream.wait();
+  }
+
+  template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+  __host__ constexpr void contains_if_async(InputIt first,
+                                            InputIt last,
+                                            StencilIt stencil,
+                                            Predicate pred,
+                                            OutputIt output_begin,
+                                            cuda::stream_ref stream) const noexcept
+  {
+    auto const num_keys = cuco::detail::distance(first, last);
+    if (num_keys == 0) { return; }
+
+    auto constexpr cg_size    = contains_optimal_cg_size();
+    auto constexpr block_size = cuco::detail::default_block_size();
+    auto const grid_size =
+      cuco::detail::grid_size(num_keys, cg_size, cuco::detail::default_stride(), block_size);
+
+    detail::contains_if_n<cg_size, block_size><<<grid_size, block_size, 0, stream.get()>>>(
+      first, num_keys, stencil, pred, output_begin, *this);
+  }
+
+  [[nodiscard]] __host__ __device__ constexpr word_type* data() noexcept { return words_; }
+
+  [[nodiscard]] __host__ __device__ constexpr word_type const* data() const noexcept
+  {
+    return words_;
+  }
+
+  [[nodiscard]] __host__ __device__ constexpr extent_type block_extent() const noexcept
+  {
+    return num_blocks_;
+  }
+
+  // TODO
+  // [[nodiscard]] __host__ double occupancy() const;
+  // [[nodiscard]] __host__ double expected_false_positive_rate(size_t unique_keys) const
+  // [[nodiscard]] __host__ __device__ static uint32_t optimal_pattern_bits(size_t num_blocks)
+  // template <typename CG, cuda::thread_scope NewScope = thread_scope>
+  // [[nodiscard]] __device__ constexpr auto make_copy(CG const& group, word_type* const
+  // memory_to_use, cuda_thread_scope<NewScope> scope = {}) const noexcept;
+
+ private:
+  template <uint32_t NumWords>
+  __device__ constexpr cuda::std::array<word_type, NumWords> vec_load_words(size_type index) const
+  {
+    return *reinterpret_cast<cuda::std::array<word_type, NumWords>*>(__builtin_assume_aligned(
+      words_ + index, cuda::std::min(sizeof(word_type) * NumWords, required_alignment())));
+  }
+
+  __host__ __device__ static constexpr size_t max_vec_bytes() noexcept
+  {
+    constexpr auto word_bytes  = sizeof(word_type);
+    constexpr auto block_bytes = word_bytes * words_per_block;
+    return cuda::std::min(cuda::std::max(word_bytes, 32ul),
+                          block_bytes);  // aiming for 2xLDG128 -> 1 sector per thread
+  }
+
+  [[nodiscard]] __host__ __device__ static constexpr int32_t add_optimal_cg_size()
+  {
+    return words_per_block;  // one thread per word so atomic updates can be coalesced
+  }
+
+  [[nodiscard]] __host__ __device__ static constexpr int32_t contains_optimal_cg_size()
+  {
+    constexpr auto word_bytes  = sizeof(word_type);
+    constexpr auto block_bytes = word_bytes * words_per_block;
+    return block_bytes / max_vec_bytes();  // one vector load per thread
+  }
+
+  __host__ __device__ static constexpr size_t required_alignment() noexcept
+  {
+    return cuda::std::min(sizeof(word_type) * words_per_block, max_vec_bytes());
+  }
+
+  word_type* words_;
+  extent_type num_blocks_;
+  policy_type policy_;
+};
+
+}  // namespace cuco::detail
\ No newline at end of file
diff --git a/include/cuco/detail/bloom_filter/bloom_filter_policy.inl b/include/cuco/detail/bloom_filter/bloom_filter_policy.inl
new file mode 100644
index 000000000..f117798f9
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/bloom_filter_policy.inl
@@ -0,0 +1,56 @@
+/*
+ * 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 Hash, class Word, uint32_t WordsPerBlock>
+__host__ __device__ constexpr bloom_filter_policy<Hash, Word, WordsPerBlock>::bloom_filter_policy(
+  uint32_t pattern_bits, Hash hash)
+  : impl_{pattern_bits, hash}
+{
+}
+
+template <class Hash, class Word, uint32_t WordsPerBlock>
+__device__ constexpr typename bloom_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type
+bloom_filter_policy<Hash, Word, WordsPerBlock>::hash(
+  typename bloom_filter_policy<Hash, Word, WordsPerBlock>::hash_argument_type const& key) const
+{
+  return impl_.hash(key);
+}
+
+template <class Hash, class Word, uint32_t WordsPerBlock>
+template <class Extent>
+__device__ constexpr auto bloom_filter_policy<Hash, Word, WordsPerBlock>::block_index(
+  typename bloom_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type hash,
+  Extent num_blocks) const
+{
+  return impl_.block_index(hash, num_blocks);
+}
+
+template <class Hash, class Word, uint32_t WordsPerBlock>
+__device__ constexpr typename bloom_filter_policy<Hash, Word, WordsPerBlock>::word_type
+bloom_filter_policy<Hash, Word, WordsPerBlock>::word_pattern(
+  bloom_filter_policy<Hash, Word, WordsPerBlock>::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/bloom_filter_policy_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh
new file mode 100644
index 000000000..d7af7667c
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/bloom_filter_policy_impl.cuh
@@ -0,0 +1,118 @@
+/*
+ * 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/error.hpp>
+
+#include <cuda/std/bit>
+#include <cuda/std/limits>
+#include <cuda/std/tuple>
+#include <cuda/std/type_traits>
+
+#include <cstdint>
+#include <nv/target>
+
+namespace cuco::detail {
+
+template <class Hash, class Word, uint32_t WordsPerBlock>
+class bloom_filter_policy_impl {
+ public:
+  using hasher             = Hash;
+  using word_type          = Word;
+  using hash_argument_type = typename hasher::argument_type;
+  using hash_result_type   = decltype(std::declval<hasher>()(std::declval<hash_argument_type>()));
+
+  static constexpr std::uint32_t words_per_block = WordsPerBlock;
+
+ private:
+  static constexpr std::uint32_t word_bits       = cuda::std::numeric_limits<word_type>::digits;
+  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)
+    : pattern_bits_{pattern_bits},
+      min_bits_per_word_{pattern_bits_ / words_per_block},
+      remainder_bits_{pattern_bits_ % words_per_block},
+      hash_{hash}
+  {
+    // This ensures each word in the block has at least one bit set; otherwise we would never use
+    // some of the words
+    constexpr uint32_t min_pattern_bits = words_per_block;
+
+    // The maximum number of bits to be set for a key is capped by the total number of bits in the
+    // filter block
+    constexpr uint32_t max_pattern_bits = word_bits * words_per_block;
+
+    constexpr uint32_t hash_bits = cuda::std::numeric_limits<hash_result_type>::digits;
+    constexpr uint32_t max_pattern_bits_from_hash = hash_bits / bit_index_width;
+
+    NV_DISPATCH_TARGET(
+      NV_IS_HOST,
+      (CUCO_EXPECTS(
+         pattern_bits <= max_pattern_bits_from_hash,
+         "`hash_result_type` too narrow to generate the requested number of `pattern_bits`");
+       CUCO_EXPECTS(pattern_bits_ >= min_pattern_bits,
+                    "`pattern_bits` must be at least `words_per_block`");
+       CUCO_EXPECTS(
+         pattern_bits_ <= max_pattern_bits,
+         "`pattern_bits` must be less than the total number of bits in a filter block");),
+      NV_IS_DEVICE,
+      (if (pattern_bits_ > max_pattern_bits_from_hash or pattern_bits_ < min_pattern_bits or
+           pattern_bits_ > max_pattern_bits) {
+        __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 <class Extent>
+  __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const
+  {
+    return hash % num_blocks;
+  }
+
+  __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const
+  {
+    word_type constexpr bit_index_mask = (word_type{1} << bit_index_width) - 1;
+
+    auto const bits_so_far = min_bits_per_word_ * word_index +
+                             (word_index < remainder_bits_ ? word_index : remainder_bits_);
+
+    hash >>= bits_so_far * bit_index_width;
+
+    word_type word        = 0;
+    int32_t bits_per_word = min_bits_per_word_ + (word_index < remainder_bits_ ? 1 : 0);
+
+    for (int32_t bit = 0; bit < bits_per_word; ++bit) {
+      word |= word_type{1} << (hash & bit_index_mask);
+      hash >>= bit_index_width;
+    }
+
+    return word;
+  }
+
+ private:
+  uint32_t pattern_bits_;
+  uint32_t min_bits_per_word_;
+  uint32_t remainder_bits_;
+  hasher hash_;
+};
+
+}  // namespace cuco::detail
\ No newline at end of file
diff --git a/include/cuco/detail/bloom_filter/bloom_filter_ref.inl b/include/cuco/detail/bloom_filter/bloom_filter_ref.inl
new file mode 100644
index 000000000..54089b2f9
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/bloom_filter_ref.inl
@@ -0,0 +1,183 @@
+/*
+ * 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/utility/cuda_thread_scope.cuh>
+
+#include <cuda/atomic>
+#include <cuda/stream_ref>
+
+namespace cuco {
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+__host__ __device__ constexpr bloom_filter_ref<Key, Extent, Scope, Policy>::bloom_filter_ref(
+  word_type* data, Extent num_blocks, cuda_thread_scope<Scope>, Policy const& policy)
+  : impl_{data, num_blocks, {}, policy}
+{
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class CG>
+__device__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::clear(CG const& group)
+{
+  impl_.clear(group);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::clear(cuda::stream_ref stream)
+{
+  impl_.clear(stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::clear_async(
+  cuda::stream_ref stream)
+{
+  impl_.clear_async(stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class ProbeKey>
+__device__ void bloom_filter_ref<Key, Extent, Scope, Policy>::add(ProbeKey const& key)
+{
+  impl_.add(key);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class CG, class ProbeKey>
+__device__ void bloom_filter_ref<Key, Extent, Scope, Policy>::add(CG const& group,
+                                                                  ProbeKey const& key)
+{
+  impl_.add(group, key);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::add(InputIt first,
+                                                                          InputIt last,
+                                                                          cuda::stream_ref stream)
+{
+  impl_.add(first, last, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::add_async(
+  InputIt first, InputIt last, cuda::stream_ref stream)
+{
+  impl_.add_async(first, last, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class StencilIt, class Predicate>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::add_if(
+  InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream)
+{
+  impl_.add_if(first, last, stencil, pred, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class StencilIt, class Predicate>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::add_if_async(
+  InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream) noexcept
+{
+  impl_.add_if_async(first, last, stencil, pred, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class ProbeKey>
+[[nodiscard]] __device__ bool bloom_filter_ref<Key, Extent, Scope, Policy>::contains(
+  ProbeKey const& key) const
+{
+  return impl_.contains(key);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class CG, class ProbeKey>
+[[nodiscard]] __device__ bool bloom_filter_ref<Key, Extent, Scope, Policy>::contains(
+  CG const& group, ProbeKey const& key) const
+{
+  return impl_.contains(group, key);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class OutputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::contains(
+  InputIt first, InputIt last, OutputIt output_begin, cuda::stream_ref stream) const
+{
+  impl_.contains(first, last, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class OutputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::contains_async(
+  InputIt first, InputIt last, OutputIt output_begin, cuda::stream_ref stream) const noexcept
+{
+  impl_.contains_async(first, last, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::contains_if(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate pred,
+  OutputIt output_begin,
+  cuda::stream_ref stream) const
+{
+  impl_.contains_if(first, last, stencil, pred, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+template <class InputIt, class StencilIt, class Predicate, class OutputIt>
+__host__ constexpr void bloom_filter_ref<Key, Extent, Scope, Policy>::contains_if_async(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate pred,
+  OutputIt output_begin,
+  cuda::stream_ref stream) const noexcept
+{
+  impl_.contains_if_async(first, last, stencil, pred, output_begin, stream);
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+[[nodiscard]] __host__ __device__ constexpr
+  typename bloom_filter_ref<Key, Extent, Scope, Policy>::word_type*
+  bloom_filter_ref<Key, Extent, Scope, Policy>::data() noexcept
+{
+  return impl_.data();
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+[[nodiscard]] __host__ __device__ constexpr
+  typename bloom_filter_ref<Key, Extent, Scope, Policy>::word_type const*
+  bloom_filter_ref<Key, Extent, Scope, Policy>::data() const noexcept
+{
+  return impl_.data();
+}
+
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+[[nodiscard]] __host__ __device__ constexpr
+  typename bloom_filter_ref<Key, Extent, Scope, Policy>::extent_type
+  bloom_filter_ref<Key, Extent, Scope, Policy>::block_extent() const noexcept
+{
+  return impl_.block_extent();
+}
+
+}  // namespace cuco
\ No newline at end of file
diff --git a/include/cuco/detail/bloom_filter/kernels.cuh b/include/cuco/detail/bloom_filter/kernels.cuh
new file mode 100644
index 000000000..1d514aa6b
--- /dev/null
+++ b/include/cuco/detail/bloom_filter/kernels.cuh
@@ -0,0 +1,92 @@
+/*
+ * 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/utility/cuda.cuh>
+
+#include <cooperative_groups.h>
+
+#include <cstdint>
+#include <iterator>
+
+namespace cuco::detail {
+
+CUCO_SUPPRESS_KERNEL_WARNINGS
+
+template <int32_t CGSize,
+          int32_t BlockSize,
+          class InputIt,
+          class StencilIt,
+          class Predicate,
+          class Ref>
+CUCO_KERNEL __launch_bounds__(BlockSize) void add_if_n(
+  InputIt first, cuco::detail::index_type n, StencilIt stencil, Predicate pred, Ref ref)
+{
+  namespace cg = cooperative_groups;
+
+  auto const loop_stride = cuco::detail::grid_stride() / CGSize;
+  auto idx               = cuco::detail::global_thread_id() / CGSize;
+
+  [[maybe_unused]] auto const tile = cg::tiled_partition<CGSize>(cg::this_thread_block());
+
+  while (idx < n) {
+    if (pred(*(stencil + idx))) {
+      typename std::iterator_traits<InputIt>::value_type const& insert_element{*(first + idx)};
+      ref.add(tile, insert_element);
+    }
+    idx += loop_stride;
+  }
+}
+
+template <int32_t CGSize,
+          int32_t BlockSize,
+          class InputIt,
+          class StencilIt,
+          class Predicate,
+          class OutputIt,
+          class Ref>
+CUCO_KERNEL __launch_bounds__(BlockSize) void contains_if_n(InputIt first,
+                                                            cuco::detail::index_type n,
+                                                            StencilIt stencil,
+                                                            Predicate pred,
+                                                            OutputIt out,
+                                                            Ref ref)
+{
+  namespace cg = cooperative_groups;
+
+  auto const loop_stride = cuco::detail::grid_stride() / CGSize;
+  auto idx               = cuco::detail::global_thread_id() / CGSize;
+
+  [[maybe_unused]] auto const tile = cg::tiled_partition<CGSize>(cg::this_thread_block());
+
+  if constexpr (CGSize == 1) {
+    while (idx < n) {
+      typename std::iterator_traits<InputIt>::value_type const& key = *(first + idx);
+      *(out + idx) = pred(*(stencil + idx)) ? ref.contains(key) : false;
+      idx += loop_stride;
+    }
+  } else {
+    auto const tile = cg::tiled_partition<CGSize>(cg::this_thread_block());
+    while (idx < n) {
+      typename std::iterator_traits<InputIt>::value_type const& key = *(first + idx);
+      auto const found = pred(*(stencil + idx)) ? ref.contains(tile, key) : false;
+      if (tile.thread_rank() == 0) { *(out + idx) = found; }
+      idx += loop_stride;
+    }
+  }
+}
+
+}  // namespace cuco::detail
\ No newline at end of file
diff --git a/include/cuco/detail/utility/cuda.cuh b/include/cuco/detail/utility/cuda.cuh
index ffc3ffd9e..4ec0f7fc5 100644
--- a/include/cuco/detail/utility/cuda.cuh
+++ b/include/cuco/detail/utility/cuda.cuh
@@ -17,6 +17,10 @@
 
 #include <cuco/detail/utility/cuda.hpp>
 
+#include <cooperative_groups.h>
+
+#include <cstdint>
+
 #if defined(CUCO_DISABLE_KERNEL_VISIBILITY_WARNING_SUPPRESSION)
 #define CUCO_SUPPRESS_KERNEL_WARNINGS
 #elif defined(__NVCC__) && (defined(__GNUC__) || defined(__clang__))
@@ -59,5 +63,27 @@ __device__ static index_type grid_stride() noexcept
   return index_type{gridDim.x} * index_type{blockDim.x};
 }
 
+/**
+ * @brief Constexpr helper to extract the size of a Cooperative Group.
+ *
+ * @tparam Tile The Cooperative Group type
+ */
+template <typename Tile>
+struct tile_size;
+
+/**
+ * @brief Specialization of `cuco::detail::tile_size` for 'cooperative_groups::thread_block_tile'.
+ *
+ * @tparam CGSize The Cooperative Group size
+ * @tparam ParentCG The Cooperative Group the tile has been created from
+ */
+template <uint32_t CGSize, class ParentCG>
+struct tile_size<cooperative_groups::thread_block_tile<CGSize, ParentCG>> {
+  static constexpr uint32_t value = CGSize;  ///< Size of the `thread_block_tile`
+};
+
+template <typename Tile>
+__device__ constexpr uint32_t tile_size_v = tile_size<Tile>::value;
+
 }  // namespace detail
 }  // namespace cuco
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index e1e036048..d34c89d98 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -135,3 +135,8 @@ ConfigureTest(HYPERLOGLOG_TEST
     hyperloglog/unique_sequence_test.cu
     hyperloglog/spark_parity_test.cu
     hyperloglog/device_ref_test.cu)
+
+###################################################################################################
+# - bloom_filter ----------------------------------------------------------------------------------
+ConfigureTest(BLOOM_FILTER_TEST
+    bloom_filter/unique_sequence_test.cu)
diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu
new file mode 100644
index 000000000..3919e77bf
--- /dev/null
+++ b/tests/bloom_filter/unique_sequence_test.cu
@@ -0,0 +1,105 @@
+/*
+ * 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 <test_utils.hpp>
+
+#include <cuco/bloom_filter.cuh>
+
+#include <cuda/functional>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/sequence.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+#include <catch2/generators/catch_generators.hpp>
+
+using size_type = int32_t;
+
+template <typename Filter>
+void test_unique_sequence(Filter& filter, size_type num_keys)
+{
+  using Key = typename Filter::key_type;
+
+  thrust::device_vector<Key> keys(num_keys);
+
+  thrust::sequence(thrust::device, keys.begin(), keys.end());
+
+  thrust::device_vector<bool> contained(num_keys, false);
+
+  auto is_even =
+    cuda::proclaim_return_type<bool>([] __device__(auto const& i) { return i % 2 == 0; });
+
+  SECTION("Non-inserted keys should not be contained.")
+  {
+    filter.contains(keys.begin(), keys.end(), contained.begin());
+    REQUIRE(cuco::test::none_of(contained.begin(), contained.end(), thrust::identity{}));
+  }
+
+  SECTION("All inserted keys should be contained.")
+  {
+    filter.add(keys.begin(), keys.end());
+    filter.contains(keys.begin(), keys.end(), contained.begin());
+    REQUIRE(cuco::test::all_of(contained.begin(), contained.end(), thrust::identity{}));
+  }
+
+  SECTION("After clearing the flter no keys should be contained.")
+  {
+    filter.clear();
+    filter.contains(keys.begin(), keys.end(), contained.begin());
+    REQUIRE(cuco::test::none_of(contained.begin(), contained.end(), thrust::identity{}));
+  }
+
+  SECTION("All conditionally inserted keys should be contained")
+  {
+    filter.add_if(keys.begin(), keys.end(), thrust::counting_iterator<std::size_t>(0), is_even);
+    filter.contains_if(keys.begin(),
+                       keys.end(),
+                       thrust::counting_iterator<std::size_t>(0),
+                       is_even,
+                       contained.begin());
+    REQUIRE(cuco::test::equal(
+      contained.begin(),
+      contained.end(),
+      thrust::counting_iterator<std::size_t>(0),
+      cuda::proclaim_return_type<bool>([] __device__(auto const& idx_contained, auto const& idx) {
+        return ((idx % 2) == 0) == idx_contained;
+      })));
+  }
+
+  // TODO test FPR but how?
+}
+
+TEMPLATE_TEST_CASE_SIG("Unique sequence",
+                       "",
+                       ((class Key, class Policy), Key, Policy),
+                       (int32_t, cuco::bloom_filter_policy<cuco::xxhash_64<int32_t>, uint32_t, 1>),
+                       (int32_t, cuco::bloom_filter_policy<cuco::xxhash_64<int32_t>, uint32_t, 8>),
+                       (int32_t, cuco::bloom_filter_policy<cuco::xxhash_64<int32_t>, uint64_t, 1>),
+                       (int32_t, cuco::bloom_filter_policy<cuco::xxhash_64<int32_t>, uint64_t, 8>))
+{
+  using filter_type =
+    cuco::bloom_filter<Key, cuco::extent<size_t>, cuda::thread_scope_device, Policy>;
+  constexpr size_type num_keys{400};
+
+  uint32_t pattern_bits =
+    GENERATE(Policy::words_per_block, Policy::words_per_block + 1, Policy::words_per_block + 2);
+
+  auto filter = filter_type{1000, {}, {pattern_bits}};
+
+  test_unique_sequence(filter, num_keys);
+}