diff --git a/include/cuco/detail/open_addressing/kernels.cuh b/include/cuco/detail/open_addressing/kernels.cuh
index 7c41e89d6..1505041fe 100644
--- a/include/cuco/detail/open_addressing/kernels.cuh
+++ b/include/cuco/detail/open_addressing/kernels.cuh
@@ -565,6 +565,98 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void count(InputIt first,
   if (threadIdx.x == 0) { count->fetch_add(block_count, cuda::std::memory_order_relaxed); }
 }
 
+/**
+ * @brief Counts the occurrences of keys in `[first, last)` contained in the container
+ * if `pred` of the corresponding stencil returns true.
+ *
+ * @tparam IsOuter Flag indicating whether it's an outer count or not
+ * @tparam CGSize Number of threads in each CG
+ * @tparam BlockSize Number of threads in each block
+ * @tparam InputIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+ * @tparam AtomicT Atomic counter type
+ * @tparam Ref Type of non-owning device container ref allowing access to storage
+ *
+ * @param first Beginning of the sequence of input elements
+ * @param n Number of input elements
+ * @param stencil Beginning of the stencil sequence
+ * @param pred Predicate to test on every element in the range `[stencil, stencil + n)`
+ * @param count Number of matches
+ * @param ref Non-owning container device ref used to access the slot storage
+ */
+template <bool IsOuter,
+          int CGSize,
+          int BlockSize,
+          typename InputIt,
+          typename StencilIt,
+          typename Predicate,
+          typename AtomicT,
+          typename Ref>
+CUCO_KERNEL __launch_bounds__(BlockSize) void count_if(InputIt first,
+                                                       cuco::detail::index_type n,
+                                                       StencilIt stencil,
+                                                       Predicate pred,
+                                                       AtomicT* count,
+                                                       Ref ref)
+{
+  using size_type = typename Ref::size_type;
+
+  size_type constexpr outer_min_count = 1;
+
+  using BlockReduce = cub::BlockReduce<size_type, BlockSize>;
+  __shared__ typename BlockReduce::TempStorage temp_storage;
+
+  size_type thread_count = 0;
+
+  auto const loop_stride = cuco::detail::grid_stride() / CGSize;
+  auto idx               = cuco::detail::global_thread_id() / CGSize;
+
+  while (idx < n) {
+    if constexpr (CGSize == 1) {
+      if (pred(*(stencil + idx))) {
+        typename cuda::std::iterator_traits<InputIt>::value_type const key = *(first + idx);
+
+        if constexpr (IsOuter) {
+          thread_count += max(ref.count(key), outer_min_count);
+        } else {
+          thread_count += ref.count(key);
+        }
+      } else if constexpr (IsOuter) {
+        thread_count += outer_min_count;
+      }
+    } else {
+      auto const tile =
+        cooperative_groups::tiled_partition<CGSize, cooperative_groups::thread_block>(
+          cooperative_groups::this_thread_block());
+
+      if (pred(*(stencil + idx))) {
+        typename cuda::std::iterator_traits<InputIt>::value_type const key = *(first + idx);
+
+        if constexpr (IsOuter) {
+          auto temp_count = ref.count(tile, key);
+
+          if (tile.all(temp_count == 0) && tile.thread_rank() == 0) { ++temp_count; }
+
+          thread_count += temp_count;
+        } else {
+          thread_count += ref.count(tile, key);
+        }
+      } else if constexpr (IsOuter) {
+        if (tile.thread_rank() == 0) { thread_count += outer_min_count; }
+      }
+    }
+
+    idx += loop_stride;
+  }
+
+  auto const block_count = BlockReduce(temp_storage).Sum(thread_count);
+
+  if (threadIdx.x == 0) { count->fetch_add(block_count, cuda::std::memory_order_relaxed); }
+}
+
 /**
  * @brief Counts the occurrences of each key in `[first, last)` contained in the container
  * and stores the counts in the output array.
@@ -699,6 +791,93 @@ CUCO_KERNEL void retrieve(InputProbeIt input_probe,
   }
 }
 
+/**
+ * @brief Retrieves the equivalent container elements of all keys in the range `[input_probe,
+ * input_probe + n)` if `pred` of the corresponding stencil returns true.
+ *
+ * If key `k = *(input_probe + i)` has one or more matches in the container  and `pred` of
+ * its corresponding stencil is true, copies `k` to `output_probe` and associated slot
+ * contents to `output_match`, respectively. The output order is unspecified.
+ *
+ * @tparam IsOuter Flag indicating whether it's an outer count or not
+ * @tparam BlockSize The size of the thread block
+ * @tparam TileStride Number of tile batches assigned to each thread block
+ * @tparam InputProbeIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+ * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+ * convertible to the `InputProbeIt`'s `value_type`
+ * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+ * convertible to the container's `value_type`
+ * @tparam AtomicCounter Integral atomic type that follows the same semantics as
+ * `cuda::(std::)atomic(_ref)`
+ * @tparam Ref Type of non-owning device ref allowing access to storage
+ *
+ * @param input_probe Beginning of the sequence of input keys
+ * @param n Number of the keys to query
+ * @param stencil Beginning of the stencil sequence
+ * @param pred Predicate to test on every element in the range `[stencil, stencil + n)`
+ * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+ * `output_match`
+ * @param output_match Beginning of the sequence of matching elements
+ * @param atomic_counter Pointer to an atomic object of integral type that is used to count the
+ * number of output elements
+ * @param ref Non-owning container device ref used to access the slot storage
+ */
+template <bool IsOuter,
+          int BlockSize,
+          int TileStride,
+          class InputProbeIt,
+          class StencilIt,
+          class Predicate,
+          class OutputProbeIt,
+          class OutputMatchIt,
+          class AtomicCounter,
+          class Ref>
+CUCO_KERNEL void retrieve_if(InputProbeIt input_probe,
+                             cuco::detail::index_type n,
+                             StencilIt stencil,
+                             Predicate pred,
+                             OutputProbeIt output_probe,
+                             OutputMatchIt output_match,
+                             AtomicCounter* atomic_counter,
+                             Ref ref)
+{
+  namespace cg = cooperative_groups;
+
+  auto const block               = cg::this_thread_block();
+  auto constexpr tiles_in_block  = BlockSize / Ref::cg_size;
+  auto constexpr tiles_per_block = TileStride * tiles_in_block;
+
+  auto const block_begin_offset = block.group_index().x * tiles_per_block;
+  auto const block_end_offset =
+    min(n, static_cast<cuco::detail::index_type>(block_begin_offset + tiles_per_block));
+
+  if (block_begin_offset < block_end_offset) {
+    if constexpr (IsOuter) {
+      ref.template retrieve_outer_if<BlockSize>(block,
+                                                input_probe + block_begin_offset,
+                                                input_probe + block_end_offset,
+                                                stencil + block_begin_offset,
+                                                pred,
+                                                output_probe,
+                                                output_match,
+                                                *atomic_counter);
+    } else {
+      ref.template retrieve_if<BlockSize>(block,
+                                          input_probe + block_begin_offset,
+                                          input_probe + block_end_offset,
+                                          stencil + block_begin_offset,
+                                          pred,
+                                          output_probe,
+                                          output_match,
+                                          *atomic_counter);
+    }
+  }
+}
+
 template <int BlockSize, typename ContainerRef, typename Predicate>
 CUCO_KERNEL __launch_bounds__(BlockSize) void rehash(
   typename ContainerRef::storage_ref_type storage_ref,
diff --git a/include/cuco/detail/open_addressing/open_addressing_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_impl.cuh
index cfae2d38f..a4d852a7c 100644
--- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh
+++ b/include/cuco/detail/open_addressing/open_addressing_impl.cuh
@@ -686,6 +686,123 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
       first, last, output_probe, output_match, container_ref, stream);
   }
 
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred( *(stencil + i) )` returns true,
+   * copies `k` to `output_probe` and associated slot contents to `output_match`, respectively.
+   * The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count()` to determine the size of the output range.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam InputProbeIt Device accessible input 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 OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   * @tparam Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param container_ref Non-owning device reference to the container
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt,
+            class Ref>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_if(InputProbeIt first,
+                                                      InputProbeIt last,
+                                                      StencilIt stencil,
+                                                      Predicate const& pred,
+                                                      OutputProbeIt output_probe,
+                                                      OutputMatchIt output_match,
+                                                      Ref container_ref,
+                                                      cuda::stream_ref stream) const
+  {
+    auto constexpr is_outer = false;
+    return this->retrieve_if_impl<is_outer>(
+      first, last, stencil, pred, output_probe, output_match, container_ref, stream);
+  }
+
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred( *(stencil + i) )` returns true,
+   * copies `k` to `output_probe` and associated slot contents to `output_match`, respectively.
+   * The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_outer()` to determine the size of the output range.
+   *
+   * If a key `k` has no matches in the container or `pred( *(stencil + i) )` returns false,
+   * then `{key, empty_slot_sentinel}` will be added to the output sequence.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam InputProbeIt Device accessible input 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 OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   * @tparam Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param container_ref Non-owning device reference to the container
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt,
+            class Ref>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_outer_if(InputProbeIt first,
+                                                            InputProbeIt last,
+                                                            StencilIt stencil,
+                                                            Predicate const& pred,
+                                                            OutputProbeIt output_probe,
+                                                            OutputMatchIt output_match,
+                                                            Ref container_ref,
+                                                            cuda::stream_ref stream) const
+  {
+    auto constexpr is_outer = true;
+    return this->retrieve_if_impl<is_outer>(
+      first, last, stencil, pred, output_probe, output_match, container_ref, stream);
+  }
+
   /**
    * @brief Counts the occurrences of keys in `[first, last)` contained in the container
    *
@@ -694,6 +811,7 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
    *
    * @param first Beginning of the sequence of keys to count
    * @param last End of the sequence of keys to count
+   * @param container_ref Non-owning device reference to the container
    * @param stream CUDA stream used for count
    *
    * @return The sum of total occurrences of all keys in `[first, last)`
@@ -708,6 +826,39 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
     return this->count<is_outer>(first, last, container_ref, stream);
   }
 
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the container
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @tparam Input Device accessible input 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 Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 container_ref Non-owning device reference to the container
+   * @param stream CUDA stream used for count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)`
+   */
+  template <typename InputIt, typename StencilIt, typename Predicate, typename Ref>
+  [[nodiscard]] size_type count_if(InputIt first,
+                                   InputIt last,
+                                   StencilIt stencil,
+                                   Predicate pred,
+                                   Ref container_ref,
+                                   cuda::stream_ref stream) const
+  {
+    auto constexpr is_outer = false;
+    return this->count_if<is_outer>(first, last, stencil, pred, container_ref, stream);
+  }
+
   /**
    * @brief Counts the occurrences of keys in `[first, last)` contained in the container
    *
@@ -718,6 +869,7 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
    *
    * @param first Beginning of the sequence of keys to count
    * @param last End of the sequence of keys to count
+   * @param container_ref Non-owning device reference to the container
    * @param stream CUDA stream used for count
    *
    * @return The sum of total occurrences of all keys in `[first, last)`
@@ -732,6 +884,43 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
     return this->count<is_outer>(first, last, container_ref, stream);
   }
 
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the container
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note If a given key has no matches, or `pred` of its corresponding stencil is false,
+   * its occurrence is 1.
+   *
+   * @tparam Input Device accessible input 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 Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 container_ref Non-owning device reference to the container
+   * @param stream CUDA stream used for count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)`
+   */
+  template <typename InputIt, typename StencilIt, typename Predicate, typename Ref>
+  [[nodiscard]] size_type count_outer_if(InputIt first,
+                                         InputIt last,
+                                         StencilIt stencil,
+                                         Predicate pred,
+                                         Ref container_ref,
+                                         cuda::stream_ref stream) const
+  {
+    auto constexpr is_outer = true;
+
+    return this->count_if<is_outer>(first, last, stencil, pred, container_ref, stream);
+  }
+
   /**
    * @brief Counts the number of occurrences of each query key in the container
    *
@@ -1191,11 +1380,12 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
    * @note If `IsOuter` is `true`, the occurrence of a non-match key is 1. Else, it's 0.
    *
    * @tparam IsOuter Flag indicating whether it's an outer count or not
-   * @tparam Input Device accessible input iterator
+   * @tparam InputIt Device accessible input iterator
    * @tparam Ref Type of non-owning device container ref allowing access to storage
    *
    * @param first Beginning of the sequence of keys to count
    * @param last End of the sequence of keys to count
+   * @param container_ref Non-owning device reference to the container
    * @param stream CUDA stream used for count
    *
    * @return The sum of total occurrences of all keys in `[first, last)`
@@ -1223,6 +1413,56 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
     return counter.load_to_host(stream);
   }
 
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the container
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note If `IsOuter` is `true`, the occurrence of a non-match key or a non-predicate
+   * key is 1. Else, it's 0.
+   *
+   * @tparam IsOuter Flag indicating whether it's an outer count or not
+   * @tparam InputIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 container_ref Non-owning device reference to the container
+   * @param stream CUDA stream used for count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)`
+   */
+  template <bool IsOuter, typename InputIt, typename StencilIt, typename Predicate, typename Ref>
+  [[nodiscard]] size_type count_if(InputIt first,
+                                   InputIt last,
+                                   StencilIt stencil,
+                                   Predicate pred,
+                                   Ref container_ref,
+                                   cuda::stream_ref stream) const
+  {
+    auto const num_keys = cuco::detail::distance(first, last);
+    if (num_keys == 0) { return 0; }
+
+    auto counter =
+      detail::counter_storage<size_type, thread_scope, allocator_type>{this->allocator(), stream};
+
+    counter.reset(stream);
+
+    auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);
+
+    detail::open_addressing_ns::count_if<IsOuter, cg_size, cuco::detail::default_block_size()>
+      <<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
+        first, num_keys, stencil, pred, counter.data(), container_ref);
+
+    return counter.load_to_host(stream);
+  }
+
   /**
    * @brief Counts the number of occurrences of each query key in the container
    *
@@ -1326,6 +1566,85 @@ class open_addressing_impl : private open_addressing_compatible<Key, Value, Prob
     return {output_probe + num_retrieved, output_match + num_retrieved};
   }
 
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`
+   * if `pred` of the corresponding stencil returns true..
+   *
+   * If key `k = *(input_probe + i)` has one or more matches in the container  and `pred` of
+   * its corresponding stencil is true, copies `k` to `output_probe` and associated slot
+   * contents to `output_match`, respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count()/count_outer()` to determine the size of the output range.
+   *
+   * If `IsOuter == true` and a key `k` has no matches in the container, or `pred` of the
+   * corresponding stencil is false, then `{key, empty_slot_sentinel}` will be added to the
+   * output sequence.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam IsOuter Flag indicating if an inner or outer retrieve operation should be performed
+   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `key_type`
+   * @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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `key_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   * @tparam Ref Type of non-owning device container ref allowing access to storage
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param container_ref Non-owning device reference to the container
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <bool IsOuter,
+            class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt,
+            class Ref>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_if_impl(InputProbeIt first,
+                                                           InputProbeIt last,
+                                                           StencilIt stencil,
+                                                           Predicate const& pred,
+                                                           OutputProbeIt output_probe,
+                                                           OutputMatchIt output_match,
+                                                           Ref container_ref,
+                                                           cuda::stream_ref stream) const
+  {
+    auto const n = detail::distance(first, last);
+    if (n == 0) { return {output_probe, output_match}; }
+
+    using counter_type = detail::counter_storage<size_type, thread_scope, allocator_type>;
+    auto counter       = counter_type{this->allocator(), stream};
+    counter.reset(stream.get());
+
+    auto constexpr block_size  = cuco::detail::default_block_size();
+    auto constexpr grid_stride = 4;
+    auto const grid_size       = cuco::detail::grid_size(n, cg_size, grid_stride, block_size);
+
+    detail::open_addressing_ns::retrieve_if<IsOuter, block_size, grid_stride>
+      <<<grid_size, block_size, 0, stream.get()>>>(
+        first, n, stencil, pred, output_probe, output_match, counter.data(), container_ref);
+
+    auto const num_retrieved = counter.load_to_host(stream.get());
+
+    return {output_probe + num_retrieved, output_match + num_retrieved};
+  }
+
   /**
    * @brief Extracts the key from a given slot.
    *
diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
index b84d07624..03466f888 100644
--- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
+++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
@@ -1175,6 +1175,67 @@ class open_addressing_ref_impl
       block, input_probe_begin, n, stencil, pred, output_probe, output_match, atomic_counter);
   }
 
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[input_probe_begin,
+   * input_probe_end)` if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred( *(stencil + i) )` returns true,
+   * copies `k` to `output_probe` and associated slot contents to `output_match`,
+   * respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_outer_if()` to determine the size of the output range.
+   *
+   * If a key `k` has no matches in the container, or `pred` of the corresponding stencil is
+   * false, then `{key, empty_slot_sentinel}` will be added to the output sequence.
+   *
+   * @tparam BlockSize Size of the thread block this operation is executed in
+   * @tparam InputProbeIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   * @tparam AtomicCounter Integral atomic counter type that follows the same semantics as
+   * `cuda::(std::)atomic(_ref)`
+   *
+   * @param block Thread block this operation is executed in
+   * @param input_probe_begin Beginning of the input sequence of keys
+   * @param input_probe_end End of the input sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil + n)`
+   * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param atomic_counter Atomic object of integral type that is used to count the
+   * number of output elements
+   */
+  template <int BlockSize,
+            class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt,
+            class AtomicCounter>
+  __device__ void retrieve_outer_if(cooperative_groups::thread_block const& block,
+                                    InputProbeIt input_probe_begin,
+                                    InputProbeIt input_probe_end,
+                                    StencilIt stencil,
+                                    Predicate pred,
+                                    OutputProbeIt output_probe,
+                                    OutputMatchIt output_match,
+                                    AtomicCounter& atomic_counter) const
+  {
+    auto constexpr is_outer = true;
+    auto const n            = cuco::detail::distance(input_probe_begin, input_probe_end);
+
+    this->retrieve_impl<is_outer, BlockSize>(
+      block, input_probe_begin, n, stencil, pred, output_probe, output_match, atomic_counter);
+  }
+
   /**
    * @brief Retrieves all the slots corresponding to all keys in the range `[input_probe_begin,
    * input_probe_end)`.
@@ -1387,6 +1448,28 @@ class open_addressing_ref_impl
           if (*probing_iter == init_idx) { running = false; }
         }  // while running
       }  // if active_flag
+      else if constexpr (IsOuter) {
+        // Predicate rejected this key. It is already known to be a miss,
+        // so do not probe the hash table. Emit the outer sentinel directly.
+        if (idx < n and probing_tile.thread_rank() == 0) {
+          auto ref = cuda::atomic_ref<cuda::std::int32_t, cuda::thread_scope_block>{
+            counters[flushing_tile_id]};
+          auto const output_idx      = ref.fetch_add(1, cuda::memory_order_relaxed);
+          probe_type const probe_key = *(input_probe + idx);
+
+          buffers[flushing_tile_id][output_idx] = {probe_key, this->empty_slot_sentinel()};
+        }
+        active_flushing_tile.sync();
+        // if the buffer has not enough empty slots for the next iteration
+        if (counters[flushing_tile_id] > (buffer_size - max_matches_per_step)) {
+          flush_buffers(active_flushing_tile);
+          active_flushing_tile.sync();
+
+          // reset buffer counter
+          if (active_flushing_tile.thread_rank() == 0) { counters[flushing_tile_id] = 0; }
+          active_flushing_tile.sync();
+        }
+      }
 
       // onto the next key
       idx += stride;
diff --git a/include/cuco/detail/static_multiset/static_multiset.inl b/include/cuco/detail/static_multiset/static_multiset.inl
index 811c07781..e39eb830f 100644
--- a/include/cuco/detail/static_multiset/static_multiset.inl
+++ b/include/cuco/detail/static_multiset/static_multiset.inl
@@ -457,6 +457,100 @@ static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>
     stream);
 }
 
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <typename InputIt, typename StencilIt, typename Predicate>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::size_type
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::count_if(
+  InputIt first, InputIt last, StencilIt stencil, Predicate const& pred, cuda::stream_ref stream)
+  const
+{
+  return this->count_if(first, last, stencil, pred, key_eq(), hash_function(), stream);
+}
+
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <typename InputIt,
+          typename StencilIt,
+          typename Predicate,
+          typename ProbeKeyEqual,
+          typename ProbeHash>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::size_type
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::count_if(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate const& pred,
+  ProbeKeyEqual const& probe_key_equal,
+  ProbeHash const& probe_hash,
+  cuda::stream_ref stream) const
+{
+  return impl_->count_if(
+    first,
+    last,
+    stencil,
+    pred,
+    ref(op::count).rebind_key_eq(probe_key_equal).rebind_hash_function(probe_hash),
+    stream);
+}
+
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <typename InputIt, typename StencilIt, typename Predicate>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::size_type
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::count_outer_if(
+  InputIt first, InputIt last, StencilIt stencil, Predicate const& pred, cuda::stream_ref stream)
+  const
+{
+  return this->count_outer_if(first, last, stencil, pred, key_eq(), hash_function(), stream);
+}
+
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <typename InputIt,
+          typename StencilIt,
+          typename Predicate,
+          typename ProbeKeyEqual,
+          typename ProbeHash>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::size_type
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::count_outer_if(
+  InputIt first,
+  InputIt last,
+  StencilIt stencil,
+  Predicate const& pred,
+  ProbeKeyEqual const& probe_key_equal,
+  ProbeHash const& probe_hash,
+  cuda::stream_ref stream) const
+{
+  return impl_->count_outer_if(
+    first,
+    last,
+    stencil,
+    pred,
+    ref(op::count).rebind_key_eq(probe_key_equal).rebind_hash_function(probe_hash),
+    stream);
+}
+
 template <class Key,
           class Extent,
           cuda::thread_scope Scope,
@@ -577,6 +671,98 @@ static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>
   return impl_->retrieve_outer(first, last, output_probe, output_match, probe_ref, stream);
 }
 
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <class InputProbeIt,
+          class StencilIt,
+          class Predicate,
+          class OutputProbeIt,
+          class OutputMatchIt>
+std::pair<OutputProbeIt, OutputMatchIt>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::retrieve_if(
+  InputProbeIt first,
+  InputProbeIt last,
+  StencilIt stencil,
+  Predicate const& pred,
+  OutputProbeIt output_probe,
+  OutputMatchIt output_match,
+  cuda::stream_ref stream) const
+{
+  return impl_->retrieve_if(
+    first, last, stencil, pred, output_probe, output_match, this->ref(op::retrieve), stream);
+}
+
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <class InputProbeIt,
+          class StencilIt,
+          class Predicate,
+          class ProbeEqual,
+          class ProbeHash,
+          class OutputProbeIt,
+          class OutputMatchIt>
+std::pair<OutputProbeIt, OutputMatchIt>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::retrieve_if(
+  InputProbeIt first,
+  InputProbeIt last,
+  StencilIt stencil,
+  Predicate const& pred,
+  ProbeEqual const& probe_equal,
+  ProbeHash const& probe_hash,
+  OutputProbeIt output_probe,
+  OutputMatchIt output_match,
+  cuda::stream_ref stream) const
+{
+  auto const probe_ref =
+    this->ref(op::retrieve).rebind_key_eq(probe_equal).rebind_hash_function(probe_hash);
+
+  return impl_->retrieve_if(
+    first, last, stencil, pred, output_probe, output_match, probe_ref, stream);
+}
+
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class KeyEqual,
+          class ProbingScheme,
+          class Allocator,
+          class Storage>
+template <class InputProbeIt,
+          class StencilIt,
+          class Predicate,
+          class ProbeEqual,
+          class ProbeHash,
+          class OutputProbeIt,
+          class OutputMatchIt>
+std::pair<OutputProbeIt, OutputMatchIt>
+static_multiset<Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::retrieve_outer_if(
+  InputProbeIt first,
+  InputProbeIt last,
+  StencilIt stencil,
+  Predicate const& pred,
+  ProbeEqual const& probe_equal,
+  ProbeHash const& probe_hash,
+  OutputProbeIt output_probe,
+  OutputMatchIt output_match,
+  cuda::stream_ref stream) const
+{
+  auto const probe_ref =
+    this->ref(op::retrieve).rebind_key_eq(probe_equal).rebind_hash_function(probe_hash);
+
+  return impl_->retrieve_outer_if(
+    first, last, stencil, pred, output_probe, output_match, probe_ref, stream);
+}
+
 template <class Key,
           class Extent,
           cuda::thread_scope Scope,
diff --git a/include/cuco/detail/static_multiset/static_multiset_ref.inl b/include/cuco/detail/static_multiset/static_multiset_ref.inl
index c34644c61..f1117d117 100644
--- a/include/cuco/detail/static_multiset/static_multiset_ref.inl
+++ b/include/cuco/detail/static_multiset/static_multiset_ref.inl
@@ -654,7 +654,7 @@ class operator_impl<
    * The output order is unspecified.
    *
    * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
-   * Use `count()` to determine the size of the output range.
+   * Use `count_if()` to determine the size of the output range.
    *
    * @tparam BlockSize Size of the thread block this operation is executed in
    * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
@@ -707,6 +707,72 @@ class operator_impl<
                                                output_match,
                                                atomic_counter);
   }
+
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[input_probe_begin,
+   * input_probe_end)` if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred` of its corresponding stencil
+   * is true, copies `k` to `output_probe` and associated slot contents to `output_match`,
+   * respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_outer_if()` to determine the size of the output range.
+   *
+   * If a key `k` has no matches in the container, or `pred` of its corresponding stencil is
+   * false, then `{key, empty_slot_sentinel}` will be added to the output sequence.
+   *
+   * @tparam BlockSize Size of the thread block this operation is executed in
+   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `key_type`
+   * @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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `key_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   * @tparam AtomicCounter Atomic counter type that follows the same semantics as
+   * `cuda::atomic(_ref)`
+   *
+   * @param block Thread block this operation is executed in
+   * @param input_probe_begin Beginning of the input sequence of keys
+   * @param input_probe_end End of the input sequence of keys
+   * @param stencil Beginning of the stencil sequence
+   * @param pred Predicate to test on every element in the range `[stencil, stencil + n)`
+   * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param atomic_counter Counter that is used to determine the next free position in the output
+   * sequences
+   */
+  template <int BlockSize,
+            class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt,
+            class AtomicCounter>
+  __device__ void retrieve_outer_if(cooperative_groups::thread_block const& block,
+                                    InputProbeIt input_probe_begin,
+                                    InputProbeIt input_probe_end,
+                                    StencilIt stencil,
+                                    Predicate pred,
+                                    OutputProbeIt output_probe,
+                                    OutputMatchIt output_match,
+                                    AtomicCounter& atomic_counter) const
+  {
+    auto const& ref_ = static_cast<ref_type const&>(*this);
+    ref_.impl_.template retrieve_outer_if<BlockSize>(block,
+                                                     input_probe_begin,
+                                                     input_probe_end,
+                                                     stencil,
+                                                     pred,
+                                                     output_probe,
+                                                     output_match,
+                                                     atomic_counter);
+  }
 };
 
 template <typename Key,
diff --git a/include/cuco/static_multiset.cuh b/include/cuco/static_multiset.cuh
index 77bb1ffb0..6f6f26eb1 100644
--- a/include/cuco/static_multiset.cuh
+++ b/include/cuco/static_multiset.cuh
@@ -654,6 +654,72 @@ class static_multiset {
                   ProbeHash const& probe_hash,
                   cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
 
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the multiset
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note This function synchronizes the given stream.
+   *
+   * @tparam Input Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)`
+   */
+  template <typename InputIt, typename StencilIt, typename Predicate>
+  size_type count_if(InputIt first,
+                     InputIt last,
+                     StencilIt stencil,
+                     Predicate const& pred,
+                     cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
+
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the multiset
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note This function synchronizes the given stream.
+   *
+   * @tparam Input Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam ProbeKeyEqual Binary callable
+   * @tparam ProbeHash Unary hash callable
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 probe_key_equal Binary callable to compare two keys for equality
+   * @param probe_hash Unary callable to hash a given key
+   * @param stream CUDA stream used for count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)`
+   */
+  template <typename InputIt,
+            typename StencilIt,
+            typename Predicate,
+            typename ProbeKeyEqual,
+            typename ProbeHash>
+  size_type count_if(InputIt first,
+                     InputIt last,
+                     StencilIt stencil,
+                     Predicate const& pred,
+                     ProbeKeyEqual const& probe_key_equal,
+                     ProbeHash const& probe_hash,
+                     cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
+
   /**
    * @brief Counts the occurrences of keys in `[first, last)` contained in the multiset
    *
@@ -680,6 +746,78 @@ class static_multiset {
                         ProbeHash const& probe_hash,
                         cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
 
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the multiset
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note This function synchronizes the given stream.
+   * @note If a given key has no matches or `pred` of its corresponding stencil is false,
+   * its occurrence is 1.
+   *
+   * @tparam Input Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)` where keys have no matches
+   * are considered to have a single occurrence.
+   */
+  template <typename InputIt, typename StencilIt, typename Predicate>
+  size_type count_outer_if(InputIt first,
+                           InputIt last,
+                           StencilIt stencil,
+                           Predicate const& pred,
+                           cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
+
+  /**
+   * @brief Counts the occurrences of keys in `[first, last)` contained in the multiset
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * @note This function synchronizes the given stream.
+   * @note If a given key has no matches or `pred` of its corresponding stencil is false,
+   * its occurrence is 1.
+   *
+   * @tparam Input Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam ProbeKeyEqual Binary callable
+   * @tparam ProbeHash Unary hash callable
+   *
+   * @param first Beginning of the sequence of keys to count
+   * @param last End of the sequence of keys to count
+   * @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 probe_key_equal Binary callable to compare two keys for equality
+   * @param probe_hash Unary callable to hash a given key
+   * @param stream CUDA stream used for count
+   *
+   * @return The sum of total occurrences of all keys in `[first, last)` where keys have no matches
+   * are considered to have a single occurrence.
+   */
+  template <typename InputIt,
+            typename StencilIt,
+            typename Predicate,
+            typename ProbeKeyEqual,
+            typename ProbeHash>
+  size_type count_outer_if(InputIt first,
+                           InputIt last,
+                           StencilIt stencil,
+                           Predicate const& pred,
+                           ProbeKeyEqual const& probe_key_equal,
+                           ProbeHash const& probe_hash,
+                           cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
+
   /**
    * @brief Counts the number of occurrences of each query key in the multiset
    *
@@ -873,6 +1011,172 @@ class static_multiset {
                                                          cuda::stream_ref stream = cuda::stream_ref{
                                                            cudaStream_t{nullptr}}) const;
 
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred` of its corresponding stencil
+   * is true, copies `k` to `output_probe` and associated slot contents to `output_match`,
+   * respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_if()` to determine the size of the output range.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam InputProbeIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class OutputProbeIt,
+            class OutputMatchIt>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_if(InputProbeIt first,
+                                                      InputProbeIt last,
+                                                      StencilIt stencil,
+                                                      Predicate const& pred,
+                                                      OutputProbeIt output_probe,
+                                                      OutputMatchIt output_match,
+                                                      cuda::stream_ref stream = cuda::stream_ref{
+                                                        cudaStream_t{nullptr}}) const;
+
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred` of its corresponding stencil
+   * is true, copies `k` to `output_probe` and associated slot contents to `output_match`,
+   * respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_if()` to determine the size of the output range.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam InputProbeIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam ProbeEqual Binary callable equal type
+   * @tparam ProbeHash Unary callable hasher type that can be constructed from
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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 probe_equal The binary function to compare set keys and probe keys for equality
+   * @param probe_hash The unary function to hash probe keys
+   * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class ProbeEqual,
+            class ProbeHash,
+            class OutputProbeIt,
+            class OutputMatchIt>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_if(InputProbeIt first,
+                                                      InputProbeIt last,
+                                                      StencilIt stencil,
+                                                      Predicate const& pred,
+                                                      ProbeEqual const& probe_equal,
+                                                      ProbeHash const& probe_hash,
+                                                      OutputProbeIt output_probe,
+                                                      OutputMatchIt output_match,
+                                                      cuda::stream_ref stream = cuda::stream_ref{
+                                                        cudaStream_t{nullptr}}) const;
+
+  /**
+   * @brief Retrieves all the slots corresponding to all keys in the range `[first, last)`.
+   * if `pred` of the corresponding stencil returns true.
+   *
+   * If key `k = *(first + i)` exists in the container and `pred` of its corresponding stencil
+   * is true, copies `k` to `output_probe` and associated slot contents to `output_match`,
+   * respectively. The output order is unspecified.
+   *
+   * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
+   * Use `count_outer_if()` to determine the size of the output range.
+   *
+   * If a key `k` has no matches in the container, or `pred` of its corresponding stencil is
+   * false, then `{key, empty_slot_sentinel}` will be added to the output sequence.
+   *
+   * This function synchronizes the given CUDA stream.
+   *
+   * @tparam InputProbeIt Device accessible input 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 `std::iterator_traits<StencilIt>::value_type`
+   * @tparam ProbeEqual Binary callable equal type
+   * @tparam ProbeHash Unary callable hasher type that can be constructed from
+   * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+   * convertible to the `InputProbeIt`'s `value_type`
+   * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+   * convertible to the container's `value_type`
+   *
+   * @param first Beginning of the input sequence of keys
+   * @param last End of the input 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 probe_equal The binary function to compare set keys and probe keys for equality
+   * @param probe_hash The unary function to hash probe keys
+   * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+   * `output_match`
+   * @param output_match Beginning of the sequence of matching elements
+   * @param stream CUDA stream this operation is executed in
+   *
+   * @return Iterator pair indicating the the end of the output sequences
+   */
+  template <class InputProbeIt,
+            class StencilIt,
+            class Predicate,
+            class ProbeEqual,
+            class ProbeHash,
+            class OutputProbeIt,
+            class OutputMatchIt>
+  std::pair<OutputProbeIt, OutputMatchIt> retrieve_outer_if(
+    InputProbeIt first,
+    InputProbeIt last,
+    StencilIt stencil,
+    Predicate const& pred,
+    ProbeEqual const& probe_equal,
+    ProbeHash const& probe_hash,
+    OutputProbeIt output_probe,
+    OutputMatchIt output_match,
+    cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
+
   /**
    * @brief Retrieves all keys contained in the multiset
    *
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index dd146341e..3a639fba4 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -100,12 +100,15 @@ ConfigureTest(DYNAMIC_MAP_TEST
 # - static_multiset tests -------------------------------------------------------------------------
 ConfigureTest(STATIC_MULTISET_TEST
     static_multiset/contains_test.cu
+    static_multiset/count_if_test.cu
     static_multiset/count_test.cu
+    static_multiset/custom_count_if_test.cu
     static_multiset/custom_count_test.cu
     static_multiset/find_test.cu
     static_multiset/insert_test.cu
     static_multiset/for_each_test.cu
     static_multiset/retrieve_test.cu
+    static_multiset/retrieve_if_impl_test.cu
     static_multiset/retrieve_if_test.cu
     static_multiset/large_input_test.cu
     static_multiset/load_factor_test.cu
diff --git a/tests/static_multiset/count_if_test.cu b/tests/static_multiset/count_if_test.cu
new file mode 100644
index 000000000..9826e36d4
--- /dev/null
+++ b/tests/static_multiset/count_if_test.cu
@@ -0,0 +1,230 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <test_utils.hpp>
+
+#include <cuco/detail/__config>
+#include <cuco/static_multiset.cuh>
+
+#include <cuda/functional>
+#include <cuda/iterator>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+
+using size_type = int32_t;
+
+template <typename Set>
+void test_count_if(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<size_type>{0},
+    cuda::proclaim_return_type<Key>([] __device__(auto i) { return Key{i}; }));
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  set.clear();
+
+  set.insert(keys_begin, keys_begin + num_keys);
+
+  SECTION("Count_if with all elements selected should match count.")
+  {
+    auto const count =
+      set.count_if(keys_begin,
+                   keys_begin + num_keys,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }));
+
+    REQUIRE(count == num_keys);
+  }
+
+  SECTION("Count_if with no elements selected should return zero.")
+  {
+    auto const count =
+      set.count_if(keys_begin,
+                   keys_begin + num_keys,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }));
+
+    REQUIRE(count == 0);
+  }
+
+  SECTION("Count_if with alternating predicate should count selected keys.")
+  {
+    auto const count = set.count_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }));
+
+    REQUIRE(count == (num_keys + 1) / 2);
+  }
+}
+
+template <typename Set>
+void test_count_outer_if(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<size_type>{0},
+    cuda::proclaim_return_type<Key>([] __device__(auto i) { return Key{i}; }));
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  set.clear();
+
+  set.insert(keys_begin, keys_begin + num_keys);
+
+  SECTION("Count_outer_if with all elements selected should match count_outer.")
+  {
+    auto const count = set.count_outer_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }));
+
+    REQUIRE(count == num_keys);
+  }
+
+  SECTION("Count_outer_if with no elements selected should return one per input.")
+  {
+    auto const count = set.count_outer_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }));
+
+    REQUIRE(count == num_keys);
+  }
+
+  SECTION(
+    "Count_outer_if with alternating predicate should count selected matches and unselected rows.")
+  {
+    auto const count = set.count_outer_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }));
+
+    REQUIRE(count == num_keys);
+  }
+}
+
+template <typename Set>
+void test_count_if_stencil(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<size_type>{0},
+    cuda::proclaim_return_type<Key>([] __device__(auto i) { return Key{i}; }));
+
+  auto stencil_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<size_type>{0},
+    cuda::proclaim_return_type<size_type>([] __device__(auto i) { return i + 1000; }));
+
+  set.clear();
+  set.insert(keys_begin, keys_begin + num_keys);
+
+  SECTION("Count_if should apply the predicate to the stencil.")
+  {
+    auto const count = set.count_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type value) { return value < 1100; }));
+
+    REQUIRE(count == 100);
+  }
+
+  SECTION("Count_outer_if should apply the predicate to the stencil.")
+  {
+    auto const count = set.count_outer_if(
+      keys_begin,
+      keys_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type value) { return value < 1100; }));
+
+    REQUIRE(count == num_keys);
+  }
+}
+
+template <typename Set>
+void test_count_if_overloads(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<size_type>{0},
+    cuda::proclaim_return_type<Key>([] __device__(auto i) { return Key{i}; }));
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  set.clear();
+  set.insert(keys_begin, keys_begin + num_keys);
+
+  auto const pred =
+    cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; });
+
+  SECTION("Count_if explicit default key equality/hash matches overload.")
+  {
+    auto const count_default = set.count_if(keys_begin, keys_begin + num_keys, stencil_begin, pred);
+
+    auto const count_explicit = set.count_if(
+      keys_begin, keys_begin + num_keys, stencil_begin, pred, set.key_eq(), set.hash_function());
+
+    REQUIRE(count_explicit == count_default);
+  }
+
+  SECTION("Count_outer_if explicit default key equality/hash matches overload.")
+  {
+    auto const count_default =
+      set.count_outer_if(keys_begin, keys_begin + num_keys, stencil_begin, pred);
+
+    auto const count_explicit = set.count_outer_if(
+      keys_begin, keys_begin + num_keys, stencil_begin, pred, set.key_eq(), set.hash_function());
+
+    REQUIRE(count_explicit == count_default);
+  }
+}
+
+TEMPLATE_TEST_CASE_SIG(
+  "static_multiset count_if tests",
+  "",
+  ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 2)
+#if defined(CUCO_HAS_128BIT_ATOMICS)
+    ,
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 1),
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 2),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 1),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 2)
+#endif
+)
+{
+  constexpr size_type num_keys{666};
+
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<CGSize, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<CGSize, cuco::default_hash_function<Key>>>;
+
+  auto set =
+    cuco::static_multiset{num_keys, cuco::empty_key<Key>{-1}, {}, probe{}, {}, cuco::storage<2>{}};
+
+  test_count_if(set, num_keys);
+  test_count_outer_if(set, num_keys);
+  test_count_if_stencil(set, num_keys);
+  test_count_if_overloads(set, num_keys);
+}
diff --git a/tests/static_multiset/custom_count_if_test.cu b/tests/static_multiset/custom_count_if_test.cu
new file mode 100644
index 000000000..d5fb6b82b
--- /dev/null
+++ b/tests/static_multiset/custom_count_if_test.cu
@@ -0,0 +1,443 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <test_utils.hpp>
+
+#include <cuco/detail/__config>
+#include <cuco/static_multiset.cuh>
+
+#include <cuda/functional>
+#include <cuda/iterator>
+
+#include <catch2/catch_template_test_macros.hpp>
+
+using size_type = std::size_t;
+
+static auto constexpr XXX = 111;
+
+template <typename T>
+struct identity_hash {
+  __host__ __device__ identity_hash() {};
+  __host__ __device__ identity_hash([[maybe_unused]] int i) {}
+  __device__ T operator()(T k) const { return k; }
+};
+
+struct custom_hash {
+  __host__ __device__ custom_hash() {}
+  __host__ __device__ custom_hash([[maybe_unused]] int i) {}
+  template <typename custom_type>
+  __device__ custom_type operator()(custom_type k) const
+  {
+    return k / XXX;
+  };
+};
+
+struct custom_key_eq {
+  template <typename lhs_type, typename rhs_type>
+  __device__ bool operator()(lhs_type lhs, rhs_type rhs) const
+  {
+    return lhs / XXX == rhs;
+  }
+};
+
+template <typename Set>
+void test_count_if_duplicates(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto constexpr multiplicity = 3;
+
+  auto duplicate_keys_begin =
+    cuda::make_transform_iterator(cuda::counting_iterator<size_type>{0},
+                                  cuda::proclaim_return_type<Key>([] __device__(size_type i) {
+                                    return static_cast<Key>(i / multiplicity);
+                                  }));
+
+  set.clear();
+  set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys);
+
+  auto query_begin      = cuda::counting_iterator<size_type>{0};
+  auto const query_size = num_keys / multiplicity;
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  SECTION("Count_if with duplicates and all keys selected returns total multiplicity.")
+  {
+    auto const count =
+      set.count_if(query_begin,
+                   query_begin + query_size,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }));
+
+    REQUIRE(count == query_size * multiplicity);
+  }
+
+  SECTION("Count_if with duplicates and no keys selected returns zero.")
+  {
+    auto const count =
+      set.count_if(query_begin,
+                   query_begin + query_size,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }));
+
+    REQUIRE(count == 0);
+  }
+
+  SECTION("Count_if with duplicates counts only selected keys.")
+  {
+    auto const count = set.count_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }));
+
+    auto const expected = ((query_size + 1) / 2) * multiplicity;
+
+    REQUIRE(count == expected);
+  }
+
+  SECTION("Count_if with duplicates counts a single selected key by its multiplicity.")
+  {
+    auto const count =
+      set.count_if(query_begin,
+                   query_begin + query_size,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type i) { return i == 0; }));
+
+    REQUIRE(count == multiplicity);
+  }
+}
+
+template <typename Set>
+void test_count_outer_if_duplicates(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto constexpr multiplicity = 3;
+
+  auto duplicate_keys_begin =
+    cuda::make_transform_iterator(cuda::counting_iterator<size_type>{0},
+                                  cuda::proclaim_return_type<Key>([] __device__(size_type i) {
+                                    return static_cast<Key>(i / multiplicity);
+                                  }));
+
+  set.clear();
+  set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys);
+
+  // Query each unique key once.
+  auto query_begin      = cuda::counting_iterator<size_type>{0};
+  auto const query_size = num_keys / multiplicity;
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  SECTION("Count_outer_if with duplicates and all keys selected returns total multiplicity.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }));
+
+    REQUIRE(count == query_size * multiplicity);
+  }
+
+  SECTION("Count_outer_if with duplicates and no keys selected returns one per query.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }));
+
+    REQUIRE(count == query_size);
+  }
+
+  SECTION(
+    "Count_outer_if with duplicates counts selected matches and one for each unselected query.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }));
+
+    auto const selected_count   = (query_size + 1) / 2;
+    auto const unselected_count = query_size / 2;
+
+    auto const expected = selected_count * multiplicity + unselected_count;
+
+    REQUIRE(count == expected);
+  }
+
+  SECTION("Count_outer_if with a selected key counts its multiplicity.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return i == 0; }));
+
+    auto const expected = multiplicity + (query_size - 1);
+
+    REQUIRE(count == expected);
+  }
+}
+
+template <typename Set>
+void test_custom_count_if(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto const hash = []() {
+    if constexpr (cuco::is_double_hashing<typename Set::probing_scheme_type>::value) {
+      return cuda::std::tuple{custom_hash{}, custom_hash{}};
+    } else {
+      return custom_hash{};
+    }
+  }();
+
+  constexpr auto multiplicity = 3;
+
+  auto duplicate_keys_begin =
+    cuda::make_transform_iterator(cuda::counting_iterator<size_type>{0},
+                                  cuda::proclaim_return_type<Key>([] __device__(size_type i) {
+                                    return static_cast<Key>(i / multiplicity);
+                                  }));
+
+  set.clear();
+  set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys);
+
+  auto query_begin      = cuda::counting_iterator<size_type>{0};
+  auto const query_size = num_keys / multiplicity;
+  auto stencil_begin    = cuda::counting_iterator<size_type>{0};
+
+  SECTION("Count_if custom key equality/hash overload counts selected duplicates.")
+  {
+    auto const count = set.count_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    auto const selected_count = (query_size + 1) / 2;
+    REQUIRE(count == selected_count * multiplicity);
+  }
+
+  SECTION("Count_outer_if custom key equality/hash overload counts selected duplicates.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    auto const selected_count   = (query_size + 1) / 2;
+    auto const unselected_count = query_size / 2;
+
+    REQUIRE(count == selected_count * multiplicity + unselected_count);
+  }
+}
+
+template <typename Set>
+void test_custom_count_if_overloads(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto const hash = []() {
+    if constexpr (cuco::is_double_hashing<typename Set::probing_scheme_type>::value) {
+      return cuda::std::tuple{custom_hash{}, custom_hash{}};
+    } else {
+      return custom_hash{};
+    }
+  }();
+
+  constexpr auto multiplicity = 3;
+
+  auto duplicate_keys_begin =
+    cuda::make_transform_iterator(cuda::counting_iterator<size_type>{0},
+                                  cuda::proclaim_return_type<Key>([] __device__(size_type i) {
+                                    return static_cast<Key>(i / multiplicity);
+                                  }));
+
+  set.clear();
+  set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys);
+
+  auto query_begin   = cuda::counting_iterator<size_type>{0};
+  auto query_size    = num_keys / multiplicity;
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  SECTION("Count_if explicit key equality/hash overload selects all duplicates.")
+  {
+    auto const count =
+      set.count_if(query_begin,
+                   query_begin + query_size,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }),
+                   custom_key_eq{},
+                   hash);
+
+    REQUIRE(count == query_size * multiplicity);
+  }
+
+  SECTION("Count_if explicit key equality/hash overload selects no duplicates.")
+  {
+    auto const count =
+      set.count_if(query_begin,
+                   query_begin + query_size,
+                   stencil_begin,
+                   cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }),
+                   custom_key_eq{},
+                   hash);
+
+    REQUIRE(count == 0);
+  }
+
+  SECTION("Count_if explicit key equality/hash overload selects alternating duplicates.")
+  {
+    auto const count = set.count_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    auto const selected_count = (query_size + 1) / 2;
+    REQUIRE(count == selected_count * multiplicity);
+  }
+
+  SECTION("Count_outer_if explicit key equality/hash overload selects all duplicates.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return true; }),
+      custom_key_eq{},
+      hash);
+
+    REQUIRE(count == query_size * multiplicity);
+  }
+
+  SECTION("Count_outer_if explicit key equality/hash overload selects no duplicates.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type) { return false; }),
+      custom_key_eq{},
+      hash);
+
+    REQUIRE(count == query_size);
+  }
+
+  SECTION("Count_outer_if explicit key equality/hash overload selects alternating duplicates.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + query_size,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    auto const selected_count   = (query_size + 1) / 2;
+    auto const unselected_count = query_size / 2;
+
+    REQUIRE(count == selected_count * multiplicity + unselected_count);
+  }
+}
+
+template <typename Set>
+void test_custom_hash_count_if(Set& set, size_type num_keys)
+{
+  using Key = typename Set::key_type;
+
+  auto const hash = []() {
+    if constexpr (cuco::is_double_hashing<typename Set::probing_scheme_type>::value) {
+      return cuda::std::tuple{custom_hash{}, custom_hash{}};
+    } else {
+      return custom_hash{};
+    }
+  }();
+
+  auto const iter = cuda::counting_iterator<Key>{0};
+  set.clear();
+  set.insert(iter, iter + num_keys);
+
+  auto query_begin = cuda::make_transform_iterator(
+    cuda::make_counting_iterator<size_type>(0),
+    cuda::proclaim_return_type<Key>([] __device__(auto i) { return static_cast<Key>(i * XXX); }));
+
+  auto stencil_begin = cuda::counting_iterator<size_type>{0};
+
+  SECTION("Count_if uses custom key equality and hash.")
+  {
+    auto const count = set.count_if(
+      query_begin,
+      query_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    REQUIRE(count == (num_keys + 1) / 2);
+  }
+
+  SECTION("Count_outer_if uses custom key equality and hash.")
+  {
+    auto const count = set.count_outer_if(
+      query_begin,
+      query_begin + num_keys,
+      stencil_begin,
+      cuda::proclaim_return_type<bool>([] __device__(size_type i) { return (i % 2) == 0; }),
+      custom_key_eq{},
+      hash);
+
+    REQUIRE(count == num_keys);
+  }
+}
+
+TEMPLATE_TEST_CASE_SIG(
+  "static_multiset custom count_if tests",
+  "",
+  ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 2)
+#if defined(CUCO_HAS_128BIT_ATOMICS)
+    ,
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 1),
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 2),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 1),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 2)
+#endif
+)
+{
+  constexpr size_type num_keys{555};
+
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<CGSize, identity_hash<Key>>,
+                                   cuco::double_hashing<CGSize, identity_hash<Key>>>;
+
+  auto set =
+    cuco::static_multiset{num_keys, cuco::empty_key<Key>{-1}, {}, probe{}, {}, cuco::storage<2>{}};
+
+  test_count_if_duplicates(set, num_keys);
+  test_count_outer_if_duplicates(set, num_keys);
+  test_custom_count_if(set, num_keys);
+  test_custom_count_if_overloads(set, num_keys);
+  test_custom_hash_count_if(set, num_keys);
+}
diff --git a/tests/static_multiset/retrieve_if_impl_test.cu b/tests/static_multiset/retrieve_if_impl_test.cu
new file mode 100644
index 000000000..fcb799e0a
--- /dev/null
+++ b/tests/static_multiset/retrieve_if_impl_test.cu
@@ -0,0 +1,757 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <test_utils.hpp>
+
+#include <cuco/detail/__config>
+#include <cuco/static_multiset.cuh>
+
+#include <cuda/functional>
+#include <cuda/iterator>
+#include <thrust/device_vector.h>
+#include <thrust/distance.h>
+#include <thrust/functional.h>
+#include <thrust/sequence.h>
+#include <thrust/sort.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+
+#include <limits>
+
+template <class Container>
+void test_multiplicity(Container& container, std::size_t num_keys, std::size_t multiplicity)
+{
+  using key_type = typename Container::key_type;
+
+  container.clear();
+
+  auto const num_unique_keys = num_keys / multiplicity;
+  REQUIRE(num_unique_keys > 0);
+
+  auto const num_actual_keys = num_unique_keys * multiplicity;
+  REQUIRE(num_actual_keys <= num_keys);
+
+  thrust::device_vector<key_type> probed_keys(num_actual_keys);
+  thrust::device_vector<key_type> matched_keys(num_actual_keys);
+
+  auto const keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<key_type>(0),
+    cuda::proclaim_return_type<key_type>([multiplicity] __device__(auto const& i) {
+      return static_cast<key_type>(i / multiplicity);
+    }));
+
+  container.insert(keys_begin, keys_begin + num_actual_keys);
+  REQUIRE(container.size() == num_actual_keys);
+
+  SECTION("All inserted keys should be contained.")
+  {
+    auto const [probed_end, matched_end] = container.retrieve(
+      keys_begin, keys_begin + num_actual_keys, probed_keys.begin(), matched_keys.begin());
+
+    thrust::sort(probed_keys.begin(), probed_end);
+    thrust::sort(matched_keys.begin(), matched_end);
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_keys.end(), keys_begin, cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::equal(
+      matched_keys.begin(), matched_keys.end(), keys_begin, cuda::std::equal_to<key_type>{}));
+  }
+}
+
+template <class Container>
+void test_outer(Container& container, std::size_t num_keys)
+{
+  using key_type                = typename Container::key_type;
+  auto const empty_key_sentinel = container.empty_key_sentinel();
+
+  container.clear();
+
+  auto const keys_begin = cuda::counting_iterator<key_type>{0};
+  auto const query_size = num_keys * 2ull;
+
+  thrust::device_vector<key_type> probed_keys(query_size);
+  thrust::device_vector<key_type> matched_keys(query_size);
+
+  SECTION("Non-inserted keys should return the empty-key sentinel.")
+  {
+    auto const [probed_end, matched_end] = container.retrieve_outer(keys_begin,
+                                                                    keys_begin + query_size,
+                                                                    container.key_eq(),
+                                                                    container.hash_function(),
+                                                                    probed_keys.begin(),
+                                                                    matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == query_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            query_size);
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_end, keys_begin, cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::all_of(
+      matched_keys.begin(),
+      matched_end,
+      cuda::proclaim_return_type<bool>([empty_key_sentinel] __device__(auto const& k) {
+        return static_cast<bool>(k == static_cast<key_type>(empty_key_sentinel));
+      })));
+  }
+
+  container.insert(keys_begin, keys_begin + num_keys);
+
+  SECTION("All inserted keys should be contained and missing keys should return the sentinel.")
+  {
+    auto const [probed_end, matched_end] = container.retrieve_outer(keys_begin,
+                                                                    keys_begin + query_size,
+                                                                    container.key_eq(),
+                                                                    container.hash_function(),
+                                                                    probed_keys.begin(),
+                                                                    matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == query_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            query_size);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_keys.end(), keys_begin, cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::equal(matched_keys.begin(),
+                              matched_keys.begin() + num_keys,
+                              keys_begin,
+                              cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::all_of(
+      matched_keys.begin() + num_keys,
+      matched_keys.end(),
+      cuda::proclaim_return_type<bool>([empty_key_sentinel] __device__(auto const& k) {
+        return static_cast<bool>(k == static_cast<key_type>(empty_key_sentinel));
+      })));
+  }
+}
+
+template <class Container>
+void test_retrieve_if(Container& container, std::size_t num_keys)
+{
+  using key_type = typename Container::key_type;
+
+  container.clear();
+
+  auto const keys_begin = cuda::counting_iterator<key_type>{0};
+
+  container.insert(keys_begin, keys_begin + num_keys);
+
+  thrust::device_vector<key_type> probed_keys(num_keys);
+  thrust::device_vector<key_type> matched_keys(num_keys);
+  thrust::device_vector<key_type> stencil(num_keys);
+
+  SECTION("retrieve_if should predicate on the stencil, not the probe.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{1});
+
+    auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const num_results =
+      static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+    auto const expected_size = num_keys / 2;
+
+    REQUIRE(num_results == expected_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            expected_size);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t i = 0; i < expected_size; ++i) {
+      auto const expected = static_cast<key_type>(i * 2 + 1);
+
+      REQUIRE(probed_keys[i] == expected);
+      REQUIRE(matched_keys[i] == expected);
+    }
+  }
+
+  SECTION("retrieve_if should retrieve only elements satisfying the predicate.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+    auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const num_results =
+      static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+    REQUIRE(num_results == (num_keys + 1) / 2);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            num_results);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t i = 0; i < num_results; ++i) {
+      auto const expected = static_cast<key_type>(i * 2);
+
+      REQUIRE(probed_keys[i] == expected);
+      REQUIRE(matched_keys[i] == expected);
+    }
+  }
+
+  SECTION("retrieve_if should return nothing when the predicate is always false.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+    auto const pred = [] __device__(key_type) { return false; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == 0);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) == 0);
+  }
+
+  SECTION("retrieve_if should retrieve everything when the predicate is always true.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+    auto const pred = [] __device__(key_type) { return true; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == num_keys);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) == num_keys);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_end, keys_begin, cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::equal(
+      matched_keys.begin(), matched_end, keys_begin, cuda::std::equal_to<key_type>{}));
+  }
+}
+
+template <class Container>
+void test_retrieve_if_stencil(Container& container, std::size_t num_keys)
+{
+  using key_type = typename Container::key_type;
+
+  container.clear();
+
+  auto const keys_begin = cuda::counting_iterator<key_type>{0};
+
+  container.insert(keys_begin, keys_begin + num_keys);
+
+  thrust::device_vector<key_type> probes{2, 1, 4, 3};
+  thrust::device_vector<key_type> stencil{1, 2, 3, 4};
+  thrust::device_vector<key_type> probed_keys(4);
+  thrust::device_vector<key_type> matched_keys(4);
+
+  SECTION("retrieve_if should predicate on the stencil, not the probe.")
+  {
+    auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(probes.begin(),
+                                                                 probes.end(),
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const num_results =
+      static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+    // stencil: [1, 2, 3, 4]
+    // probe:   [2, 1, 4, 3]
+    //
+    // Only stencil values 2 and 4 satisfy the predicate, corresponding
+    // to probes 1 and 3.
+    REQUIRE(num_results == 2);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) == 2);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    REQUIRE(probed_keys[0] == key_type{1});
+    REQUIRE(probed_keys[1] == key_type{3});
+    REQUIRE(matched_keys[0] == key_type{1});
+    REQUIRE(matched_keys[1] == key_type{3});
+  }
+}
+
+template <class Container>
+void test_retrieve_if_with_probe(Container& container, std::size_t num_keys)
+{
+  using key_type = typename Container::key_type;
+
+  container.clear();
+
+  auto const keys_begin = cuda::counting_iterator<key_type>{0};
+
+  container.insert(keys_begin, keys_begin + num_keys);
+
+  thrust::device_vector<key_type> probed_keys(num_keys);
+  thrust::device_vector<key_type> matched_keys(num_keys);
+  thrust::device_vector<key_type> stencil(num_keys);
+
+  thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+  SECTION("retrieve_if should accept explicit equality and hash functions.")
+  {
+    auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 container.key_eq(),
+                                                                 container.hash_function(),
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const num_results =
+      static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+    REQUIRE(num_results == (num_keys + 1) / 2);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            num_results);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t i = 0; i < num_results; ++i) {
+      auto const expected = static_cast<key_type>(i * 2);
+
+      REQUIRE(probed_keys[i] == expected);
+      REQUIRE(matched_keys[i] == expected);
+    }
+  }
+
+  SECTION("retrieve_if with explicit equality and hash should return nothing for false predicate.")
+  {
+    auto const pred = [] __device__(key_type) { return false; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 container.key_eq(),
+                                                                 container.hash_function(),
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == 0);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) == 0);
+  }
+}
+
+template <class Container>
+void test_retrieve_if_multiplicity(Container& container, std::size_t num_keys)
+{
+  using key_type = typename Container::key_type;
+
+  constexpr std::size_t multiplicity = 2;
+
+  container.clear();
+
+  auto const num_unique_keys = num_keys / multiplicity;
+  auto const num_actual_keys = num_unique_keys * multiplicity;
+
+  auto const keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<key_type>(0),
+    cuda::proclaim_return_type<key_type>([multiplicity] __device__(auto const& i) {
+      return static_cast<key_type>(i / multiplicity);
+    }));
+
+  container.insert(keys_begin, keys_begin + num_actual_keys);
+  REQUIRE(container.size() == num_actual_keys);
+
+  thrust::device_vector<key_type> stencil(num_actual_keys);
+
+  thrust::device_vector<key_type> probed_keys(num_actual_keys * multiplicity);
+  thrust::device_vector<key_type> matched_keys(num_actual_keys * multiplicity);
+
+  thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+  SECTION("retrieve_if should filter duplicate matches using the stencil predicate.")
+  {
+    auto const pred = [] __device__(key_type value) { return value % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_actual_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const num_results =
+      static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+    auto const expected_results = (num_actual_keys / 2) * multiplicity;
+
+    REQUIRE(num_results == expected_results);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            expected_results);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t key = 0; key < num_unique_keys; ++key) {
+      auto const expected_key  = static_cast<key_type>(key);
+      auto const output_offset = key * multiplicity;
+
+      for (std::size_t j = 0; j < multiplicity; ++j) {
+        REQUIRE(probed_keys[output_offset + j] == expected_key);
+        REQUIRE(matched_keys[output_offset + j] == expected_key);
+      }
+    }
+  }
+
+  SECTION("retrieve_if should return nothing when the predicate is always false.")
+  {
+    auto const pred = [] __device__(key_type) { return false; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_actual_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == 0);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) == 0);
+  }
+
+  SECTION("retrieve_if should return all matches when the predicate is always true.")
+  {
+    auto const pred = [] __device__(key_type) { return true; };
+
+    auto const [probed_end, matched_end] = container.retrieve_if(keys_begin,
+                                                                 keys_begin + num_actual_keys,
+                                                                 stencil.begin(),
+                                                                 pred,
+                                                                 probed_keys.begin(),
+                                                                 matched_keys.begin());
+
+    auto const expected_results = num_actual_keys * multiplicity;
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) ==
+            expected_results);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            expected_results);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t key = 0; key < num_unique_keys; ++key) {
+      auto const expected_key   = static_cast<key_type>(key);
+      auto const expected_count = multiplicity * multiplicity;
+
+      for (std::size_t j = 0; j < expected_count; ++j) {
+        auto const output_index = key * expected_count + j;
+
+        REQUIRE(probed_keys[output_index] == expected_key);
+        REQUIRE(matched_keys[output_index] == expected_key);
+      }
+    }
+  }
+}
+
+template <class Container>
+void test_retrieve_outer_if(Container& container, std::size_t num_keys)
+{
+  using key_type                = typename Container::key_type;
+  auto const empty_key_sentinel = container.empty_key_sentinel();
+
+  container.clear();
+
+  auto const keys_begin = cuda::counting_iterator<key_type>{0};
+  auto const query_size = num_keys * 2ull;
+
+  container.insert(keys_begin, keys_begin + num_keys);
+
+  thrust::device_vector<key_type> probes(query_size);
+  thrust::device_vector<key_type> stencil(query_size);
+  thrust::device_vector<key_type> probed_keys(query_size);
+  thrust::device_vector<key_type> matched_keys(query_size);
+
+  thrust::sequence(probes.begin(), probes.end(), key_type{0});
+
+  SECTION("retrieve_outer_if should return matches and sentinels for misses.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+    auto const pred = [] __device__(key_type) { return true; };
+
+    auto const [probed_end, matched_end] = container.retrieve_outer_if(probes.begin(),
+                                                                       probes.end(),
+                                                                       stencil.begin(),
+                                                                       pred,
+                                                                       container.key_eq(),
+                                                                       container.hash_function(),
+                                                                       probed_keys.begin(),
+                                                                       matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == query_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            query_size);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_keys.end(), probes.begin(), cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::equal(matched_keys.begin(),
+                              matched_keys.begin() + num_keys,
+                              keys_begin,
+                              cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::all_of(
+      matched_keys.begin() + num_keys,
+      matched_keys.end(),
+      cuda::proclaim_return_type<bool>([empty_key_sentinel] __device__(auto const& k) {
+        return static_cast<bool>(k == static_cast<key_type>(empty_key_sentinel));
+      })));
+  }
+
+  SECTION("retrieve_outer_if should predicate on the stencil, not the probe.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{1});
+
+    auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+    auto const [probed_end, matched_end] = container.retrieve_outer_if(probes.begin(),
+                                                                       probes.end(),
+                                                                       stencil.begin(),
+                                                                       pred,
+                                                                       container.key_eq(),
+                                                                       container.hash_function(),
+                                                                       probed_keys.begin(),
+                                                                       matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == query_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            query_size);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+    for (std::size_t i = 0; i < query_size; ++i) {
+      auto const expected_probe = static_cast<key_type>(i);
+
+      REQUIRE(probed_keys[i] == expected_probe);
+
+      if (i % 2 == 0) {
+        // stencil[i] = i + 1 is odd -> predicate false
+        REQUIRE(matched_keys[i] == static_cast<key_type>(empty_key_sentinel));
+      } else {
+        // stencil[i] = i + 1 is even -> predicate true
+        if (i < num_keys) {
+          REQUIRE(matched_keys[i] == expected_probe);
+        } else {
+          REQUIRE(matched_keys[i] == static_cast<key_type>(empty_key_sentinel));
+        }
+      }
+    }
+  }
+
+  SECTION(
+    "retrieve_outer_if should return every probe with the sentinel for an always-false predicate.")
+  {
+    thrust::sequence(stencil.begin(), stencil.end(), key_type{0});
+
+    auto const pred = [] __device__(key_type) { return false; };
+
+    auto const [probed_end, matched_end] = container.retrieve_outer_if(probes.begin(),
+                                                                       probes.end(),
+                                                                       stencil.begin(),
+                                                                       pred,
+                                                                       container.key_eq(),
+                                                                       container.hash_function(),
+                                                                       probed_keys.begin(),
+                                                                       matched_keys.begin());
+
+    REQUIRE(static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end)) == query_size);
+    REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+            query_size);
+
+    thrust::sort_by_key(
+      probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>{});
+
+    REQUIRE(cuco::test::equal(
+      probed_keys.begin(), probed_end, probes.begin(), cuda::std::equal_to<key_type>{}));
+
+    REQUIRE(cuco::test::all_of(
+      matched_keys.begin(),
+      matched_end,
+      cuda::proclaim_return_type<bool>([empty_key_sentinel] __device__(auto const& k) {
+        return k == static_cast<key_type>(empty_key_sentinel);
+      })));
+  }
+}
+
+template <class Container>
+void test_retrieve_outer_if_multiplicity(Container& container, std::size_t num_keys)
+{
+  using key_type                = typename Container::key_type;
+  auto const empty_key_sentinel = container.empty_key_sentinel();
+
+  constexpr std::size_t multiplicity = 2;
+
+  container.clear();
+
+  auto const num_unique_keys = num_keys / multiplicity;
+  auto const num_actual_keys = num_unique_keys * multiplicity;
+
+  auto const keys_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator<key_type>(0),
+    cuda::proclaim_return_type<key_type>([multiplicity] __device__(auto const& i) {
+      return static_cast<key_type>(i / multiplicity);
+    }));
+
+  container.insert(keys_begin, keys_begin + num_actual_keys);
+
+  auto const query_size = num_unique_keys * 2ull;
+
+  thrust::device_vector<key_type> probes(query_size);
+  thrust::device_vector<key_type> stencil(query_size);
+
+  thrust::device_vector<key_type> probed_keys(query_size * multiplicity);
+  thrust::device_vector<key_type> matched_keys(query_size * multiplicity);
+
+  thrust::sequence(probes.begin(), probes.end(), key_type{0});
+  thrust::sequence(stencil.begin(), stencil.end(), key_type{1});
+
+  auto const pred = [] __device__(key_type key) { return key % 2 == 0; };
+
+  auto const [probed_end, matched_end] = container.retrieve_outer_if(probes.begin(),
+                                                                     probes.end(),
+                                                                     stencil.begin(),
+                                                                     pred,
+                                                                     container.key_eq(),
+                                                                     container.hash_function(),
+                                                                     probed_keys.begin(),
+                                                                     matched_keys.begin());
+
+  auto const num_predicate_false_probes = query_size / 2;
+  auto const num_predicate_true_probes  = query_size / 2;
+
+  auto const num_matching_unique_keys = (num_unique_keys > 1) ? (num_unique_keys / 2) : 0;
+
+  auto const num_missing_selected_probes = num_predicate_true_probes - num_matching_unique_keys;
+
+  auto const expected_results = num_predicate_false_probes +
+                                num_matching_unique_keys * multiplicity +
+                                num_missing_selected_probes;
+
+  auto const num_results = static_cast<std::size_t>(std::distance(probed_keys.begin(), probed_end));
+
+  REQUIRE(num_results == expected_results);
+  REQUIRE(static_cast<std::size_t>(std::distance(matched_keys.begin(), matched_end)) ==
+          expected_results);
+
+  thrust::sort_by_key(
+    probed_keys.begin(), probed_end, matched_keys.begin(), cuda::std::less<key_type>());
+
+  std::size_t output_index = 0;
+
+  for (std::size_t probe = 0; probe < query_size; ++probe) {
+    auto const expected_probe = static_cast<key_type>(probe);
+
+    if (probe % 2 == 0) {
+      // Predicate is false for even stencil values.
+      REQUIRE(probed_keys[output_index] == expected_probe);
+      REQUIRE(matched_keys[output_index] == static_cast<key_type>(empty_key_sentinel));
+      ++output_index;
+    } else {
+      // Predicate is true for odd probe positions.
+      if (probe < num_unique_keys) {
+        for (std::size_t j = 0; j < multiplicity; ++j) {
+          REQUIRE(probed_keys[output_index] == expected_probe);
+          REQUIRE(matched_keys[output_index] == expected_probe);
+          ++output_index;
+        }
+      } else {
+        REQUIRE(probed_keys[output_index] == expected_probe);
+        REQUIRE(matched_keys[output_index] == static_cast<key_type>(empty_key_sentinel));
+        ++output_index;
+      }
+    }
+  }
+
+  REQUIRE(output_index == expected_results);
+}
+
+TEMPLATE_TEST_CASE_SIG(
+  "static_multiset retrieve if tests",
+  "",
+  ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int32_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int64_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int32_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int64_t, cuco::test::probe_sequence::linear_probing, 2)
+#if defined(CUCO_HAS_128BIT_ATOMICS)
+    ,
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 1),
+  (__int128_t, cuco::test::probe_sequence::double_hashing, 2),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 1),
+  (__int128_t, cuco::test::probe_sequence::linear_probing, 2)
+#endif
+)
+{
+  constexpr std::size_t num_keys{400};
+  constexpr double desired_load_factor = 0.5;
+  constexpr auto empty_key_sentinel    = std::numeric_limits<Key>::max();
+
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<CGSize, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<CGSize, cuco::default_hash_function<Key>>>;
+
+  auto set = cuco::static_multiset{
+    num_keys, desired_load_factor, cuco::empty_key<Key>{empty_key_sentinel}, {}, probe{}};
+
+  test_multiplicity(set, num_keys, 1);
+  test_multiplicity(set, num_keys, 2);
+  test_multiplicity(set, num_keys, 11);
+
+  test_outer(set, num_keys);
+
+  test_retrieve_if(set, num_keys);
+  test_retrieve_if_with_probe(set, num_keys);
+  test_retrieve_if_stencil(set, num_keys);
+  test_retrieve_if_multiplicity(set, num_keys);
+
+  test_retrieve_outer_if(set, num_keys);
+  test_retrieve_outer_if_multiplicity(set, num_keys);
+}