Dispatch make_bucket_extent based on probing scheme type

include/cuco/detail/extent/extent.inl

@@ -132,20 +132,20 @@ template <int32_t CGSize, int32_t BucketSize, typename SizeType>
 template <typename ProbingScheme, typename Storage, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_bucket_extent(extent<SizeType, N> ext)
 {
-  return make_bucket_extent<ProbingScheme::cg_size, Storage::bucket_size, SizeType, N>(ext);
-  /*
-   // TODO
-if constexpr (cuco::is_double_hashing<ProbingScheme>::value) {
-} else {
-  auto const size = cuco::detail::int_div_ceil(
-    cuda::std::max(static_cast<SizeType>(ext), static_cast<SizeType>(1)),
-    ProbingScheme::cg_size * Storage::bucket_size) + cuda::std::min(static_cast<SizeType>(ext),
-static_cast<SizeType>(1)); if constexpr (N == dynamic_extent) { return bucket_extent<SizeType>{size
-* ProbingScheme::cg_size}; } else { return bucket_extent<SizeType, size * ProbingScheme::cg_size>{};
+  if constexpr (cuco::is_double_hashing<ProbingScheme>::value) {
+    return make_bucket_extent<ProbingScheme::cg_size, Storage::bucket_size, SizeType, N>(ext);
+  } else {
+    auto const size = cuco::detail::int_div_ceil(
+                        cuda::std::max(static_cast<SizeType>(ext), static_cast<SizeType>(1)),
+                        ProbingScheme::cg_size * Storage::bucket_size) +
+                      cuda::std::min(static_cast<SizeType>(ext), static_cast<SizeType>(1));
+    if constexpr (N == dynamic_extent) {
+      return bucket_extent<SizeType>{size * ProbingScheme::cg_size};
+    } else {
+      return bucket_extent<SizeType, size * ProbingScheme::cg_size>{};
+    }
   }
 }
-*/
-}
 
 template <typename ProbingScheme, typename Storage, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size)

tests/static_map/capacity_test.cu

@@ -159,7 +159,7 @@ TEST_CASE("Static map capacity", "")
 
   SECTION("Dynamic extent is evaluated at run time.")
   {
-    auto constexpr gold_capacity = 412;  // 103 x 2 x 2
+    auto constexpr gold_capacity = 404;  // 101 x 2 x 2
 
     using probe = cuco::linear_probing<2, cuco::default_hash_function<Key>>;
     auto map =

tests/static_map/erase_test.cu

@@ -119,8 +119,10 @@ TEMPLATE_TEST_CASE_SIG(
                               thrust::equal_to<Key>,
                               probe,
                               cuco::cuda_allocator<cuda::std::byte>,
-                              cuco::storage<2>>{
-    num_keys, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}, cuco::erased_key<Key>{-2}};
+                              cuco::storage<2>>{num_keys * 2,
+                                                cuco::empty_key<Key>{-1},
+                                                cuco::empty_value<Value>{-1},
+                                                cuco::erased_key<Key>{-2}};
 
   test_erase(map, num_keys);
 }

tests/static_map/unique_sequence_test.cu

@@ -166,24 +166,32 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
   constexpr size_type num_keys{400};
-  constexpr size_type gold_capacity = CGSize == 1 ? 422   // 211 x 1 x 2
-                                                  : 412;  // 103 x 2 x 2
+
+  using probe = std::conditional_t<
+    Probe == cuco::test::probe_sequence::linear_probing,
+    cuco::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
+    cuco::double_hashing<CGSize, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+
+  auto constexpr gold_capacity = [&]() {
+    if constexpr (cuco::is_double_hashing<probe>::value) {
+      return CGSize == 1 ? 422   // 211 x 1 x 2
+                         : 412;  // 103 x 2 x 2
+    } else {
+      return CGSize == 1 ? 402   // 201 x 1 x 2
+                         : 404;  // 101 x 2 x 2
+    }
+  }();
 
   // XXX: testing static extent is intended, DO NOT CHANGE
   using extent_type = cuco::extent<size_type, num_keys>;
-  using probe       = std::conditional_t<
-          Probe == cuco::test::probe_sequence::linear_probing,
-          cuco::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
-          cuco::double_hashing<CGSize, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
-
-  auto map = cuco::static_map<Key,
-                              Value,
-                              extent_type,
-                              cuda::thread_scope_device,
-                              thrust::equal_to<Key>,
-                              probe,
-                              cuco::cuda_allocator<cuda::std::byte>,
-                              cuco::storage<2>>{
+  auto map          = cuco::static_map<Key,
+                                       Value,
+                                       extent_type,
+                                       cuda::thread_scope_device,
+                                       thrust::equal_to<Key>,
+                                       probe,
+                                       cuco::cuda_allocator<cuda::std::byte>,
+                                       cuco::storage<2>>{
     extent_type{}, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
   REQUIRE(map.capacity() == gold_capacity);

tests/static_multiset/insert_test.cu

@@ -89,13 +89,21 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
   constexpr size_type num_keys{400};
-  constexpr size_type gold_capacity = CGSize == 1 ? 422   // 211 x 1 x 2
-                                                  : 412;  // 103 x 2 x 2
 
   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 constexpr gold_capacity = [&]() {
+    if constexpr (cuco::is_double_hashing<probe>::value) {
+      return CGSize == 1 ? 422   // 211 x 1 x 2
+                         : 412;  // 103 x 2 x 2
+    } else {
+      return CGSize == 1 ? 402   // 201 x 1 x 2
+                         : 404;  // 101 x 2 x 2
+    }
+  }();
+
   auto set =
     cuco::static_multiset{num_keys, cuco::empty_key<Key>{-1}, {}, probe{}, {}, cuco::storage<2>{}};
 

tests/static_set/capacity_test.cu

@@ -18,7 +18,7 @@
 
 #include <catch2/catch_test_macros.hpp>
 
-TEST_CASE("Static set capacity", "")
+TEST_CASE("static_set capacity tests", "")
 {
   using Key        = int32_t;
   using ProbeT     = cuco::double_hashing<1, cuco::default_hash_function<Key>>;
@@ -120,9 +120,9 @@ TEST_CASE("Static set capacity", "")
     REQUIRE(ref_capacity == gold_capacity);
   }
 
-  SECTION("Dynamic extent is evaluated at run time.")
+  SECTION("Dynamic extent for linear porbing.")
   {
-    auto constexpr gold_capacity = 412;  // 103 x 2 x 2
+    auto constexpr gold_capacity = 404;  // 101 x 2 x 2
 
     using probe = cuco::linear_probing<2, cuco::default_hash_function<Key>>;
     auto set    = cuco::static_set<Key,

tests/static_set/retrieve_all_test.cu

@@ -58,7 +58,7 @@ void test_unique_sequence(Set& set, std::size_t num_keys)
 }
 
 TEMPLATE_TEST_CASE_SIG(
-  "Retrieve all",
+  "static_set retrieve_all tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -72,14 +72,21 @@ TEMPLATE_TEST_CASE_SIG(
 {
   constexpr std::size_t num_keys{400};
   constexpr double desired_load_factor = 1.;
-  auto constexpr gold_capacity         = CGSize == 1 ? 409  // 409 x 1 x 1
-                                                     : 422  // 211 x 2 x 1
-    ;
 
   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 constexpr gold_capacity = [&]() {
+    if constexpr (cuco::is_double_hashing<probe>::value) {
+      return CGSize == 1 ? 409   // 409 x 1 x 1
+                         : 422;  // 211 x 2 x 1
+    } else {
+      return CGSize == 1 ? 401   // 401 x 1 x 1
+                         : 402;  // 201 x 2 x 1
+    }
+  }();
+
   auto set = cuco::static_set{num_keys, desired_load_factor, cuco::empty_key<Key>{-1}, {}, probe{}};
 
   REQUIRE(set.capacity() == gold_capacity);

tests/static_set/unique_sequence_test.cu

@@ -132,14 +132,21 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
   constexpr size_type num_keys{400};
-  constexpr size_type gold_capacity = CGSize == 1 ? 422  // 211 x 1 x 2
-                                                  : 412  // 103 x 2 x 2
-    ;
 
   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 constexpr gold_capacity = [&]() {
+    if constexpr (cuco::is_double_hashing<probe>::value) {
+      return CGSize == 1 ? 422   // 211 x 1 x 2
+                         : 412;  // 103 x 2 x 2
+    } else {
+      return CGSize == 1 ? 402   // 201 x 1 x 2
+                         : 404;  // 101 x 2 x 2
+    }
+  }();
+
   auto set =
     cuco::static_set{num_keys, cuco::empty_key<Key>{-1}, {}, probe{}, {}, cuco::storage<2>{}};