Resolve conflicts

NVIDIA · Oct 31, 2024 · aa0fe6a · aa0fe6a
1 parent f4b764c
commit aa0fe6a
Show file tree

Hide file tree

Showing 14 changed files with 69 additions and 99 deletions.
diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl
@@ -132,20 +132,24 @@ 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)
 {
-  if constexpr (cuco::is_double_hashing<ProbingScheme>::value) {
-    return make_bucket_extent<ProbingScheme::cg_size, Storage::bucket_size, SizeType, N>(ext);
+  return make_bucket_extent<ProbingScheme::cg_size, Storage::bucket_size, SizeType, N>(ext);
+  /*
+// TODO fix linear probing with exact capacity
+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 {
-    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>{};
-    }
+    return bucket_extent<SizeType, size * ProbingScheme::cg_size>{};
   }
 }
+*/
+}
 
 template <typename ProbingScheme, typename Storage, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size)

diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl
@@ -1384,7 +1384,7 @@ class operator_impl<
   using size_type  = typename base_type::size_type;
 
   static constexpr auto cg_size     = base_type::cg_size;
-  static constexpr auto window_size = base_type::window_size;
+  static constexpr auto bucket_size = base_type::bucket_size;
 
  public:
   /**

diff --git a/include/cuco/detail/static_set/static_set_ref.inl b/include/cuco/detail/static_set/static_set_ref.inl
@@ -644,7 +644,7 @@ class operator_impl<op::count_tag,
   using size_type  = typename base_type::size_type;
 
   static constexpr auto cg_size     = base_type::cg_size;
-  static constexpr auto window_size = base_type::window_size;
+  static constexpr auto bucket_size = base_type::bucket_size;
 
  public:
   /**

diff --git a/tests/static_map/capacity_test.cu b/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 = 404;  // 101 x 2 x 2
+    auto constexpr gold_capacity = 412;  // 103 x 2 x 2
 
     using probe = cuco::linear_probing<2, cuco::default_hash_function<Key>>;
     auto map =

diff --git a/tests/static_map/erase_test.cu b/tests/static_map/erase_test.cu
@@ -105,7 +105,7 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, int32_t, cuco::test::probe_sequence::linear_probing, 2),
   (int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
-  constexpr size_type num_keys{10'000};
+  constexpr size_type num_keys{1'000'000};
 
   using probe = std::conditional_t<
     Probe == cuco::test::probe_sequence::linear_probing,
@@ -119,10 +119,8 @@ TEMPLATE_TEST_CASE_SIG(
                               thrust::equal_to<Key>,
                               probe,
                               cuco::cuda_allocator<cuda::std::byte>,
-                              cuco::storage<2>>{num_keys * 2,
-                                                cuco::empty_key<Key>{-1},
-                                                cuco::empty_value<Value>{-1},
-                                                cuco::erased_key<Key>{-2}};
+                              cuco::storage<2>>{
+    num_keys, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}, cuco::erased_key<Key>{-2}};
 
   test_erase(map, num_keys);
 }
diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu
@@ -30,7 +30,7 @@
 
 #include <limits>
 
-template <std::size_t NumBuckets, typename Ref>
+template <std::size_t NumWindows, typename Ref>
 __global__ void shared_memory_test_kernel(Ref* maps,
                                           typename Ref::key_type const* const insterted_keys,
                                           typename Ref::mapped_type const* const inserted_values,
@@ -42,7 +42,7 @@ __global__ void shared_memory_test_kernel(Ref* maps,
   const size_t map_id = blockIdx.x;
   const size_t offset = map_id * number_of_elements;
 
-  __shared__ typename Ref::window_type sm_buffer[NumBuckets];
+  __shared__ typename Ref::window_type sm_buffer[NumWindows];
 
   auto g          = cuco::test::cg::this_thread_block();
   auto insert_ref = maps[map_id].make_copy(g, sm_buffer, cuco::thread_scope_block);
@@ -169,16 +169,16 @@ TEMPLATE_TEST_CASE_SIG("Shared memory static map",
 auto constexpr cg_size     = 1;
 auto constexpr window_size = 1;
 
-template <std::size_t NumBuckets>
+template <std::size_t NumWindows>
 __global__ void shared_memory_hash_table_kernel(bool* key_found)
 {
   using Key       = int32_t;
   using Value     = int32_t;
   using slot_type = cuco::pair<Key, Value>;
 
-  __shared__ cuco::window<slot_type, window_size> map[NumBuckets];
+  __shared__ cuco::window<slot_type, window_size> map[NumWindows];
 
-  using extent_type      = cuco::extent<std::size_t, NumBuckets>;
+  using extent_type      = cuco::extent<std::size_t, NumWindows>;
   using storage_ref_type = cuco::aow_storage_ref<slot_type, window_size, extent_type>;
 
   auto raw_ref =

diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu
@@ -170,32 +170,24 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
   constexpr size_type num_keys{400};
-
-  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
-    }
-  }();
+  constexpr size_type gold_capacity = CGSize == 1 ? 422   // 211 x 1 x 2
+                                                  : 412;  // 103 x 2 x 2
 
   // XXX: testing static extent is intended, DO NOT CHANGE
   using extent_type = cuco::extent<size_type, num_keys>;
-  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>>{
+  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>>{
     extent_type{}, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
   REQUIRE(map.capacity() == gold_capacity);

diff --git a/tests/static_multiset/insert_test.cu b/tests/static_multiset/insert_test.cu
@@ -89,21 +89,13 @@ 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>{}};
 

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

diff --git a/tests/static_set/retrieve_all_test.cu b/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(
-  "static_set retrieve_all tests",
+  "Retrieve all",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -72,21 +72,14 @@ 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);

diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
@@ -30,7 +30,7 @@
 
 #include <limits>
 
-template <std::size_t NumBuckets, typename Ref>
+template <std::size_t NumWindows, typename Ref>
 __global__ void shared_memory_test_kernel(Ref* sets,
                                           typename Ref::key_type const* const insterted_keys,
                                           size_t number_of_elements,
@@ -41,7 +41,7 @@ __global__ void shared_memory_test_kernel(Ref* sets,
   const size_t set_id = blockIdx.x;
   const size_t offset = set_id * number_of_elements;
 
-  __shared__ typename Ref::window_type sm_buffer[NumBuckets];
+  __shared__ typename Ref::window_type sm_buffer[NumWindows];
 
   auto g          = cuco::test::cg::this_thread_block();
   auto insert_ref = sets[set_id].make_copy(g, sm_buffer, cuco::thread_scope_block);
@@ -151,15 +151,15 @@ TEMPLATE_TEST_CASE_SIG("Shared memory static set", "", ((typename Key), Key), (i
 auto constexpr cg_size     = 1;
 auto constexpr window_size = 1;
 
-template <std::size_t NumBuckets>
+template <std::size_t NumWindows>
 __global__ void shared_memory_hash_set_kernel(bool* key_found)
 {
   using Key       = int32_t;
   using slot_type = Key;
 
-  __shared__ cuco::window<slot_type, window_size> set[NumBuckets];
+  __shared__ cuco::window<slot_type, window_size> set[NumWindows];
 
-  using extent_type      = cuco::extent<std::size_t, NumBuckets>;
+  using extent_type      = cuco::extent<std::size_t, NumWindows>;
   using storage_ref_type = cuco::aow_storage_ref<slot_type, window_size, extent_type>;
 
   auto raw_ref =

diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
@@ -132,21 +132,14 @@ 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>{}};
 

diff --git a/tests/utility/probing_scheme_test.cu b/tests/utility/probing_scheme_test.cu
@@ -79,7 +79,7 @@ __global__ void generate_cg_probing_sequence(Key key,
 TEMPLATE_TEST_CASE_SIG(
   "probing_scheme scalar vs CGSize 1 test",
   "",
-  ((typename Key, cuco::test::probe_sequence Probe, int32_t BucketSize), Key, Probe, BucketSize),
+  ((typename Key, cuco::test::probe_sequence Probe, int32_t WindowSize), Key, Probe, WindowSize),
   (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),
@@ -89,7 +89,7 @@ TEMPLATE_TEST_CASE_SIG(
   (int64_t, cuco::test::probe_sequence::linear_probing, 1),
   (int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
-  auto const upper_bound = cuco::make_window_extent<1, BucketSize>(cuco::extent<std::size_t>{10});
+  auto const upper_bound = cuco::make_window_extent<1, WindowSize>(cuco::extent<std::size_t>{10});
   constexpr size_t seq_length{8};
   constexpr Key key{42};
 

diff --git a/tests/utility/storage_test.cu b/tests/utility/storage_test.cu
@@ -39,20 +39,18 @@ TEMPLATE_TEST_CASE_SIG("Storage tests",
 
   SECTION("Initialize empty storage is allowed.")
   {
-    auto s = cuco::bucket_storage<cuco::pair<Key, Value>,
-                                  bucket_size,
-                                  cuco::extent<std::size_t>,
-                                  allocator_type>{cuco::extent<std::size_t>{0}, allocator};
+    auto s = cuco::
+      aow_storage<cuco::pair<Key, Value>, bucket_size, cuco::extent<std::size_t>, allocator_type>{
+        cuco::extent<std::size_t>{0}, allocator};
 
     s.initialize(cuco::pair<Key, Value>{1, 1});
   }
 
   SECTION("Allocate array of pairs with AoS storage.")
   {
-    auto s                 = cuco::bucket_storage<cuco::pair<Key, Value>,
-                                                  bucket_size,
-                                                  cuco::extent<std::size_t>,
-                                                  allocator_type>(cuco::extent{size}, allocator);
+    auto s = cuco::
+      aow_storage<cuco::pair<Key, Value>, bucket_size, cuco::extent<std::size_t>, allocator_type>(
+        cuco::extent{size}, allocator);
     auto const num_buckets = s.num_buckets();
     auto const capacity    = s.capacity();
 
@@ -63,7 +61,7 @@ TEMPLATE_TEST_CASE_SIG("Storage tests",
   SECTION("Allocate array of pairs with AoS storage with static extent.")
   {
     using extent_type = cuco::extent<std::size_t, size>;
-    auto s = cuco::bucket_storage<cuco::pair<Key, Value>, bucket_size, extent_type, allocator_type>(
+    auto s = cuco::aow_storage<cuco::pair<Key, Value>, bucket_size, extent_type, allocator_type>(
       extent_type{}, allocator);
     auto const num_buckets = s.num_buckets();
     auto const capacity    = s.capacity();
@@ -74,7 +72,7 @@ TEMPLATE_TEST_CASE_SIG("Storage tests",
 
   SECTION("Allocate array of keys with AoS storage.")
   {
-    auto s = cuco::bucket_storage<Key, bucket_size, cuco::extent<std::size_t>, allocator_type>(
+    auto s = cuco::aow_storage<Key, bucket_size, cuco::extent<std::size_t>, allocator_type>(
       cuco::extent{size}, allocator);
     auto const num_buckets = s.num_buckets();
     auto const capacity    = s.capacity();
@@ -87,7 +85,7 @@ TEMPLATE_TEST_CASE_SIG("Storage tests",
   {
     using extent_type = cuco::extent<std::size_t, size>;
     auto s =
-      cuco::bucket_storage<Key, bucket_size, extent_type, allocator_type>(extent_type{}, allocator);
+      cuco::aow_storage<Key, bucket_size, extent_type, allocator_type>(extent_type{}, allocator);
     auto const num_buckets = s.num_buckets();
     auto const capacity    = s.capacity();