Add accuracy test for Bloom filter

README.md

@@ -256,4 +256,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries.
 
 #### Examples:
-- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw))
+- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVSJhEg0SXfRViQLaauqVMjjMYyVwZ76AUFR_vteezxkCKRaLZGSYN_HueceX_sp0kxrLoWO-t-fIp5G_V4c5UQsLVmyqB9Rm5IojrS0irrvnXczAe_gShZbxZeZgQZtwmn39I8W_jqLYfxldD0awtXt5O52Mrwf3Y7bzsE7feKUCc1SsCJlCkzGYFgQin_CTgxfmHJo4LTdhYYzmEVhbxY1z32UrbSwIlsQ0oDVDMNwDQueM2CPlBUGuAAqV0XOiaAMNtxkPlWI4-HAtxBEJoagPUGPAr8t6pZAzA66-2TGFP1OZ7PZtImH3ZZq2clLY935NLq6GU9vWgh95_ZZ5EgvKPbTcoWFJ1sgBSKjJEG8OdmAVECWiuGekQ75RnHDxTIGLRdmQxTzcVKujeKJNXvkVTix_roB0kcEEjecwmg6i-DDcDqaxj7O19H9X7ef7-HrcDIZju9HN1O4nWCzxtcj1yr89hGG42_w92h8HQND6jAVeyyUqwKhckcrS0sOp4ztwVjIEpYuGOULTqGSESzlmimBZUHB1IqXgkOQqY-T8xU3xPi1g-J8qs5MzMRvXNDcpgwuqKWyk-RSrubYd8NUm9rsct_GZMpq06HSCtN2mwdbKVtjivmaUSPVcRP2yKh1wOaFxKZtj1tp7C5DqbVfY-ASm8LIyi9zYVBxXDTWkqfNmXjCwsAtUqzbOI5B2NX8gW21E9sAet3fu93uOew-nU7nAv5kgiliWNgGZ388kilKt8FL3HfQbZ-dh0gjZFcZz_WCK20gI_nCx3PBZNjw9L6RQBwkaIXM54ce2iahWxo9Tn1lHgf-6zZb1SZKBHxzQ3ZHnjN10MBsi1Jo-xaAk8DJC03m3mTgsp_vXKcF2bhTXvcqJ0PKFsTmBsoGO02-BuSrd5Lr9-uau6hyXYZ4T7UKn0NqbdJ-HzVo4OICT-QHmz8gMM-75_htPAssh6lCORKXZc_dXCxh9jHWTEDtg-F9LibSPOQu1dnv7-m8hto1rFF1zs_WnUul6IbbaidsibJtxt6jjRnc_71mSEOs00ox91bIe93lvG6Ajl4sO9uTulhKK7ELUzrsbVUBXmC8NPi4mMM89w7OMMiZpGmjQhGHTM1fspZImV86SxyE2J1GCTyGBck12-fuqKOoO-6dlFqMUMo_lqlt7fTt5ipO4TWXVufboCGc9bvS0PHeXYQ6kzZPoUwH_mozyrJWITVeLGsGfnggMfd3k0GvzgoeVXcX6gNq4pfCq8b-F7i-TJMRA3iL-dvaXSNMeFSvhos7otzd1vUDob1joWRCcrze8EJJiSGoc2WpsRgrroUJUdhjxhNu0LXk9VXdH-8ml108YwWW5eaFLJFgKxJEjqx4X17VVqvKDYZ904zojGn3BEn9zMJ6j9MpdnSKQKd4g85FLpEux7YDOyiPuF9sVPLyF1pjX2xHGlRfCifWyaDZhE4IWMqv1G6Zd_E_8r6upL70i7yieXxAoii_rxjR2N30xwDfLrha0eENXAcPLALwg4kYzPH6wcNYMy8bVC3PvVDmlVB8qIqdY3NVMRSfgC5-fcaXsXtq4uNTvTygI7GmtHd6Znu4LQtTvq6jFgYa0JOT3ntoEUWzgV7N33eh1UIgBn8ZzMHSVk5WiX9y5zypxaSU5ri4Lt_HuIC0iofoOa728YLa28cWRc8__M-_9KAQpw==))

benchmarks/bloom_filter/utils.hpp

@@ -62,13 +62,24 @@ void add_fpr_summary(nvbench::state& state, FilterType& filter)
   filter.add(tp_begin, tp_end);
   filter.contains(tn_begin, tn_end, result.begin());
 
-  float fp = thrust::count(thrust::device, result.begin(), result.end(), true);
+  auto const fp_emp =
+    static_cast<double>(thrust::count(thrust::device, result.begin(), result.end(), true)) /
+    static_cast<double>(num_keys);
 
-  auto& summ = state.add_summary("FalsePositiveRate");
-  summ.set_string("hint", "FPR");
-  summ.set_string("short_name", "FPR");
-  summ.set_string("description", "False-positive rate of the bloom filter.");
-  summ.set_float64("value", fp / num_keys);
+  auto& summ_fp = state.add_summary("FalsePositiveRate");
+  summ_fp.set_string("hint", "FPR");
+  summ_fp.set_string("short_name", "FPR");
+  summ_fp.set_string("description", "False-positive rate of the bloom filter.");
+  summ_fp.set_float64("value", fp_emp);
+
+  auto const fp_theo = filter.expected_false_positive_rate(num_keys);
+  auto const fp_dev  = fp_emp - fp_theo;
+
+  auto& summ_dev = state.add_summary("FalsePositiveRateDeviation");
+  summ_dev.set_string("hint", "FPRDev");
+  summ_dev.set_string("short_name", "FPRDev");
+  summ_dev.set_string("description", "Deviation of false-positive rate over theoretical value.");
+  summ_dev.set_float64("value", fp_dev);
 
   filter.clear();
 }

examples/bloom_filter/host_bulk_example.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -68,7 +68,8 @@ int main(void)
   float fp_rate =
     float(thrust::count(thrust::device, tn_result.begin(), tn_result.end(), true)) / float(num_tn);
 
-  std::cout << "TPR=" << tp_rate << " FPR=" << fp_rate << std::endl;
+  std::cout << "TPR[measured]=" << tp_rate << " FPR[measured]=" << fp_rate
+            << " FPR[expected]=" << filter.expected_false_positive_rate(num_tp) << std::endl;
 
   return 0;
 }

include/cuco/bloom_filter.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -312,6 +312,17 @@ class bloom_filter {
                                             OutputIt output_begin,
                                             cuda::stream_ref stream = {}) const noexcept;
 
+  /**
+   * @brief Computes the expected false-positive rate of a blocked Bloom filter
+   * using the Poisson-based formula (Eq. 3) from Putze et.al. "Cache-, Hash- and Space-Efficient
+   * Bloom Filters".
+   *
+   * @param num_items Number of inserted distinct elements
+   *
+   * @return Approximation of the expected false-positive rate
+   */
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items) const;
+
   /**
    * @brief Gets a pointer to the underlying filter storage.
    *

include/cuco/bloom_filter_policies.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -40,7 +40,7 @@ using arrow_filter_policy = detail::arrow_filter_policy<Key, XXHash64>;
  * @brief The default policy that defines how a Blocked Bloom Filter generates and stores a key's
  * fingerprint.
  *
- * @note `Word` type must be an atomically updatable integral type. `WordsPerBlock` must
+ * @note `Word` type must be an atomically updatable unsigned integral type. `WordsPerBlock` must
  * be a power-of-two.
  *
  * @tparam Hash Hash function used to generate a key's fingerprint
@@ -92,7 +92,7 @@ class default_filter_policy {
    *
    * @return The hash value of the key
    */
-  __device__ constexpr hash_result_type hash(hash_argument_type const& key) const;
+  [[nodiscard]] __device__ constexpr hash_result_type hash(hash_argument_type const& key) const;
 
   /**
    * @brief Determines the filter block a key is added into.
@@ -108,7 +108,8 @@ class default_filter_policy {
    * @return The block index for the given key's hash value
    */
   template <class Extent>
-  __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const;
+  [[nodiscard]] __device__ constexpr auto block_index(hash_result_type hash,
+                                                      Extent num_blocks) const;
 
   /**
    * @brief Determines the fingerprint pattern for a word/segment within the filter block for a
@@ -122,8 +123,21 @@ class default_filter_policy {
    *
    * @return The bit pattern for the word/segment in the filter block
    */
-  __device__ constexpr word_type word_pattern(hash_result_type hash,
-                                              std::uint32_t word_index) const;
+  [[nodiscard]] __device__ constexpr word_type word_pattern(hash_result_type hash,
+                                                            std::uint32_t word_index) const;
+
+  /**
+   * @brief Computes the expected false-positive rate of a blocked Bloom filter
+   * using the Poisson-based formula (Eq. 3) from Putze et.al. "Cache-, Hash- and Space-Efficient
+   * Bloom Filters".
+   *
+   * @param num_items Number of inserted distinct elements
+   * @param num_blocks The total number of blocks in the filter
+   *
+   * @return Approximation of the expected false-positive rate
+   */
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items,
+                                                             size_t num_blocks) const;
 
  private:
   impl_type impl_;  ///< Policy implementation

include/cuco/bloom_filter_ref.cuh

@@ -369,6 +369,17 @@ class bloom_filter_ref {
                                             OutputIt output_begin,
                                             cuda::stream_ref stream = {}) const noexcept;
 
+  /**
+   * @brief Computes the expected false-positive rate of a blocked Bloom filter
+   * using the Poisson-based formula (Eq. 3) from Putze et.al. "Cache-, Hash- and Space-Efficient
+   * Bloom Filters".
+   *
+   * @param num_items Number of inserted distinct elements
+   *
+   * @return Approximation of the expected false-positive rate
+   */
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items) const;
+
   /**
    * @brief Gets a pointer to the underlying filter storage.
    *

include/cuco/detail/bloom_filter/arrow_filter_policy.cuh

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cuco/detail/bloom_filter/utils.hpp>
 #include <cuco/hash_functions.cuh>
 
 #include <cuda/functional>
@@ -135,7 +136,8 @@ class arrow_filter_policy {
    * @return The block index for the given key's hash value
    */
   template <class Extent>
-  __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const
+  [[nodiscard]] __device__ constexpr auto block_index(hash_result_type hash,
+                                                      Extent num_blocks) const
   {
     constexpr auto hash_bits = cuda::std::numeric_limits<word_type>::digits;
     // TODO: assert if num_blocks > max_filter_blocks
@@ -153,7 +155,8 @@ class arrow_filter_policy {
    *
    * @return The bit pattern for the word/segment in the filter block
    */
-  __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const
+  [[nodiscard]] __device__ constexpr word_type word_pattern(hash_result_type hash,
+                                                            std::uint32_t word_index) const
   {
     word_type const key = static_cast<word_type>(hash);
     std::uint32_t salt;
@@ -182,6 +185,27 @@ class arrow_filter_policy {
     return word_type{1} << ((key * salt) >> 27);
   }
 
+  /**
+   * @brief Computes the expected false-positive rate of a blocked Bloom filter
+   * using the Poisson-based formula (Eq. 3) from Putze et.al. "Cache-, Hash- and Space-Efficient
+   * Bloom Filters".
+   *
+   * @param num_items Number of inserted distinct elements
+   * @param num_blocks The total number of blocks in the filter
+   *
+   * @return Approximation of the expected false-positive rate
+   */
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items,
+                                                             size_t num_blocks) const
+  {
+    return blocked_bloom_filter_expected_fpr(num_items,
+                                             num_blocks * words_per_block * sizeof(word_type),
+                                             cuda::std::numeric_limits<word_type>::digits,
+                                             words_per_block,
+                                             cuda::std::numeric_limits<hash_result_type>::digits,
+                                             words_per_block);
+  }
+
  private:
   hasher hash_;
 };

include/cuco/detail/bloom_filter/bloom_filter.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -129,6 +129,14 @@ __host__ constexpr void bloom_filter<Key, Extent, Scope, Policy, Allocator>::con
   ref_.contains_if_async(first, last, stencil, pred, output_begin, stream);
 }
 
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
+[[nodiscard]] __host__ double
+bloom_filter<Key, Extent, Scope, Policy, Allocator>::expected_false_positive_rate(
+  size_t num_items) const
+{
+  return ref_.expected_false_positive_rate(num_items);
+}
+
 template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
 [[nodiscard]] __host__ constexpr
   typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::word_type*

include/cuco/detail/bloom_filter/bloom_filter_impl.cuh

@@ -414,6 +414,11 @@ class bloom_filter_impl {
         first, num_keys, stencil, pred, output_begin, *this);
   }
 
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items) const
+  {
+    return policy_.expected_false_positive_rate(num_items, num_blocks_);
+  }
+
   [[nodiscard]] __host__ __device__ constexpr word_type* data() noexcept { return words_; }
 
   [[nodiscard]] __host__ __device__ constexpr word_type const* data() const noexcept
@@ -428,7 +433,6 @@ class bloom_filter_impl {
 
   // TODO
   // [[nodiscard]] __host__ double occupancy() const;
-  // [[nodiscard]] __host__ double expected_false_positive_rate(size_t unique_keys) const
   // [[nodiscard]] __host__ __device__ static uint32_t optimal_pattern_bits(size_t num_blocks)
   // template <typename CG, cuda::thread_scope NewScope = thread_scope>
   // [[nodiscard]] __device__ constexpr auto make_copy(CG const& group, word_type* const

include/cuco/detail/bloom_filter/bloom_filter_ref.inl

@@ -180,6 +180,13 @@ template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
   return impl_.data();
 }
 
+template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
+[[nodiscard]] __host__ double
+bloom_filter_ref<Key, Extent, Scope, Policy>::expected_false_positive_rate(size_t num_items) const
+{
+  return impl_.expected_false_positive_rate(num_items);
+}
+
 template <class Key, class Extent, cuda::thread_scope Scope, class Policy>
 [[nodiscard]] __host__ __device__ constexpr
   typename bloom_filter_ref<Key, Extent, Scope, Policy>::word_type const*

include/cuco/detail/bloom_filter/default_filter_policy.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -29,29 +29,40 @@ __host__
 }
 
 template <class Hash, class Word, uint32_t WordsPerBlock>
-__device__ constexpr typename default_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type
-default_filter_policy<Hash, Word, WordsPerBlock>::hash(
-  typename default_filter_policy<Hash, Word, WordsPerBlock>::hash_argument_type const& key) const
+[[nodiscard]] __device__ constexpr
+  typename default_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type
+  default_filter_policy<Hash, Word, WordsPerBlock>::hash(
+    typename default_filter_policy<Hash, Word, WordsPerBlock>::hash_argument_type const& key) const
 {
   return impl_.hash(key);
 }
 
 template <class Hash, class Word, uint32_t WordsPerBlock>
 template <class Extent>
-__device__ constexpr auto default_filter_policy<Hash, Word, WordsPerBlock>::block_index(
+[[nodiscard]] __device__ constexpr auto
+default_filter_policy<Hash, Word, WordsPerBlock>::block_index(
   typename default_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type hash,
   Extent num_blocks) const
 {
   return impl_.block_index(hash, num_blocks);
 }
 
 template <class Hash, class Word, uint32_t WordsPerBlock>
-__device__ constexpr typename default_filter_policy<Hash, Word, WordsPerBlock>::word_type
-default_filter_policy<Hash, Word, WordsPerBlock>::word_pattern(
-  default_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type hash,
-  std::uint32_t word_index) const
+[[nodiscard]] __device__ constexpr
+  typename default_filter_policy<Hash, Word, WordsPerBlock>::word_type
+  default_filter_policy<Hash, Word, WordsPerBlock>::word_pattern(
+    default_filter_policy<Hash, Word, WordsPerBlock>::hash_result_type hash,
+    std::uint32_t word_index) const
 {
   return impl_.word_pattern(hash, word_index);
 }
 
+template <class Hash, class Word, uint32_t WordsPerBlock>
+[[nodiscard]] __host__ double
+default_filter_policy<Hash, Word, WordsPerBlock>::expected_false_positive_rate(
+  size_t num_items, size_t num_blocks) const
+{
+  return impl_.expected_false_positive_rate(num_items, num_blocks);
+}
+
 }  // namespace cuco

include/cuco/detail/bloom_filter/default_filter_policy_impl.cuh

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cuco/detail/bloom_filter/utils.hpp>
 #include <cuco/detail/error.hpp>
 
 #include <cuda/std/bit>
@@ -36,6 +37,14 @@ class default_filter_policy_impl {
   using hash_argument_type = typename hasher::argument_type;
   using hash_result_type   = decltype(std::declval<hasher>()(std::declval<hash_argument_type>()));
 
+  static_assert(cuda::std::is_integral<word_type>::value &&
+                  cuda::std::is_unsigned<word_type>::value,
+                "Word type must be an unsigned integral type");
+
+  static_assert(cuda::std::is_integral<hash_result_type>::value &&
+                  cuda::std::is_unsigned<hash_result_type>::value,
+                "Hash result type must be an unsigned integral type");
+
   static constexpr std::uint32_t words_per_block = WordsPerBlock;
 
  private:
@@ -105,6 +114,17 @@ class default_filter_policy_impl {
     return word;
   }
 
+  [[nodiscard]] __host__ double expected_false_positive_rate(size_t num_items,
+                                                             size_t num_blocks) const
+  {
+    return blocked_bloom_filter_expected_fpr(num_items,
+                                             num_blocks * words_per_block * sizeof(word_type),
+                                             cuda::std::numeric_limits<word_type>::digits,
+                                             words_per_block,
+                                             cuda::std::numeric_limits<hash_result_type>::digits,
+                                             pattern_bits_);
+  }
+
  private:
   uint32_t pattern_bits_;
   uint32_t min_bits_per_word_;